* [igt-dev] [PATCH i-g-t v2 0/4] Make xe_compute test more generic
@ 2023-04-04 7:38 Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 1/4] compute_square_kernel.cl: add CL file used at xe_compute.c Mauro Carvalho Chehab
` (5 more replies)
0 siblings, 6 replies; 8+ messages in thread
From: Mauro Carvalho Chehab @ 2023-04-04 7:38 UTC (permalink / raw)
To: igt-dev
From: Mauro Carvalho Chehab <mchehab@kernel.org>
Currently, xe_compute test runs only on an specific TGL PCI ID.
Make the test more generic, preparing to add support for more
platforms.
---
v2:
- removed the RFC example patch:
https://patchwork.freedesktop.org/patch/528910/?series=115670&rev=1
(no need to re-send, as this is just an example that won't be
merged)
- run_xe_compute_kernel() now returns true if suceeded, false
otherwise;
- opencl/gen_opencl_kernel is now compatible with legacy versions
of xxd;
- added some notes at opencl/README about Intel compute mode
versions.
Mauro Carvalho Chehab (4):
compute_square_kernel.cl: add CL file used at xe_compute.c
xe/xe_compute: place OpenCL kernel on a separate file
lib/xe/xe_compute: use registers defs from intel_gpu_commands.h
gen_opencl_kernel: add script to dynamically create OpenCL kernels
lib/meson.build | 1 +
lib/xe/xe_compute.c | 236 +++++++++++++++++++----------
lib/xe/xe_compute.h | 31 ++--
lib/xe/xe_compute_square_kernels.c | 71 +++++++++
opencl/README | 30 ++++
opencl/compute_square_kernel.cl | 5 +
opencl/gen_opencl_kernel | 103 +++++++++++++
tests/xe/xe_compute.c | 108 +------------
8 files changed, 391 insertions(+), 194 deletions(-)
create mode 100644 lib/xe/xe_compute_square_kernels.c
create mode 100644 opencl/README
create mode 100644 opencl/compute_square_kernel.cl
create mode 100755 opencl/gen_opencl_kernel
--
2.39.2
^ permalink raw reply [flat|nested] 8+ messages in thread
* [igt-dev] [PATCH i-g-t v2 1/4] compute_square_kernel.cl: add CL file used at xe_compute.c
2023-04-04 7:38 [igt-dev] [PATCH i-g-t v2 0/4] Make xe_compute test more generic Mauro Carvalho Chehab
@ 2023-04-04 7:38 ` Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 2/4] xe/xe_compute: place OpenCL kernel on a separate file Mauro Carvalho Chehab
` (4 subsequent siblings)
5 siblings, 0 replies; 8+ messages in thread
From: Mauro Carvalho Chehab @ 2023-04-04 7:38 UTC (permalink / raw)
To: igt-dev
From: Mauro Carvalho Chehab <mchehab@kernel.org>
Provide the cl file that it was used to produce the OpenCL
Kernel used by xe_compute, and document how the binary at
xe_compute.h was produced.
Reviewed-by: Zbigniew Kempczyński <zbigniew.kempczynski@intel.com>
Signed-off-by: Mauro Carvalho Chehab <mchehab@kernel.org>
---
lib/xe/xe_compute.c | 3 ++-
opencl/compute_square_kernel.cl | 5 +++++
2 files changed, 7 insertions(+), 1 deletion(-)
create mode 100644 opencl/compute_square_kernel.cl
diff --git a/lib/xe/xe_compute.c b/lib/xe/xe_compute.c
index 8c0f8c87d50f..2165eada8931 100644
--- a/lib/xe/xe_compute.c
+++ b/lib/xe/xe_compute.c
@@ -18,7 +18,8 @@
#define GPGPU_WALKER 0x7105000d
#define MI_BATCH_BUFFER_END (0xA << 23)
-// generated with: ocloc -file kernel.cl -device tgllp && xxd -i kernel_Gen12LPlp.gen
+// generated with:
+// ocloc -file opencl/compute_square_kernel.cl -device tgllp && xxd -i compute_square_kernel_Gen12LPlp.bin
unsigned char tgllp_kernel_square_bin[] = {
0x61, 0x00, 0x03, 0x80, 0x20, 0x02, 0x05, 0x03, 0x04, 0x00, 0x10, 0x00,
0x00, 0x00, 0x00, 0x00, 0x66, 0x01, 0x00, 0x80, 0x20, 0x82, 0x01, 0x80,
diff --git a/opencl/compute_square_kernel.cl b/opencl/compute_square_kernel.cl
new file mode 100644
index 000000000000..f6260fb934dc
--- /dev/null
+++ b/opencl/compute_square_kernel.cl
@@ -0,0 +1,5 @@
+__kernel void square(__global float* input, __global float* output, const unsigned int count) {
+ int i = get_global_id(0);
+ if(i < count)
+ output[i] = input[i] * input[i];
+}
--
2.39.2
^ permalink raw reply related [flat|nested] 8+ messages in thread
* [igt-dev] [PATCH i-g-t v2 2/4] xe/xe_compute: place OpenCL kernel on a separate file
2023-04-04 7:38 [igt-dev] [PATCH i-g-t v2 0/4] Make xe_compute test more generic Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 1/4] compute_square_kernel.cl: add CL file used at xe_compute.c Mauro Carvalho Chehab
@ 2023-04-04 7:38 ` Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 3/4] lib/xe/xe_compute: use registers defs from intel_gpu_commands.h Mauro Carvalho Chehab
` (3 subsequent siblings)
5 siblings, 0 replies; 8+ messages in thread
From: Mauro Carvalho Chehab @ 2023-04-04 7:38 UTC (permalink / raw)
To: igt-dev
From: Mauro Carvalho Chehab <mchehab@kernel.org>
In order to prepare for supporting multiple Kernels, move
the tgllp to a separate file.
While here, address a few coding style nitpicks.
Reviewed-by: Zbigniew Kempczyński <zbigniew.kempczynski@intel.com>
Signed-off-by: Mauro Carvalho Chehab <mchehab@kernel.org>
---
lib/meson.build | 1 +
lib/xe/xe_compute.c | 234 ++++++++++++++++++++---------
lib/xe/xe_compute.h | 31 ++--
lib/xe/xe_compute_square_kernels.c | 71 +++++++++
tests/xe/xe_compute.c | 108 +------------
5 files changed, 256 insertions(+), 189 deletions(-)
create mode 100644 lib/xe/xe_compute_square_kernels.c
diff --git a/lib/meson.build b/lib/meson.build
index ad9e2abef4c3..ad68089dcf43 100644
--- a/lib/meson.build
+++ b/lib/meson.build
@@ -99,6 +99,7 @@ lib_sources = [
'igt_msm.c',
'igt_dsc.c',
'xe/xe_compute.c',
+ 'xe/xe_compute_square_kernels.c',
'xe/xe_ioctl.c',
'xe/xe_query.c',
'xe/xe_spin.c'
diff --git a/lib/xe/xe_compute.c b/lib/xe/xe_compute.c
index 2165eada8931..fb11b8bc7770 100644
--- a/lib/xe/xe_compute.c
+++ b/lib/xe/xe_compute.c
@@ -6,71 +6,51 @@
* Francois Dugast <francois.dugast@intel.com>
*/
+#include <stdint.h>
+
+#include "igt.h"
+#include "xe_drm.h"
+#include "lib/igt_syncobj.h"
+#include "lib/intel_reg.h"
+
#include "xe_compute.h"
+#include "xe/xe_ioctl.h"
+#include "xe/xe_query.h"
#define PIPE_CONTROL 0x7a000004
-#define MI_LOAD_REGISTER_IMM 0x11000001
-#define PIPELINE_SELECT 0x69040302
+#define MEDIA_STATE_FLUSH 0x0
+#define MAX(X, Y) (((X) > (Y)) ? (X) : (Y))
+#define SIZE_DATA 64
+#define SIZE_BATCH 0x1000
+#define SIZE_BUFFER_INPUT MAX(sizeof(float) * SIZE_DATA, 0x1000)
+#define SIZE_BUFFER_OUTPUT MAX(sizeof(float) * SIZE_DATA, 0x1000)
+#define ADDR_BATCH 0x100000
+#define ADDR_INPUT 0x200000UL
+#define ADDR_OUTPUT 0x300000UL
+#define ADDR_SURFACE_STATE_BASE 0x400000UL
+#define ADDR_DYNAMIC_STATE_BASE 0x500000UL
+#define ADDR_INDIRECT_OBJECT_BASE 0x800100000000
+#define OFFSET_INDIRECT_DATA_START 0xFFFDF000
+#define OFFSET_KERNEL 0xFFFEF000
+
+#undef MEDIA_VFE_STATE
#define MEDIA_VFE_STATE 0x70000007
+#undef STATE_BASE_ADDRESS
#define STATE_BASE_ADDRESS 0x61010014
-#define MEDIA_STATE_FLUSH 0x0
+#undef MEDIA_INTERFACE_DESCRIPTOR_LOAD
#define MEDIA_INTERFACE_DESCRIPTOR_LOAD 0x70020002
+#undef GPGPU_WALKER
#define GPGPU_WALKER 0x7105000d
-#define MI_BATCH_BUFFER_END (0xA << 23)
-
-// generated with:
-// ocloc -file opencl/compute_square_kernel.cl -device tgllp && xxd -i compute_square_kernel_Gen12LPlp.bin
-unsigned char tgllp_kernel_square_bin[] = {
- 0x61, 0x00, 0x03, 0x80, 0x20, 0x02, 0x05, 0x03, 0x04, 0x00, 0x10, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x66, 0x01, 0x00, 0x80, 0x20, 0x82, 0x01, 0x80,
- 0x00, 0x80, 0x00, 0x01, 0xc0, 0x04, 0xc0, 0x04, 0x41, 0x01, 0x20, 0x22,
- 0x16, 0x09, 0x11, 0x03, 0x49, 0x00, 0x04, 0xa2, 0x12, 0x09, 0x11, 0x03,
- 0x40, 0x01, 0x04, 0x00, 0x60, 0x06, 0x05, 0x05, 0x04, 0x04, 0x00, 0x01,
- 0x05, 0x01, 0x58, 0x00, 0x40, 0x00, 0x24, 0x00, 0x60, 0x06, 0x05, 0x0a,
- 0x04, 0x04, 0x00, 0x01, 0x05, 0x02, 0x58, 0x00, 0x40, 0x02, 0x0c, 0xa0,
- 0x02, 0x05, 0x10, 0x07, 0x40, 0x02, 0x0e, 0xa6, 0x02, 0x0a, 0x10, 0x07,
- 0x70, 0x02, 0x04, 0x00, 0x60, 0x02, 0x01, 0x00, 0x05, 0x0c, 0x46, 0x52,
- 0x84, 0x08, 0x00, 0x00, 0x70, 0x02, 0x24, 0x00, 0x60, 0x02, 0x01, 0x00,
- 0x05, 0x0e, 0x46, 0x52, 0x84, 0x08, 0x00, 0x00, 0x72, 0x00, 0x02, 0x80,
- 0x50, 0x0d, 0x04, 0x00, 0x05, 0x00, 0x05, 0x1d, 0x05, 0x00, 0x05, 0x00,
- 0x22, 0x00, 0x05, 0x01, 0x00, 0xc0, 0x00, 0x00, 0x90, 0x00, 0x00, 0x00,
- 0x90, 0x00, 0x00, 0x00, 0x69, 0x00, 0x10, 0x60, 0x02, 0x0c, 0x20, 0x00,
- 0x69, 0x00, 0x12, 0x66, 0x02, 0x0e, 0x20, 0x00, 0x40, 0x02, 0x14, 0xa0,
- 0x32, 0x10, 0x10, 0x08, 0x40, 0x02, 0x16, 0xa6, 0x32, 0x12, 0x10, 0x08,
- 0x31, 0xa0, 0x04, 0x00, 0x00, 0x00, 0x14, 0x18, 0x14, 0x14, 0x00, 0xcc,
- 0x00, 0x00, 0x16, 0x00, 0x31, 0x91, 0x24, 0x00, 0x00, 0x00, 0x14, 0x1a,
- 0x14, 0x16, 0x00, 0xcc, 0x00, 0x00, 0x16, 0x00, 0x40, 0x00, 0x10, 0xa0,
- 0x4a, 0x10, 0x10, 0x08, 0x40, 0x00, 0x12, 0xa6, 0x4a, 0x12, 0x10, 0x08,
- 0x41, 0x20, 0x18, 0x20, 0x00, 0x18, 0x00, 0x18, 0x41, 0x21, 0x1a, 0x26,
- 0x00, 0x1a, 0x00, 0x1a, 0x31, 0xa2, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x14, 0x10, 0x02, 0xcc, 0x14, 0x18, 0x96, 0x00, 0x31, 0x93, 0x24, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x14, 0x12, 0x02, 0xcc, 0x14, 0x1a, 0x96, 0x00,
- 0x25, 0x00, 0x05, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x10, 0x00, 0x00, 0x00, 0x61, 0x00, 0x7f, 0x64, 0x00, 0x03, 0x10, 0x00,
- 0x31, 0x44, 0x03, 0x80, 0x00, 0x00, 0x0c, 0x1c, 0x0c, 0x03, 0x00, 0xa0,
- 0x00, 0x00, 0x78, 0x02, 0x61, 0x24, 0x03, 0x80, 0x20, 0x02, 0x01, 0x00,
- 0x05, 0x1c, 0x46, 0x00, 0x00, 0x00, 0x00, 0x00, 0x61, 0x00, 0x04, 0x80,
- 0xa0, 0x4a, 0x01, 0x20, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x31, 0x01, 0x03, 0x80, 0x04, 0x00, 0x00, 0x00, 0x0c, 0x7f, 0x20, 0x70,
- 0x00, 0x00, 0x00, 0x00, 0x60, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
- 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00
+
+struct bo_dict_entry {
+ uint64_t addr;
+ uint32_t size;
+ void *data;
};
-unsigned int tgllp_kernel_square_length = sizeof(tgllp_kernel_square_bin);
+
+/*
+ * TGL compatible batch
+ */
/**
* tgllp_create_indirect_data:
@@ -80,8 +60,9 @@ unsigned int tgllp_kernel_square_length = sizeof(tgllp_kernel_square_bin);
*
* Prepares indirect data for compute pipeline.
*/
-void tgllp_create_indirect_data(uint32_t *addr_bo_buffer_batch,
- uint64_t addr_input, uint64_t addr_output)
+static void tgllp_create_indirect_data(uint32_t *addr_bo_buffer_batch,
+ uint64_t addr_input,
+ uint64_t addr_output)
{
int b = 0;
@@ -183,8 +164,9 @@ void tgllp_create_indirect_data(uint32_t *addr_bo_buffer_batch,
*
* Prepares surface state for compute pipeline.
*/
-void tgllp_create_surface_state(uint32_t *addr_bo_buffer_batch,
- uint64_t addr_input, uint64_t addr_output)
+static void tgllp_create_surface_state(uint32_t *addr_bo_buffer_batch,
+ uint64_t addr_input,
+ uint64_t addr_output)
{
int b = 0;
@@ -261,8 +243,8 @@ void tgllp_create_surface_state(uint32_t *addr_bo_buffer_batch,
*
* Prepares dynamic state for compute pipeline.
*/
-void tgllp_create_dynamic_state(uint32_t *addr_bo_buffer_batch,
- uint64_t offset_kernel)
+static void tgllp_create_dynamic_state(uint32_t *addr_bo_buffer_batch,
+ uint64_t offset_kernel)
{
int b = 0;
@@ -280,7 +262,7 @@ void tgllp_create_dynamic_state(uint32_t *addr_bo_buffer_batch,
}
/**
- * tgllp_create_batch_compute:
+ * tgllp_compute_exec_compute:
* @addr_bo_buffer_batch: pointer to batch buffer
* @addr_surface_state_base: gpu offset of surface state data
* @addr_dynamic_state_base: gpu offset of dynamic state data
@@ -289,19 +271,19 @@ void tgllp_create_dynamic_state(uint32_t *addr_bo_buffer_batch,
*
* Prepares compute pipeline.
*/
-void tgllp_create_batch_compute(uint32_t *addr_bo_buffer_batch,
- uint64_t addr_surface_state_base,
- uint64_t addr_dynamic_state_base,
- uint64_t addr_indirect_object_base,
- uint64_t offset_indirect_data_start)
+static void tgllp_compute_exec_compute(uint32_t *addr_bo_buffer_batch,
+ uint64_t addr_surface_state_base,
+ uint64_t addr_dynamic_state_base,
+ uint64_t addr_indirect_object_base,
+ uint64_t offset_indirect_data_start)
{
int b = 0;
- addr_bo_buffer_batch[b++] = MI_LOAD_REGISTER_IMM;
+ addr_bo_buffer_batch[b++] = MI_LOAD_REGISTER_IMM(1);
addr_bo_buffer_batch[b++] = 0x00002580;
addr_bo_buffer_batch[b++] = 0x00060002;
addr_bo_buffer_batch[b++] = PIPELINE_SELECT;
- addr_bo_buffer_batch[b++] = MI_LOAD_REGISTER_IMM;
+ addr_bo_buffer_batch[b++] = MI_LOAD_REGISTER_IMM(1);
addr_bo_buffer_batch[b++] = 0x00007034;
addr_bo_buffer_batch[b++] = 0x60000321;
addr_bo_buffer_batch[b++] = PIPE_CONTROL;
@@ -310,7 +292,7 @@ void tgllp_create_batch_compute(uint32_t *addr_bo_buffer_batch,
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00000000;
- addr_bo_buffer_batch[b++] = MI_LOAD_REGISTER_IMM;
+ addr_bo_buffer_batch[b++] = MI_LOAD_REGISTER_IMM(1);
addr_bo_buffer_batch[b++] = 0x0000E404;
addr_bo_buffer_batch[b++] = 0x00000100;
addr_bo_buffer_batch[b++] = PIPE_CONTROL;
@@ -405,3 +387,111 @@ void tgllp_create_batch_compute(uint32_t *addr_bo_buffer_batch,
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = MI_BATCH_BUFFER_END;
}
+
+/**
+ * tgl_compute_exec - run a pipeline compatible with Tiger Lake
+ *
+ * @fd: file descriptor of the opened DRM device
+ * @kernel: GPU Kernel binary to be executed
+ * @size: size of @kernel.
+ */
+static void tgl_compute_exec(int fd, const unsigned char *kernel,
+ unsigned int size)
+{
+ uint32_t vm, engine;
+ float *dinput;
+ struct drm_xe_sync sync = { 0 };
+#define TGL_BO_DICT_ENTRIES 7
+ struct bo_dict_entry bo_dict[TGL_BO_DICT_ENTRIES] = {
+ { .addr = ADDR_INDIRECT_OBJECT_BASE + OFFSET_KERNEL}, // kernel
+ { .addr = ADDR_DYNAMIC_STATE_BASE, .size = 0x1000}, // dynamic state
+ { .addr = ADDR_SURFACE_STATE_BASE, .size = 0x1000}, // surface state
+ { .addr = ADDR_INDIRECT_OBJECT_BASE + OFFSET_INDIRECT_DATA_START, .size = 0x10000}, // indirect data
+ { .addr = ADDR_INPUT, .size = SIZE_BUFFER_INPUT }, // input
+ { .addr = ADDR_OUTPUT, .size = SIZE_BUFFER_OUTPUT }, // output
+ { .addr = ADDR_BATCH, .size = SIZE_BATCH }, // batch
+ };
+
+ /* Sets Kernel size */
+ bo_dict[0].size = ALIGN(size, 0x1000);
+
+ vm = xe_vm_create(fd, DRM_XE_VM_CREATE_ASYNC_BIND_OPS, 0);
+ engine = xe_engine_create_class(fd, vm, DRM_XE_ENGINE_CLASS_RENDER);
+ sync.flags = DRM_XE_SYNC_SYNCOBJ | DRM_XE_SYNC_SIGNAL;
+ sync.handle = syncobj_create(fd, 0);
+
+ for (int i = 0; i < TGL_BO_DICT_ENTRIES; i++) {
+ bo_dict[i].data = aligned_alloc(xe_get_default_alignment(fd), bo_dict[i].size);
+ xe_vm_bind_userptr_async(fd, vm, 0, to_user_pointer(bo_dict[i].data), bo_dict[i].addr, bo_dict[i].size, &sync, 1);
+ syncobj_wait(fd, &sync.handle, 1, INT64_MAX, 0, NULL);
+ memset(bo_dict[i].data, 0, bo_dict[i].size);
+ }
+ memcpy(bo_dict[0].data, kernel, size);
+ tgllp_create_dynamic_state(bo_dict[1].data, OFFSET_KERNEL);
+ tgllp_create_surface_state(bo_dict[2].data, ADDR_INPUT, ADDR_OUTPUT);
+ tgllp_create_indirect_data(bo_dict[3].data, ADDR_INPUT, ADDR_OUTPUT);
+ dinput = (float *)bo_dict[4].data;
+ srand(time(NULL));
+
+ for (int i = 0; i < SIZE_DATA; i++)
+ ((float *)dinput)[i] = rand() / (float)RAND_MAX;
+
+ tgllp_compute_exec_compute(bo_dict[6].data, ADDR_SURFACE_STATE_BASE, ADDR_DYNAMIC_STATE_BASE, ADDR_INDIRECT_OBJECT_BASE, OFFSET_INDIRECT_DATA_START);
+
+ xe_exec_wait(fd, engine, ADDR_BATCH);
+
+ for (int i = 0; i < SIZE_DATA; i++)
+ igt_assert(((float *)bo_dict[5].data)[i] == ((float *)bo_dict[4].data)[i] * ((float *) bo_dict[4].data)[i]);
+
+ for (int i = 0; i < TGL_BO_DICT_ENTRIES; i++) {
+ xe_vm_unbind_async(fd, vm, 0, 0, bo_dict[i].addr, bo_dict[i].size, &sync, 1);
+ syncobj_wait(fd, &sync.handle, 1, INT64_MAX, 0, NULL);
+ free(bo_dict[i].data);
+ }
+
+ syncobj_destroy(fd, sync.handle);
+ xe_engine_destroy(fd, engine);
+ xe_vm_destroy(fd, vm);
+}
+
+/*
+ * Generic code
+ */
+
+static const struct {
+ unsigned int ip_ver;
+ void (*compute_exec)(int fd, const unsigned char *kernel,
+ unsigned int size);
+} xe_compute_batches[] = {
+ {
+ .ip_ver = IP_VER(12, 0),
+ .compute_exec = tgl_compute_exec,
+ },
+};
+
+bool run_xe_compute_kernel(int fd)
+{
+ unsigned int ip_ver = intel_graphics_ver(intel_get_drm_devid(fd));
+ unsigned int batch;
+ const struct xe_compute_kernels *kernels = xe_compute_square_kernels;
+
+ for (batch = 0; batch < ARRAY_SIZE(xe_compute_batches); batch++) {
+ if (ip_ver == xe_compute_batches[batch].ip_ver)
+ break;
+ }
+ if (batch == ARRAY_SIZE(xe_compute_batches))
+ return false;
+
+ while (kernels->kernel) {
+ if (ip_ver == kernels->ip_ver)
+ break;
+ kernels++;
+ }
+ if (!kernels->kernel)
+ return 1;
+
+ xe_compute_batches[batch].compute_exec(fd, kernels->kernel,
+ kernels->size);
+
+ return true;
+}
diff --git a/lib/xe/xe_compute.h b/lib/xe/xe_compute.h
index de763101da90..b2e7e9827836 100644
--- a/lib/xe/xe_compute.h
+++ b/lib/xe/xe_compute.h
@@ -9,21 +9,24 @@
#ifndef XE_COMPUTE_H
#define XE_COMPUTE_H
-#include <stdint.h>
+/*
+ * OpenCL Kernels are generated using:
+ *
+ * GPU=tgllp && \
+ * ocloc -file opencl/compute_square_kernel.cl -device $GPU && \
+ * xxd -i compute_square_kernel_Gen12LPlp.bin
+ *
+ * For each GPU model desired. A list of supported models can be obtained with: ocloc compile --help
+ */
+
+struct xe_compute_kernels {
+ int ip_ver;
+ unsigned int size;
+ const unsigned char *kernel;
+};
-void tgllp_create_indirect_data(uint32_t *addr_bo_buffer_batch,
- uint64_t addr_input, uint64_t addr_output);
-void tgllp_create_surface_state(uint32_t *addr_bo_buffer_batch,
- uint64_t addr_input, uint64_t addr_output);
-void tgllp_create_dynamic_state(uint32_t *addr_bo_buffer_batch,
- uint64_t offset_kernel);
-void tgllp_create_batch_compute(uint32_t *addr_bo_buffer_batch,
- uint64_t addr_surface_state_base,
- uint64_t addr_dynamic_state_base,
- uint64_t addr_indirect_object_base,
- uint64_t offset_indirect_data_start);
+extern const struct xe_compute_kernels xe_compute_square_kernels[];
-extern unsigned char tgllp_kernel_square_bin[];
-extern unsigned int tgllp_kernel_square_length;
+bool run_xe_compute_kernel(int fd);
#endif /* XE_COMPUTE_H */
diff --git a/lib/xe/xe_compute_square_kernels.c b/lib/xe/xe_compute_square_kernels.c
new file mode 100644
index 000000000000..f9c07dc778bd
--- /dev/null
+++ b/lib/xe/xe_compute_square_kernels.c
@@ -0,0 +1,71 @@
+/* SPDX-License-Identifier: MIT */
+
+/*
+ * Copyright © 2022 Intel Corporation
+ *
+ * Authors:
+ * Francois Dugast <francois.dugast@intel.com>
+ */
+
+#include "intel_chipset.h"
+#include "lib/xe/xe_compute.h"
+
+static const unsigned char tgllp_kernel_square_bin[] = {
+ 0x61, 0x00, 0x03, 0x80, 0x20, 0x02, 0x05, 0x03, 0x04, 0x00, 0x10, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x66, 0x01, 0x00, 0x80, 0x20, 0x82, 0x01, 0x80,
+ 0x00, 0x80, 0x00, 0x01, 0xc0, 0x04, 0xc0, 0x04, 0x41, 0x01, 0x20, 0x22,
+ 0x16, 0x09, 0x11, 0x03, 0x49, 0x00, 0x04, 0xa2, 0x12, 0x09, 0x11, 0x03,
+ 0x40, 0x01, 0x04, 0x00, 0x60, 0x06, 0x05, 0x05, 0x04, 0x04, 0x00, 0x01,
+ 0x05, 0x01, 0x58, 0x00, 0x40, 0x00, 0x24, 0x00, 0x60, 0x06, 0x05, 0x0a,
+ 0x04, 0x04, 0x00, 0x01, 0x05, 0x02, 0x58, 0x00, 0x40, 0x02, 0x0c, 0xa0,
+ 0x02, 0x05, 0x10, 0x07, 0x40, 0x02, 0x0e, 0xa6, 0x02, 0x0a, 0x10, 0x07,
+ 0x70, 0x02, 0x04, 0x00, 0x60, 0x02, 0x01, 0x00, 0x05, 0x0c, 0x46, 0x52,
+ 0x84, 0x08, 0x00, 0x00, 0x70, 0x02, 0x24, 0x00, 0x60, 0x02, 0x01, 0x00,
+ 0x05, 0x0e, 0x46, 0x52, 0x84, 0x08, 0x00, 0x00, 0x72, 0x00, 0x02, 0x80,
+ 0x50, 0x0d, 0x04, 0x00, 0x05, 0x00, 0x05, 0x1d, 0x05, 0x00, 0x05, 0x00,
+ 0x22, 0x00, 0x05, 0x01, 0x00, 0xc0, 0x00, 0x00, 0x90, 0x00, 0x00, 0x00,
+ 0x90, 0x00, 0x00, 0x00, 0x69, 0x00, 0x10, 0x60, 0x02, 0x0c, 0x20, 0x00,
+ 0x69, 0x00, 0x12, 0x66, 0x02, 0x0e, 0x20, 0x00, 0x40, 0x02, 0x14, 0xa0,
+ 0x32, 0x10, 0x10, 0x08, 0x40, 0x02, 0x16, 0xa6, 0x32, 0x12, 0x10, 0x08,
+ 0x31, 0xa0, 0x04, 0x00, 0x00, 0x00, 0x14, 0x18, 0x14, 0x14, 0x00, 0xcc,
+ 0x00, 0x00, 0x16, 0x00, 0x31, 0x91, 0x24, 0x00, 0x00, 0x00, 0x14, 0x1a,
+ 0x14, 0x16, 0x00, 0xcc, 0x00, 0x00, 0x16, 0x00, 0x40, 0x00, 0x10, 0xa0,
+ 0x4a, 0x10, 0x10, 0x08, 0x40, 0x00, 0x12, 0xa6, 0x4a, 0x12, 0x10, 0x08,
+ 0x41, 0x20, 0x18, 0x20, 0x00, 0x18, 0x00, 0x18, 0x41, 0x21, 0x1a, 0x26,
+ 0x00, 0x1a, 0x00, 0x1a, 0x31, 0xa2, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x14, 0x10, 0x02, 0xcc, 0x14, 0x18, 0x96, 0x00, 0x31, 0x93, 0x24, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x14, 0x12, 0x02, 0xcc, 0x14, 0x1a, 0x96, 0x00,
+ 0x25, 0x00, 0x05, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x10, 0x00, 0x00, 0x00, 0x61, 0x00, 0x7f, 0x64, 0x00, 0x03, 0x10, 0x00,
+ 0x31, 0x44, 0x03, 0x80, 0x00, 0x00, 0x0c, 0x1c, 0x0c, 0x03, 0x00, 0xa0,
+ 0x00, 0x00, 0x78, 0x02, 0x61, 0x24, 0x03, 0x80, 0x20, 0x02, 0x01, 0x00,
+ 0x05, 0x1c, 0x46, 0x00, 0x00, 0x00, 0x00, 0x00, 0x61, 0x00, 0x04, 0x80,
+ 0xa0, 0x4a, 0x01, 0x20, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x31, 0x01, 0x03, 0x80, 0x04, 0x00, 0x00, 0x00, 0x0c, 0x7f, 0x20, 0x70,
+ 0x00, 0x00, 0x00, 0x00, 0x60, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00
+};
+
+const struct xe_compute_kernels xe_compute_square_kernels[] = {
+ {
+ .ip_ver = IP_VER(12, 0),
+ .size = sizeof(tgllp_kernel_square_bin),
+ .kernel = tgllp_kernel_square_bin,
+ },
+ {}
+};
diff --git a/tests/xe/xe_compute.c b/tests/xe/xe_compute.c
index 138d80671435..7ac64dfe3199 100644
--- a/tests/xe/xe_compute.c
+++ b/tests/xe/xe_compute.c
@@ -14,117 +14,21 @@
#include <string.h>
#include "igt.h"
-#include "lib/igt_syncobj.h"
-#include "xe_drm.h"
-#include "xe/xe_ioctl.h"
#include "xe/xe_query.h"
#include "xe/xe_compute.h"
-#define MAX(X, Y) (((X) > (Y)) ? (X) : (Y))
-#define SIZE_DATA 64
-#define SIZE_BATCH 0x1000
-#define SIZE_KERNEL 0x1000
-#define SIZE_BUFFER_INPUT MAX(sizeof(float)*SIZE_DATA, 0x1000)
-#define SIZE_BUFFER_OUTPUT MAX(sizeof(float)*SIZE_DATA, 0x1000)
-#define ADDR_BATCH 0x100000
-#define ADDR_INPUT (unsigned long)0x200000
-#define ADDR_OUTPUT (unsigned long)0x300000
-#define ADDR_SURFACE_STATE_BASE (unsigned long)0x400000
-#define ADDR_DYNAMIC_STATE_BASE (unsigned long)0x500000
-#define ADDR_INDIRECT_OBJECT_BASE 0x800100000000
-#define OFFSET_INDIRECT_DATA_START 0xFFFDF000
-#define OFFSET_KERNEL 0xFFFEF000
-
-struct bo_dict_entry {
- uint64_t addr;
- uint32_t size;
- void *data;
-};
-
/**
* SUBTEST: compute-square
- * GPU requirement: only works on TGL_GT2 with device ID: 0x9a49
+ * GPU requirement: only works on TGL
* Description:
- * This test shows how to create a batch to execute a
- * compute kernel. For now it supports tgllp only.
+ * Run an openCL Kernel that returns output[i] = input[i] * input[i],
+ * for an input dataset..
* TODO: extend test to cover other platforms
*/
static void
test_compute_square(int fd)
{
- uint32_t vm, engine;
- float *dinput;
- struct drm_xe_sync sync = { 0 };
-
-#define BO_DICT_ENTRIES 7
- struct bo_dict_entry bo_dict[BO_DICT_ENTRIES] = {
- { .addr = ADDR_INDIRECT_OBJECT_BASE + OFFSET_KERNEL, .size = SIZE_KERNEL }, // kernel
- { .addr = ADDR_DYNAMIC_STATE_BASE, .size = 0x1000}, // dynamic state
- { .addr = ADDR_SURFACE_STATE_BASE, .size = 0x1000}, // surface state
- { .addr = ADDR_INDIRECT_OBJECT_BASE + OFFSET_INDIRECT_DATA_START, .size = 0x10000}, // indirect data
- { .addr = ADDR_INPUT, .size = SIZE_BUFFER_INPUT }, // input
- { .addr = ADDR_OUTPUT, .size = SIZE_BUFFER_OUTPUT }, // output
- { .addr = ADDR_BATCH, .size = SIZE_BATCH }, // batch
- };
-
- vm = xe_vm_create(fd, DRM_XE_VM_CREATE_ASYNC_BIND_OPS, 0);
- engine = xe_engine_create_class(fd, vm, DRM_XE_ENGINE_CLASS_RENDER);
- sync.flags = DRM_XE_SYNC_SYNCOBJ | DRM_XE_SYNC_SIGNAL;
- sync.handle = syncobj_create(fd, 0);
-
- for(int i = 0; i < BO_DICT_ENTRIES; i++) {
- bo_dict[i].data = aligned_alloc(xe_get_default_alignment(fd), bo_dict[i].size);
- xe_vm_bind_userptr_async(fd, vm, 0, to_user_pointer(bo_dict[i].data), bo_dict[i].addr, bo_dict[i].size, &sync, 1);
- syncobj_wait(fd, &sync.handle, 1, INT64_MAX, 0, NULL);
- memset(bo_dict[i].data, 0, bo_dict[i].size);
- }
- memcpy(bo_dict[0].data, tgllp_kernel_square_bin, tgllp_kernel_square_length);
- tgllp_create_dynamic_state(bo_dict[1].data, OFFSET_KERNEL);
- tgllp_create_surface_state(bo_dict[2].data, ADDR_INPUT, ADDR_OUTPUT);
- tgllp_create_indirect_data(bo_dict[3].data, ADDR_INPUT, ADDR_OUTPUT);
- dinput = (float *)bo_dict[4].data;
- srand(time(NULL));
- for(int i=0; i < SIZE_DATA; i++) {
- ((float*) dinput)[i] = rand()/(float)RAND_MAX;
- }
- tgllp_create_batch_compute(bo_dict[6].data, ADDR_SURFACE_STATE_BASE, ADDR_DYNAMIC_STATE_BASE, ADDR_INDIRECT_OBJECT_BASE, OFFSET_INDIRECT_DATA_START);
-
- xe_exec_wait(fd, engine, ADDR_BATCH);
- for(int i = 0; i < SIZE_DATA; i++) {
- igt_assert(((float*) bo_dict[5].data)[i] == ((float*) bo_dict[4].data)[i] * ((float*) bo_dict[4].data)[i]);
- }
-
- for(int i = 0; i < BO_DICT_ENTRIES; i++) {
- xe_vm_unbind_async(fd, vm, 0, 0, bo_dict[i].addr, bo_dict[i].size, &sync, 1);
- syncobj_wait(fd, &sync.handle, 1, INT64_MAX, 0, NULL);
- free(bo_dict[i].data);
- }
-
- syncobj_destroy(fd, sync.handle);
- xe_engine_destroy(fd, engine);
- xe_vm_destroy(fd, vm);
-}
-
-static bool
-is_device_supported(int fd)
-{
- struct drm_xe_query_config *config;
- struct drm_xe_device_query query = {
- .extensions = 0,
- .query = DRM_XE_DEVICE_QUERY_CONFIG,
- .size = 0,
- .data = 0,
- };
-
- igt_assert_eq(igt_ioctl(fd, DRM_IOCTL_XE_DEVICE_QUERY, &query), 0);
-
- config = malloc(query.size);
- igt_assert(config);
-
- query.data = to_user_pointer(config);
- igt_assert_eq(igt_ioctl(fd, DRM_IOCTL_XE_DEVICE_QUERY, &query), 0);
-
- return (config->info[XE_QUERY_CONFIG_REV_AND_DEVICE_ID] & 0xffff) == 0x9a49;
+ igt_require_f(run_xe_compute_kernel(fd), "GPU not supported\n");
}
igt_main
@@ -136,10 +40,8 @@ igt_main
xe_device_get(xe);
}
- igt_subtest("compute-square") {
- igt_skip_on(!is_device_supported(xe));
+ igt_subtest("compute-square")
test_compute_square(xe);
- }
igt_fixture {
xe_device_put(xe);
--
2.39.2
^ permalink raw reply related [flat|nested] 8+ messages in thread
* [igt-dev] [PATCH i-g-t v2 3/4] lib/xe/xe_compute: use registers defs from intel_gpu_commands.h
2023-04-04 7:38 [igt-dev] [PATCH i-g-t v2 0/4] Make xe_compute test more generic Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 1/4] compute_square_kernel.cl: add CL file used at xe_compute.c Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 2/4] xe/xe_compute: place OpenCL kernel on a separate file Mauro Carvalho Chehab
@ 2023-04-04 7:38 ` Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 4/4] gen_opencl_kernel: add script to dynamically create OpenCL kernels Mauro Carvalho Chehab
` (2 subsequent siblings)
5 siblings, 0 replies; 8+ messages in thread
From: Mauro Carvalho Chehab @ 2023-04-04 7:38 UTC (permalink / raw)
To: igt-dev
From: Mauro Carvalho Chehab <mchehab@kernel.org>
There are some register definitions that are already defined inside
intel_gpu_commands.h with a different concept.
Change the code to re-use the definitions there.
Reviewed-by: Zbigniew Kempczyński <zbigniew.kempczynski@intel.com>
Signed-off-by: Mauro Carvalho Chehab <mchehab@kernel.org>
---
lib/xe/xe_compute.c | 17 ++++-------------
1 file changed, 4 insertions(+), 13 deletions(-)
diff --git a/lib/xe/xe_compute.c b/lib/xe/xe_compute.c
index fb11b8bc7770..2a3686a1bee6 100644
--- a/lib/xe/xe_compute.c
+++ b/lib/xe/xe_compute.c
@@ -33,15 +33,6 @@
#define OFFSET_INDIRECT_DATA_START 0xFFFDF000
#define OFFSET_KERNEL 0xFFFEF000
-#undef MEDIA_VFE_STATE
-#define MEDIA_VFE_STATE 0x70000007
-#undef STATE_BASE_ADDRESS
-#define STATE_BASE_ADDRESS 0x61010014
-#undef MEDIA_INTERFACE_DESCRIPTOR_LOAD
-#define MEDIA_INTERFACE_DESCRIPTOR_LOAD 0x70020002
-#undef GPGPU_WALKER
-#define GPGPU_WALKER 0x7105000d
-
struct bo_dict_entry {
uint64_t addr;
uint32_t size;
@@ -301,7 +292,7 @@ static void tgllp_compute_exec_compute(uint32_t *addr_bo_buffer_batch,
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00000000;
- addr_bo_buffer_batch[b++] = MEDIA_VFE_STATE;
+ addr_bo_buffer_batch[b++] = MEDIA_VFE_STATE | (9 - 2);
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00A70100;
@@ -316,7 +307,7 @@ static void tgllp_compute_exec_compute(uint32_t *addr_bo_buffer_batch,
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00000000;
- addr_bo_buffer_batch[b++] = STATE_BASE_ADDRESS;
+ addr_bo_buffer_batch[b++] = STATE_BASE_ADDRESS | (16 - 2);
addr_bo_buffer_batch[b++] = 0x00000001;
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00040000;
@@ -352,11 +343,11 @@ static void tgllp_compute_exec_compute(uint32_t *addr_bo_buffer_batch,
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = MEDIA_STATE_FLUSH;
addr_bo_buffer_batch[b++] = 0x00000000;
- addr_bo_buffer_batch[b++] = MEDIA_INTERFACE_DESCRIPTOR_LOAD;
+ addr_bo_buffer_batch[b++] = MEDIA_INTERFACE_DESCRIPTOR_LOAD | (4 - 2);
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00000020;
addr_bo_buffer_batch[b++] = 0x00000000;
- addr_bo_buffer_batch[b++] = GPGPU_WALKER;
+ addr_bo_buffer_batch[b++] = GPGPU_WALKER | 13;
addr_bo_buffer_batch[b++] = 0x00000000;
addr_bo_buffer_batch[b++] = 0x00000c80;
addr_bo_buffer_batch[b++] = offset_indirect_data_start;
--
2.39.2
^ permalink raw reply related [flat|nested] 8+ messages in thread
* [igt-dev] [PATCH i-g-t v2 4/4] gen_opencl_kernel: add script to dynamically create OpenCL kernels
2023-04-04 7:38 [igt-dev] [PATCH i-g-t v2 0/4] Make xe_compute test more generic Mauro Carvalho Chehab
` (2 preceding siblings ...)
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 3/4] lib/xe/xe_compute: use registers defs from intel_gpu_commands.h Mauro Carvalho Chehab
@ 2023-04-04 7:38 ` Mauro Carvalho Chehab
2023-04-04 14:28 ` Kamil Konieczny
2023-04-04 10:41 ` [igt-dev] ✓ Fi.CI.BAT: success for Make xe_compute test more generic (rev2) Patchwork
2023-04-04 16:53 ` [igt-dev] ✓ Fi.CI.IGT: " Patchwork
5 siblings, 1 reply; 8+ messages in thread
From: Mauro Carvalho Chehab @ 2023-04-04 7:38 UTC (permalink / raw)
To: igt-dev
From: Mauro Carvalho Chehab <mchehab@kernel.org>
Compute tests can be produced by using OpenCL, by calling ocloc.
While this can be part of IGT building system, for now, let's add
a script for such purpose.
Signed-off-by: Mauro Carvalho Chehab <mchehab@kernel.org>
---
opencl/README | 30 ++++++++++++
opencl/gen_opencl_kernel | 103 +++++++++++++++++++++++++++++++++++++++
2 files changed, 133 insertions(+)
create mode 100644 opencl/README
create mode 100755 opencl/gen_opencl_kernel
diff --git a/opencl/README b/opencl/README
new file mode 100644
index 000000000000..2fd0687a299b
--- /dev/null
+++ b/opencl/README
@@ -0,0 +1,30 @@
+This directory contains some OpenCL compute files, and a script to be used
+to produce a header file containing the binaries for the CL against
+multiple platforms.
+
+For instance, to generate compute square Kernel binaries for TGL and ADL
+variants, use this:
+
+ opencl/gen_opencl_kernel xe_compute_square opencl/compute_square_kernel.cl \
+ xe_compute_square_kernels.c build/opencl tgllp adl-s adl-p adl-n
+
+ cp build/opencl/xe_compute_square_kernels.c lib/xe/
+
+The opencl/gen_opencl_kernel requires the Intel compute runtime[1].
+
+[1] https://github.com/intel/compute-runtime
+
+This is usually shipped with different names on different distributions.
+The above command generate Kernels for both TGL and ADL platforms.
+Modern packages for Ubuntu are provided at github, under releases
+tag.
+
+Please notice that the GPU platforms supported by Intel ICD tools depend
+on its version. In order to know what's supported, you can run:
+
+ $ ocloc compile --help 2>&1|grep -A1 'Target device.'
+ -device <device_type> Target device.
+ <device_type> can be: bdw, skl, kbl, cfl, apl, bxt, glk, whl, aml, cml, icllp, lkf, ehl, jsl, tgllp, rkl, adl-s, adl-p, adl-n, dg1, acm-g10, ats-m150, dg2-g10, acm-g11, ats-m75, dg2-g11, acm-g12, dg2-g12, pvc-sdv, pvc, gen11, gen12lp, gen8, gen9, xe, xe-hp, xe-hpc, xe-hpg, version or hexadecimal value with 0x prefix
+
+The above results are for Intel ICD version 22.43.24558, which supports
+both TGL and ADL platforms, plus other newer GPU models.
diff --git a/opencl/gen_opencl_kernel b/opencl/gen_opencl_kernel
new file mode 100755
index 000000000000..e6f9601e0edb
--- /dev/null
+++ b/opencl/gen_opencl_kernel
@@ -0,0 +1,103 @@
+#!/bin/bash
+
+trap 'catch $LINENO' ERR
+
+catch() {
+ echo "error in line $1"
+ exit 1
+}
+
+
+# Parse arguments
+if [ $# -lt 5 ]; then
+ echo -e 'Usage:\n\t$0: <Kernel name> <kernel.cl> <header name> <dest_dir> <GPU models>' >&2
+ echo -e "Example:\n\t$0 kernel_foo kernel.cl kernels.c ../build/opencl tgllp rkl\n" >&2
+ exit 1
+fi
+
+if [ "$(xxd --help 2>&1|grep '\-n')" == "" ]; then
+ # Old versions have its own criteria to generate names.
+ # In this specific case, names will be like:
+ # "build_opencl_${GPU_DEVICE}_${kernel_name}_bin"
+ # Not fancy but it works.
+ USE_NAME_PARM=
+else
+ # Remove bloatware from the names, calling Kernels as:
+ # "${GPU_DEVICE}_${kernel_name}"
+ USE_NAME_PARM=1
+fi
+
+
+kernel_name=$1
+shift
+
+kernel_cl=$1
+shift
+
+output_fname=$1
+shift
+
+dest_dir=$1
+shift
+
+mkdir -p $dest_dir
+
+args=( "$@" )
+
+echo $args
+
+out_files=""
+for i in "${args[@]}"; do
+ name="$dest_dir/${i}_${kernel_name}"
+ out="$name.h"
+ echo "Generating $out"
+ ocloc compile -q -file ${kernel_cl} -device ${i} -output ${name}_bin -output_no_suffix
+ if [ "$USE_NAME_PARM" != "" ]; then
+ xxd -n "${i}_${kernel_name}" -i ${name}_bin >$out
+ else
+ xxd -i ${name}_bin >$out
+ fi
+ sed "s, ,\t,;s,.*unsigned int.*,,;s,\-,_,g;s,unsigned,static const unsigned," -i $out
+ sed "1 i// Match ID: $(ocloc ids $i|grep -v "Matched ids:")" -i $out
+ out_files+=" $out"
+done
+
+output_fname="$dest_dir/$output_fname"
+echo "Generating $output_fname"
+
+cat << PREFIX >$output_fname
+/* SPDX-License-Identifier: MIT */
+/*
+ * This file is auto-generated from $kernel_cl:
+ *
+PREFIX
+
+cat $kernel_cl |sed s,"^"," * ," >>$output_fname
+
+cat << INCLUDES >>$output_fname
+ */
+
+#include "intel_chipset.h"
+#include "lib/xe/xe_compute.h"
+
+INCLUDES
+
+cat $out_files >>$output_fname
+
+echo "const struct xe_compute_kernels ${kernel_name}_kernels[] = {" >>$output_fname
+
+for i in "${args[@]}"; do
+ out="$dest_dir/${i}_${kernel_name}.h"
+ echo -e "\t{" >>$output_fname; \
+ grep "Match ID:" $out|sed -E "s/.*\s([0-9]+)\.([0-9]+).*/\t\t.ip_ver = IP_VER(\1, \2),/" >>$output_fname;
+ grep unsigned $out|sed -E "s/.*\s+([_a-zA-Z0-9]+)\[\].*/\t\t.size = sizeof(\1),/" >>$output_fname;
+ grep unsigned $out|sed -E "s/.*\s+([_a-zA-Z0-9]+)\[\].*/\t\t.kernel = \1,/" >>$output_fname;
+ echo -e "\t}," >>$output_fname;
+done
+
+cat << SUFFIX >>$output_fname
+ {}
+};
+SUFFIX
+
+echo "Done."
--
2.39.2
^ permalink raw reply related [flat|nested] 8+ messages in thread
* [igt-dev] ✓ Fi.CI.BAT: success for Make xe_compute test more generic (rev2)
2023-04-04 7:38 [igt-dev] [PATCH i-g-t v2 0/4] Make xe_compute test more generic Mauro Carvalho Chehab
` (3 preceding siblings ...)
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 4/4] gen_opencl_kernel: add script to dynamically create OpenCL kernels Mauro Carvalho Chehab
@ 2023-04-04 10:41 ` Patchwork
2023-04-04 16:53 ` [igt-dev] ✓ Fi.CI.IGT: " Patchwork
5 siblings, 0 replies; 8+ messages in thread
From: Patchwork @ 2023-04-04 10:41 UTC (permalink / raw)
To: Mauro Carvalho Chehab; +Cc: igt-dev
[-- Attachment #1: Type: text/plain, Size: 4207 bytes --]
== Series Details ==
Series: Make xe_compute test more generic (rev2)
URL : https://patchwork.freedesktop.org/series/115670/
State : success
== Summary ==
CI Bug Log - changes from IGT_7234 -> IGTPW_8748
====================================================
Summary
-------
**SUCCESS**
No regressions found.
External URL: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/index.html
Participating hosts (37 -> 36)
------------------------------
Missing (1): fi-snb-2520m
Known issues
------------
Here are the changes found in IGTPW_8748 that come from known issues:
### IGT changes ###
#### Issues hit ####
* igt@kms_chamelium_hpd@common-hpd-after-suspend:
- bat-rpls-2: NOTRUN -> [SKIP][1] ([i915#7828])
[1]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/bat-rpls-2/igt@kms_chamelium_hpd@common-hpd-after-suspend.html
- bat-adln-1: NOTRUN -> [SKIP][2] ([i915#7828])
[2]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/bat-adln-1/igt@kms_chamelium_hpd@common-hpd-after-suspend.html
* igt@kms_pipe_crc_basic@suspend-read-crc:
- bat-rpls-2: NOTRUN -> [SKIP][3] ([i915#1845])
[3]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/bat-rpls-2/igt@kms_pipe_crc_basic@suspend-read-crc.html
#### Possible fixes ####
* igt@dmabuf@all-tests@dma_fence:
- bat-adln-1: [DMESG-FAIL][4] ([i915#8143]) -> [PASS][5]
[4]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/bat-adln-1/igt@dmabuf@all-tests@dma_fence.html
[5]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/bat-adln-1/igt@dmabuf@all-tests@dma_fence.html
* igt@dmabuf@all-tests@sanitycheck:
- bat-adln-1: [ABORT][6] ([i915#8058] / [i915#8144]) -> [PASS][7]
[6]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/bat-adln-1/igt@dmabuf@all-tests@sanitycheck.html
[7]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/bat-adln-1/igt@dmabuf@all-tests@sanitycheck.html
* igt@gem_exec_suspend@basic-s0@smem:
- bat-rpls-2: [ABORT][8] ([i915#6687]) -> [PASS][9]
[8]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/bat-rpls-2/igt@gem_exec_suspend@basic-s0@smem.html
[9]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/bat-rpls-2/igt@gem_exec_suspend@basic-s0@smem.html
* igt@kms_pipe_crc_basic@nonblocking-crc-frame-sequence@pipe-c-dp-1:
- bat-dg2-8: [FAIL][10] ([i915#7932]) -> [PASS][11] +1 similar issue
[10]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/bat-dg2-8/igt@kms_pipe_crc_basic@nonblocking-crc-frame-sequence@pipe-c-dp-1.html
[11]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/bat-dg2-8/igt@kms_pipe_crc_basic@nonblocking-crc-frame-sequence@pipe-c-dp-1.html
#### Warnings ####
* igt@i915_selftest@live@slpc:
- bat-rpls-1: [DMESG-FAIL][12] ([i915#6997]) -> [DMESG-FAIL][13] ([i915#6367] / [i915#6997])
[12]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/bat-rpls-1/igt@i915_selftest@live@slpc.html
[13]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/bat-rpls-1/igt@i915_selftest@live@slpc.html
[i915#1845]: https://gitlab.freedesktop.org/drm/intel/issues/1845
[i915#6367]: https://gitlab.freedesktop.org/drm/intel/issues/6367
[i915#6687]: https://gitlab.freedesktop.org/drm/intel/issues/6687
[i915#6997]: https://gitlab.freedesktop.org/drm/intel/issues/6997
[i915#7828]: https://gitlab.freedesktop.org/drm/intel/issues/7828
[i915#7932]: https://gitlab.freedesktop.org/drm/intel/issues/7932
[i915#8058]: https://gitlab.freedesktop.org/drm/intel/issues/8058
[i915#8143]: https://gitlab.freedesktop.org/drm/intel/issues/8143
[i915#8144]: https://gitlab.freedesktop.org/drm/intel/issues/8144
Build changes
-------------
* CI: CI-20190529 -> None
* IGT: IGT_7234 -> IGTPW_8748
CI-20190529: 20190529
CI_DRM_12963: 9ce564e0e7b78a777aade183f3deed30b5b83336 @ git://anongit.freedesktop.org/gfx-ci/linux
IGTPW_8748: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/index.html
IGT_7234: 70802fb59c65164f3587e71376ebed1ed5e91fd1 @ https://gitlab.freedesktop.org/drm/igt-gpu-tools.git
== Logs ==
For more details see: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/index.html
[-- Attachment #2: Type: text/html, Size: 5118 bytes --]
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [igt-dev] [PATCH i-g-t v2 4/4] gen_opencl_kernel: add script to dynamically create OpenCL kernels
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 4/4] gen_opencl_kernel: add script to dynamically create OpenCL kernels Mauro Carvalho Chehab
@ 2023-04-04 14:28 ` Kamil Konieczny
0 siblings, 0 replies; 8+ messages in thread
From: Kamil Konieczny @ 2023-04-04 14:28 UTC (permalink / raw)
To: igt-dev
Hi Mauro,
On 2023-04-04 at 09:38:35 +0200, Mauro Carvalho Chehab wrote:
> From: Mauro Carvalho Chehab <mchehab@kernel.org>
>
> Compute tests can be produced by using OpenCL, by calling ocloc.
>
> While this can be part of IGT building system, for now, let's add
> a script for such purpose.
>
> Signed-off-by: Mauro Carvalho Chehab <mchehab@kernel.org>
It works on my Ubuntu 22.10,
Acked-by: Kamil Konieczny <kamil.konieczny@linux.intel.com>
Regards,
Kamil
> ---
> opencl/README | 30 ++++++++++++
> opencl/gen_opencl_kernel | 103 +++++++++++++++++++++++++++++++++++++++
> 2 files changed, 133 insertions(+)
> create mode 100644 opencl/README
> create mode 100755 opencl/gen_opencl_kernel
>
> diff --git a/opencl/README b/opencl/README
> new file mode 100644
> index 000000000000..2fd0687a299b
> --- /dev/null
> +++ b/opencl/README
> @@ -0,0 +1,30 @@
> +This directory contains some OpenCL compute files, and a script to be used
> +to produce a header file containing the binaries for the CL against
> +multiple platforms.
> +
> +For instance, to generate compute square Kernel binaries for TGL and ADL
> +variants, use this:
> +
> + opencl/gen_opencl_kernel xe_compute_square opencl/compute_square_kernel.cl \
> + xe_compute_square_kernels.c build/opencl tgllp adl-s adl-p adl-n
> +
> + cp build/opencl/xe_compute_square_kernels.c lib/xe/
> +
> +The opencl/gen_opencl_kernel requires the Intel compute runtime[1].
> +
> +[1] https://github.com/intel/compute-runtime
> +
> +This is usually shipped with different names on different distributions.
> +The above command generate Kernels for both TGL and ADL platforms.
> +Modern packages for Ubuntu are provided at github, under releases
> +tag.
> +
> +Please notice that the GPU platforms supported by Intel ICD tools depend
> +on its version. In order to know what's supported, you can run:
> +
> + $ ocloc compile --help 2>&1|grep -A1 'Target device.'
> + -device <device_type> Target device.
> + <device_type> can be: bdw, skl, kbl, cfl, apl, bxt, glk, whl, aml, cml, icllp, lkf, ehl, jsl, tgllp, rkl, adl-s, adl-p, adl-n, dg1, acm-g10, ats-m150, dg2-g10, acm-g11, ats-m75, dg2-g11, acm-g12, dg2-g12, pvc-sdv, pvc, gen11, gen12lp, gen8, gen9, xe, xe-hp, xe-hpc, xe-hpg, version or hexadecimal value with 0x prefix
> +
> +The above results are for Intel ICD version 22.43.24558, which supports
> +both TGL and ADL platforms, plus other newer GPU models.
> diff --git a/opencl/gen_opencl_kernel b/opencl/gen_opencl_kernel
> new file mode 100755
> index 000000000000..e6f9601e0edb
> --- /dev/null
> +++ b/opencl/gen_opencl_kernel
> @@ -0,0 +1,103 @@
> +#!/bin/bash
> +
> +trap 'catch $LINENO' ERR
> +
> +catch() {
> + echo "error in line $1"
> + exit 1
> +}
> +
> +
> +# Parse arguments
> +if [ $# -lt 5 ]; then
> + echo -e 'Usage:\n\t$0: <Kernel name> <kernel.cl> <header name> <dest_dir> <GPU models>' >&2
> + echo -e "Example:\n\t$0 kernel_foo kernel.cl kernels.c ../build/opencl tgllp rkl\n" >&2
> + exit 1
> +fi
> +
> +if [ "$(xxd --help 2>&1|grep '\-n')" == "" ]; then
> + # Old versions have its own criteria to generate names.
> + # In this specific case, names will be like:
> + # "build_opencl_${GPU_DEVICE}_${kernel_name}_bin"
> + # Not fancy but it works.
> + USE_NAME_PARM=
> +else
> + # Remove bloatware from the names, calling Kernels as:
> + # "${GPU_DEVICE}_${kernel_name}"
> + USE_NAME_PARM=1
> +fi
> +
> +
> +kernel_name=$1
> +shift
> +
> +kernel_cl=$1
> +shift
> +
> +output_fname=$1
> +shift
> +
> +dest_dir=$1
> +shift
> +
> +mkdir -p $dest_dir
> +
> +args=( "$@" )
> +
> +echo $args
> +
> +out_files=""
> +for i in "${args[@]}"; do
> + name="$dest_dir/${i}_${kernel_name}"
> + out="$name.h"
> + echo "Generating $out"
> + ocloc compile -q -file ${kernel_cl} -device ${i} -output ${name}_bin -output_no_suffix
> + if [ "$USE_NAME_PARM" != "" ]; then
> + xxd -n "${i}_${kernel_name}" -i ${name}_bin >$out
> + else
> + xxd -i ${name}_bin >$out
> + fi
> + sed "s, ,\t,;s,.*unsigned int.*,,;s,\-,_,g;s,unsigned,static const unsigned," -i $out
> + sed "1 i// Match ID: $(ocloc ids $i|grep -v "Matched ids:")" -i $out
> + out_files+=" $out"
> +done
> +
> +output_fname="$dest_dir/$output_fname"
> +echo "Generating $output_fname"
> +
> +cat << PREFIX >$output_fname
> +/* SPDX-License-Identifier: MIT */
> +/*
> + * This file is auto-generated from $kernel_cl:
> + *
> +PREFIX
> +
> +cat $kernel_cl |sed s,"^"," * ," >>$output_fname
> +
> +cat << INCLUDES >>$output_fname
> + */
> +
> +#include "intel_chipset.h"
> +#include "lib/xe/xe_compute.h"
> +
> +INCLUDES
> +
> +cat $out_files >>$output_fname
> +
> +echo "const struct xe_compute_kernels ${kernel_name}_kernels[] = {" >>$output_fname
> +
> +for i in "${args[@]}"; do
> + out="$dest_dir/${i}_${kernel_name}.h"
> + echo -e "\t{" >>$output_fname; \
> + grep "Match ID:" $out|sed -E "s/.*\s([0-9]+)\.([0-9]+).*/\t\t.ip_ver = IP_VER(\1, \2),/" >>$output_fname;
> + grep unsigned $out|sed -E "s/.*\s+([_a-zA-Z0-9]+)\[\].*/\t\t.size = sizeof(\1),/" >>$output_fname;
> + grep unsigned $out|sed -E "s/.*\s+([_a-zA-Z0-9]+)\[\].*/\t\t.kernel = \1,/" >>$output_fname;
> + echo -e "\t}," >>$output_fname;
> +done
> +
> +cat << SUFFIX >>$output_fname
> + {}
> +};
> +SUFFIX
> +
> +echo "Done."
> --
> 2.39.2
>
^ permalink raw reply [flat|nested] 8+ messages in thread
* [igt-dev] ✓ Fi.CI.IGT: success for Make xe_compute test more generic (rev2)
2023-04-04 7:38 [igt-dev] [PATCH i-g-t v2 0/4] Make xe_compute test more generic Mauro Carvalho Chehab
` (4 preceding siblings ...)
2023-04-04 10:41 ` [igt-dev] ✓ Fi.CI.BAT: success for Make xe_compute test more generic (rev2) Patchwork
@ 2023-04-04 16:53 ` Patchwork
5 siblings, 0 replies; 8+ messages in thread
From: Patchwork @ 2023-04-04 16:53 UTC (permalink / raw)
To: Mauro Carvalho Chehab; +Cc: igt-dev
[-- Attachment #1: Type: text/plain, Size: 15296 bytes --]
== Series Details ==
Series: Make xe_compute test more generic (rev2)
URL : https://patchwork.freedesktop.org/series/115670/
State : success
== Summary ==
CI Bug Log - changes from IGT_7234_full -> IGTPW_8748_full
====================================================
Summary
-------
**SUCCESS**
No regressions found.
External URL: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/index.html
Participating hosts (7 -> 7)
------------------------------
No changes in participating hosts
Possible new issues
-------------------
Here are the unknown changes that may have been introduced in IGTPW_8748_full:
### IGT changes ###
#### Suppressed ####
The following results come from untrusted machines, tests, or statuses.
They do not affect the overall result.
* igt@kms_rotation_crc@primary-y-tiled-reflect-x-0:
- {shard-rkl}: [PASS][1] -> [ABORT][2]
[1]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-rkl-6/igt@kms_rotation_crc@primary-y-tiled-reflect-x-0.html
[2]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-rkl-6/igt@kms_rotation_crc@primary-y-tiled-reflect-x-0.html
Known issues
------------
Here are the changes found in IGTPW_8748_full that come from known issues:
### IGT changes ###
#### Issues hit ####
* igt@gem_exec_fair@basic-pace-share@rcs0:
- shard-apl: [PASS][3] -> [FAIL][4] ([i915#2842])
[3]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-apl2/igt@gem_exec_fair@basic-pace-share@rcs0.html
[4]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-apl4/igt@gem_exec_fair@basic-pace-share@rcs0.html
* igt@gem_lmem_swapping@heavy-verify-random-ccs:
- shard-apl: NOTRUN -> [SKIP][5] ([fdo#109271] / [i915#4613])
[5]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-apl4/igt@gem_lmem_swapping@heavy-verify-random-ccs.html
* igt@gen9_exec_parse@allowed-all:
- shard-apl: [PASS][6] -> [ABORT][7] ([i915#5566])
[6]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-apl6/igt@gen9_exec_parse@allowed-all.html
[7]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-apl2/igt@gen9_exec_parse@allowed-all.html
* igt@i915_module_load@reload-with-fault-injection:
- shard-snb: [PASS][8] -> [ABORT][9] ([i915#4528])
[8]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-snb2/igt@i915_module_load@reload-with-fault-injection.html
[9]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-snb4/igt@i915_module_load@reload-with-fault-injection.html
* igt@kms_ccs@pipe-b-crc-sprite-planes-basic-y_tiled_gen12_mc_ccs:
- shard-apl: NOTRUN -> [SKIP][10] ([fdo#109271] / [i915#3886]) +3 similar issues
[10]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-apl2/igt@kms_ccs@pipe-b-crc-sprite-planes-basic-y_tiled_gen12_mc_ccs.html
* igt@kms_ccs@pipe-c-ccs-on-another-bo-y_tiled_ccs:
- shard-apl: NOTRUN -> [SKIP][11] ([fdo#109271]) +46 similar issues
[11]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-apl3/igt@kms_ccs@pipe-c-ccs-on-another-bo-y_tiled_ccs.html
* igt@kms_content_protection@atomic@pipe-a-dp-1:
- shard-apl: NOTRUN -> [TIMEOUT][12] ([i915#7173])
[12]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-apl6/igt@kms_content_protection@atomic@pipe-a-dp-1.html
* igt@kms_cursor_legacy@flip-vs-cursor-atomic-transitions-varying-size:
- shard-apl: [PASS][13] -> [FAIL][14] ([i915#2346])
[13]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-apl6/igt@kms_cursor_legacy@flip-vs-cursor-atomic-transitions-varying-size.html
[14]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-apl2/igt@kms_cursor_legacy@flip-vs-cursor-atomic-transitions-varying-size.html
#### Possible fixes ####
* igt@drm_fdinfo@most-busy-idle-check-all@rcs0:
- {shard-rkl}: [FAIL][15] ([i915#7742]) -> [PASS][16]
[15]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-rkl-6/igt@drm_fdinfo@most-busy-idle-check-all@rcs0.html
[16]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-rkl-1/igt@drm_fdinfo@most-busy-idle-check-all@rcs0.html
* igt@gem_exec_fair@basic-pace-share@rcs0:
- shard-glk: [FAIL][17] ([i915#2842]) -> [PASS][18] +2 similar issues
[17]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-glk7/igt@gem_exec_fair@basic-pace-share@rcs0.html
[18]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-glk9/igt@gem_exec_fair@basic-pace-share@rcs0.html
* igt@gem_exec_suspend@basic-s4-devices@smem:
- {shard-tglu}: [ABORT][19] ([i915#7975]) -> [PASS][20]
[19]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-tglu-10/igt@gem_exec_suspend@basic-s4-devices@smem.html
[20]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-tglu-2/igt@gem_exec_suspend@basic-s4-devices@smem.html
* igt@i915_pm_dc@dc9-dpms:
- shard-apl: [SKIP][21] ([fdo#109271]) -> [PASS][22]
[21]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-apl3/igt@i915_pm_dc@dc9-dpms.html
[22]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-apl1/igt@i915_pm_dc@dc9-dpms.html
* igt@i915_pm_rpm@modeset-lpsp-stress-no-wait:
- {shard-rkl}: [SKIP][23] ([i915#1397]) -> [PASS][24]
[23]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-rkl-6/igt@i915_pm_rpm@modeset-lpsp-stress-no-wait.html
[24]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-rkl-7/igt@i915_pm_rpm@modeset-lpsp-stress-no-wait.html
* igt@kms_cursor_legacy@single-move@pipe-b:
- {shard-rkl}: [INCOMPLETE][25] ([i915#8011]) -> [PASS][26] +1 similar issue
[25]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-rkl-7/igt@kms_cursor_legacy@single-move@pipe-b.html
[26]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-rkl-2/igt@kms_cursor_legacy@single-move@pipe-b.html
* igt@kms_flip@2x-plain-flip-fb-recreate-interruptible@bc-hdmi-a1-hdmi-a2:
- shard-glk: [FAIL][27] ([i915#2122]) -> [PASS][28]
[27]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-glk5/igt@kms_flip@2x-plain-flip-fb-recreate-interruptible@bc-hdmi-a1-hdmi-a2.html
[28]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-glk3/igt@kms_flip@2x-plain-flip-fb-recreate-interruptible@bc-hdmi-a1-hdmi-a2.html
* igt@kms_plane_scaling@i915-max-src-size@pipe-a-hdmi-a-2:
- {shard-rkl}: [FAIL][29] -> [PASS][30]
[29]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-rkl-3/igt@kms_plane_scaling@i915-max-src-size@pipe-a-hdmi-a-2.html
[30]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-rkl-2/igt@kms_plane_scaling@i915-max-src-size@pipe-a-hdmi-a-2.html
#### Warnings ####
* igt@kms_content_protection@atomic-dpms@pipe-a-dp-1:
- shard-apl: [TIMEOUT][31] ([i915#7173]) -> [FAIL][32] ([i915#7173])
[31]: https://intel-gfx-ci.01.org/tree/drm-tip/IGT_7234/shard-apl4/igt@kms_content_protection@atomic-dpms@pipe-a-dp-1.html
[32]: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/shard-apl2/igt@kms_content_protection@atomic-dpms@pipe-a-dp-1.html
{name}: This element is suppressed. This means it is ignored when computing
the status of the difference (SUCCESS, WARNING, or FAILURE).
[IGT#2]: https://gitlab.freedesktop.org/drm/igt-gpu-tools/issues/2
[fdo#109271]: https://bugs.freedesktop.org/show_bug.cgi?id=109271
[fdo#109274]: https://bugs.freedesktop.org/show_bug.cgi?id=109274
[fdo#109280]: https://bugs.freedesktop.org/show_bug.cgi?id=109280
[fdo#109289]: https://bugs.freedesktop.org/show_bug.cgi?id=109289
[fdo#109302]: https://bugs.freedesktop.org/show_bug.cgi?id=109302
[fdo#109303]: https://bugs.freedesktop.org/show_bug.cgi?id=109303
[fdo#109315]: https://bugs.freedesktop.org/show_bug.cgi?id=109315
[fdo#110189]: https://bugs.freedesktop.org/show_bug.cgi?id=110189
[fdo#110723]: https://bugs.freedesktop.org/show_bug.cgi?id=110723
[fdo#111068]: https://bugs.freedesktop.org/show_bug.cgi?id=111068
[fdo#111615]: https://bugs.freedesktop.org/show_bug.cgi?id=111615
[fdo#111825]: https://bugs.freedesktop.org/show_bug.cgi?id=111825
[fdo#111827]: https://bugs.freedesktop.org/show_bug.cgi?id=111827
[fdo#112054]: https://bugs.freedesktop.org/show_bug.cgi?id=112054
[fdo#112283]: https://bugs.freedesktop.org/show_bug.cgi?id=112283
[i915#1072]: https://gitlab.freedesktop.org/drm/intel/issues/1072
[i915#1257]: https://gitlab.freedesktop.org/drm/intel/issues/1257
[i915#1397]: https://gitlab.freedesktop.org/drm/intel/issues/1397
[i915#1825]: https://gitlab.freedesktop.org/drm/intel/issues/1825
[i915#1839]: https://gitlab.freedesktop.org/drm/intel/issues/1839
[i915#1937]: https://gitlab.freedesktop.org/drm/intel/issues/1937
[i915#2122]: https://gitlab.freedesktop.org/drm/intel/issues/2122
[i915#2346]: https://gitlab.freedesktop.org/drm/intel/issues/2346
[i915#2437]: https://gitlab.freedesktop.org/drm/intel/issues/2437
[i915#2527]: https://gitlab.freedesktop.org/drm/intel/issues/2527
[i915#2575]: https://gitlab.freedesktop.org/drm/intel/issues/2575
[i915#2587]: https://gitlab.freedesktop.org/drm/intel/issues/2587
[i915#2672]: https://gitlab.freedesktop.org/drm/intel/issues/2672
[i915#2681]: https://gitlab.freedesktop.org/drm/intel/issues/2681
[i915#280]: https://gitlab.freedesktop.org/drm/intel/issues/280
[i915#284]: https://gitlab.freedesktop.org/drm/intel/issues/284
[i915#2842]: https://gitlab.freedesktop.org/drm/intel/issues/2842
[i915#2856]: https://gitlab.freedesktop.org/drm/intel/issues/2856
[i915#3023]: https://gitlab.freedesktop.org/drm/intel/issues/3023
[i915#315]: https://gitlab.freedesktop.org/drm/intel/issues/315
[i915#3281]: https://gitlab.freedesktop.org/drm/intel/issues/3281
[i915#3282]: https://gitlab.freedesktop.org/drm/intel/issues/3282
[i915#3297]: https://gitlab.freedesktop.org/drm/intel/issues/3297
[i915#3299]: https://gitlab.freedesktop.org/drm/intel/issues/3299
[i915#3359]: https://gitlab.freedesktop.org/drm/intel/issues/3359
[i915#3458]: https://gitlab.freedesktop.org/drm/intel/issues/3458
[i915#3469]: https://gitlab.freedesktop.org/drm/intel/issues/3469
[i915#3539]: https://gitlab.freedesktop.org/drm/intel/issues/3539
[i915#3555]: https://gitlab.freedesktop.org/drm/intel/issues/3555
[i915#3591]: https://gitlab.freedesktop.org/drm/intel/issues/3591
[i915#3638]: https://gitlab.freedesktop.org/drm/intel/issues/3638
[i915#3689]: https://gitlab.freedesktop.org/drm/intel/issues/3689
[i915#3708]: https://gitlab.freedesktop.org/drm/intel/issues/3708
[i915#3734]: https://gitlab.freedesktop.org/drm/intel/issues/3734
[i915#3742]: https://gitlab.freedesktop.org/drm/intel/issues/3742
[i915#3743]: https://gitlab.freedesktop.org/drm/intel/issues/3743
[i915#3840]: https://gitlab.freedesktop.org/drm/intel/issues/3840
[i915#3886]: https://gitlab.freedesktop.org/drm/intel/issues/3886
[i915#3936]: https://gitlab.freedesktop.org/drm/intel/issues/3936
[i915#3938]: https://gitlab.freedesktop.org/drm/intel/issues/3938
[i915#4036]: https://gitlab.freedesktop.org/drm/intel/issues/4036
[i915#4077]: https://gitlab.freedesktop.org/drm/intel/issues/4077
[i915#4078]: https://gitlab.freedesktop.org/drm/intel/issues/4078
[i915#4079]: https://gitlab.freedesktop.org/drm/intel/issues/4079
[i915#4083]: https://gitlab.freedesktop.org/drm/intel/issues/4083
[i915#4103]: https://gitlab.freedesktop.org/drm/intel/issues/4103
[i915#4212]: https://gitlab.freedesktop.org/drm/intel/issues/4212
[i915#4213]: https://gitlab.freedesktop.org/drm/intel/issues/4213
[i915#4270]: https://gitlab.freedesktop.org/drm/intel/issues/4270
[i915#4349]: https://gitlab.freedesktop.org/drm/intel/issues/4349
[i915#4387]: https://gitlab.freedesktop.org/drm/intel/issues/4387
[i915#4528]: https://gitlab.freedesktop.org/drm/intel/issues/4528
[i915#4538]: https://gitlab.freedesktop.org/drm/intel/issues/4538
[i915#4579]: https://gitlab.freedesktop.org/drm/intel/issues/4579
[i915#4613]: https://gitlab.freedesktop.org/drm/intel/issues/4613
[i915#4771]: https://gitlab.freedesktop.org/drm/intel/issues/4771
[i915#4812]: https://gitlab.freedesktop.org/drm/intel/issues/4812
[i915#4833]: https://gitlab.freedesktop.org/drm/intel/issues/4833
[i915#4852]: https://gitlab.freedesktop.org/drm/intel/issues/4852
[i915#4854]: https://gitlab.freedesktop.org/drm/intel/issues/4854
[i915#4859]: https://gitlab.freedesktop.org/drm/intel/issues/4859
[i915#4860]: https://gitlab.freedesktop.org/drm/intel/issues/4860
[i915#4873]: https://gitlab.freedesktop.org/drm/intel/issues/4873
[i915#4884]: https://gitlab.freedesktop.org/drm/intel/issues/4884
[i915#5099]: https://gitlab.freedesktop.org/drm/intel/issues/5099
[i915#5176]: https://gitlab.freedesktop.org/drm/intel/issues/5176
[i915#5235]: https://gitlab.freedesktop.org/drm/intel/issues/5235
[i915#5286]: https://gitlab.freedesktop.org/drm/intel/issues/5286
[i915#5288]: https://gitlab.freedesktop.org/drm/intel/issues/5288
[i915#5289]: https://gitlab.freedesktop.org/drm/intel/issues/5289
[i915#5354]: https://gitlab.freedesktop.org/drm/intel/issues/5354
[i915#5439]: https://gitlab.freedesktop.org/drm/intel/issues/5439
[i915#5461]: https://gitlab.freedesktop.org/drm/intel/issues/5461
[i915#5563]: https://gitlab.freedesktop.org/drm/intel/issues/5563
[i915#5566]: https://gitlab.freedesktop.org/drm/intel/issues/5566
[i915#5784]: https://gitlab.freedesktop.org/drm/intel/issues/5784
[i915#6095]: https://gitlab.freedesktop.org/drm/intel/issues/6095
[i915#6230]: https://gitlab.freedesktop.org/drm/intel/issues/6230
[i915#6334]: https://gitlab.freedesktop.org/drm/intel/issues/6334
[i915#6493]: https://gitlab.freedesktop.org/drm/intel/issues/6493
[i915#658]: https://gitlab.freedesktop.org/drm/intel/issues/658
[i915#7116]: https://gitlab.freedesktop.org/drm/intel/issues/7116
[i915#7173]: https://gitlab.freedesktop.org/drm/intel/issues/7173
[i915#7561]: https://gitlab.freedesktop.org/drm/intel/issues/7561
[i915#7697]: https://gitlab.freedesktop.org/drm/intel/issues/7697
[i915#7701]: https://gitlab.freedesktop.org/drm/intel/issues/7701
[i915#7711]: https://gitlab.freedesktop.org/drm/intel/issues/7711
[i915#7742]: https://gitlab.freedesktop.org/drm/intel/issues/7742
[i915#7828]: https://gitlab.freedesktop.org/drm/intel/issues/7828
[i915#7975]: https://gitlab.freedesktop.org/drm/intel/issues/7975
[i915#8011]: https://gitlab.freedesktop.org/drm/intel/issues/8011
[i915#8150]: https://gitlab.freedesktop.org/drm/intel/issues/8150
[i915#8155]: https://gitlab.freedesktop.org/drm/intel/issues/8155
[i915#8308]: https://gitlab.freedesktop.org/drm/intel/issues/8308
Build changes
-------------
* CI: CI-20190529 -> None
* IGT: IGT_7234 -> IGTPW_8748
CI-20190529: 20190529
CI_DRM_12963: 9ce564e0e7b78a777aade183f3deed30b5b83336 @ git://anongit.freedesktop.org/gfx-ci/linux
IGTPW_8748: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/index.html
IGT_7234: 70802fb59c65164f3587e71376ebed1ed5e91fd1 @ https://gitlab.freedesktop.org/drm/igt-gpu-tools.git
== Logs ==
For more details see: https://intel-gfx-ci.01.org/tree/drm-tip/IGTPW_8748/index.html
[-- Attachment #2: Type: text/html, Size: 10158 bytes --]
^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2023-04-04 16:53 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz follow: Atom feed
-- links below jump to the message on this page --
2023-04-04 7:38 [igt-dev] [PATCH i-g-t v2 0/4] Make xe_compute test more generic Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 1/4] compute_square_kernel.cl: add CL file used at xe_compute.c Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 2/4] xe/xe_compute: place OpenCL kernel on a separate file Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 3/4] lib/xe/xe_compute: use registers defs from intel_gpu_commands.h Mauro Carvalho Chehab
2023-04-04 7:38 ` [igt-dev] [PATCH i-g-t v2 4/4] gen_opencl_kernel: add script to dynamically create OpenCL kernels Mauro Carvalho Chehab
2023-04-04 14:28 ` Kamil Konieczny
2023-04-04 10:41 ` [igt-dev] ✓ Fi.CI.BAT: success for Make xe_compute test more generic (rev2) Patchwork
2023-04-04 16:53 ` [igt-dev] ✓ Fi.CI.IGT: " Patchwork
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox