* [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
* 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.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
* [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