Description
The Intercept Layer incorrectly deduces the type of 1D images and 1D image arrays. When an application creates a CL_MEM_OBJECT_IMAGE1D or CL_MEM_OBJECT_IMAGE1D_ARRAY, the internal metadata stored in m_ImageInfoMap categorizes them as IMAGE2D and IMAGE2D_ARRAY respectively. This leads to corrupted metadata when generating replay captures (via CLI_CaptureReplay=1), causing subsequent failures when parsing or replaying the serialized objects.
Root Cause
In intercept/src/clintercept_core.cpp, inside the CLIntercept::addImage(cl_mem image) function, the image type is deduced based on its dimensions (specifically checking for height == 0 and depth == 0). For a valid 1D image, both height == 0 and depth == 0 are true. The code correctly sets the ImageType to 1D in the first block, but immediately enters the second block and overwrites it.
Steps to Repraoduce
- Compile the following minimal reproducer that creates only 1D and 1D Array images:
#include <CL/cl.h>
#include <iostream>
int main() {
cl_platform_id platform;
cl_device_id device;
cl_uint num_platforms, num_devices;
clGetPlatformIDs(1, &platform, &num_platforms);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, &num_devices);
cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
cl_command_queue queue = clCreateCommandQueue(context, device, 0, NULL);
cl_image_format format = { CL_RGBA, CL_FLOAT };
cl_int err;
// 1. IMAGE1D
cl_image_desc desc1d = {};
desc1d.image_type = CL_MEM_OBJECT_IMAGE1D;
desc1d.image_width = 1024;
// 2. IMAGE1D_ARRAY
cl_image_desc desc1da = {};
desc1da.image_type = CL_MEM_OBJECT_IMAGE1D_ARRAY;
desc1da.image_width = 1024;
desc1da.image_array_size = 5;
cl_mem img1d = clCreateImage(context, CL_MEM_READ_WRITE, &format, &desc1d, NULL, &err);
cl_mem img1da = clCreateImage(context, CL_MEM_READ_WRITE, &format, &desc1da, NULL, &err);
const char* source =
"__kernel void test_1d(__write_only image1d_t i1, __write_only image1d_array_t i1a) {}";
cl_program program = clCreateProgramWithSource(context, 1, &source, NULL, NULL);
clBuildProgram(program, 1, &device, NULL, NULL, NULL);
cl_kernel kernel = clCreateKernel(program, "test_1d", NULL);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &img1d);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &img1da);
size_t global_work_size[1] = { 1024 };
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);
clFinish(queue);
std::cout << "[+] Kernel executed!" << std::endl;
clReleaseKernel(kernel); clReleaseProgram(program);
clReleaseMemObject(img1d); clReleaseMemObject(img1da);
clReleaseCommandQueue(queue); clReleaseContext(context);
return 0;
}
- Run the executable with Replay Capture enabled:
set CLI_CaptureReplay=1
cliloader.exe ./test_app.exe
- Inspect the exported metadata files for the kernel arguments in the replay directory.
The metadata exported for the IMAGE1D argument shows it has been incorrectly overwritten to CL_MEM_OBJECT_IMAGE2D (0x10F1 / 4337). Output extracted from the IMAGE1D metadata:
1024 // Width
1 // Height (Mutated)
1 // Depth (Mutated)
16 // ElementSize
16384 // RowPitch
0 // SlicePitch
4318 // Format (CL_RGBA)
4277 // Type (CL_FLOAT)
4337 // ImageType -> 4337 is CL_MEM_OBJECT_IMAGE2D!
The exact same behavior applies to IMAGE1D_ARRAY (Arg_1), which gets overwritten to CL_MEM_OBJECT_IMAGE2D_ARRAY (4339) because it also has depth == 0.
Environment
OS: Windows on ARM (ARM64)
Hardware: Qualcomm Snapdragon X Elite
Description
The Intercept Layer incorrectly deduces the type of 1D images and 1D image arrays. When an application creates a
CL_MEM_OBJECT_IMAGE1DorCL_MEM_OBJECT_IMAGE1D_ARRAY, the internal metadata stored inm_ImageInfoMapcategorizes them asIMAGE2DandIMAGE2D_ARRAYrespectively. This leads to corrupted metadata when generating replay captures (viaCLI_CaptureReplay=1), causing subsequent failures when parsing or replaying the serialized objects.Root Cause
In
intercept/src/clintercept_core.cpp, inside theCLIntercept::addImage(cl_mem image)function, the image type is deduced based on its dimensions (specifically checking for height == 0 and depth == 0). For a valid 1D image, both height == 0 and depth == 0 are true. The code correctly sets the ImageType to 1D in the first block, but immediately enters the second block and overwrites it.Steps to Repraoduce
set CLI_CaptureReplay=1 cliloader.exe ./test_app.exeThe metadata exported for the IMAGE1D argument shows it has been incorrectly overwritten to
CL_MEM_OBJECT_IMAGE2D(0x10F1 / 4337). Output extracted from theIMAGE1Dmetadata:The exact same behavior applies to
IMAGE1D_ARRAY(Arg_1), which gets overwritten toCL_MEM_OBJECT_IMAGE2D_ARRAY(4339) because it also has depth == 0.Environment
OS: Windows on ARM (ARM64)
Hardware: Qualcomm Snapdragon X Elite