Computer bench test LocalToneMapping issues

[alfreddong]

It is running  properly when I use the computer bench 2.0 AP test Intel GPU(UHD Graphics 630) bench-mark score. However, it return the wrong number(-59) at clEnqueueNDRangeKernel() when executing the kernel named createRipMapX and createRipMapY of LocalToneMapping script case.

[alfreddong]

I want to know why these two kernels in LocalToneMapping case can not running on intel GPU UHD 630 separately, but the whole LocalToneMapping can run normally......

__kernel
void createRipMapX(__read_write image2d_t ripmap, const int offset, const int sizeX)
{
const uint idX = get_global_id(0);
const uint idY = get_global_id(1);
const uint globalSizeX = get_global_size(0);

float A = read_imagef(ripmap, (int2)(offset + 2*idX + 0, idY)).x;
float B = read_imagef(ripmap, (int2)(offset + 2*idX + 1, idY)).x;
write_imagef(ripmap, (int2)(offset + 2*globalSizeX + idX, idY), (A+B));

if(globalSizeX == 1)
return;

if (idX == 0 && idY == 0)
{
const int _offset = offset + 2*globalSizeX;
const size_t grid[2] = {globalSizeX/2, get_global_size(1)};

enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_2D(grid), ^{ createRipMapX(ripmap, _offset, sizeX); });
}
}

__kernel
void createRipMapY(__read_write image2d_t ripmap, const int offset, const int sizeX)
{
const uint idX = get_global_id(0);
const uint idY = get_global_id(1);
const uint globalSizeY = get_global_size(1);

float A = read_imagef(ripmap, (int2)(idX, offset + 2*idY + 0)).x;
float B = read_imagef(ripmap, (int2)(idX, offset + 2*idY + 1)).x;
write_imagef(ripmap, (int2)(idX, offset + 2*globalSizeY + idY), (A+B));

if(globalSizeY == 1)
return;

if (idX == 0 && idY == 0)
{
const int _offset = offset + 2*globalSizeY;
const size_t grid[2] = {get_global_size(0), globalSizeY/2};

enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_2D(grid), ^{ createRipMapY(ripmap, _offset, sizeX); });
}
}

[Sergey_I_Intel1]

Hi,

On what system did you run it?

How did you call createRipMapY and createRipMapX separately?

Best regards,

Sergey

[alfreddong]

Hi Sergey,

My computer's  os is Win10.

 I think createRipMapX and  createRipMapY are two symmetrical kernels. We can consider both of them as irrelevant two cases. Considering the kernel's first argument, It is initialized with the following code ......

But the clEnqueueNDRangeKernel API returns -50 error Num. I would like to know if there is an initialization problem or if the graphics card does not support this test?

int error = CL_SUCCESS;
cl_image_format format;
format.image_channel_data_type = CL_FLOAT;
format.image_channel_order = CL_R;
//cl_image_desc clImageDesc;
//memset(&clImageDesc, 0, sizeof(cl_image_desc));
//clImageDesc.image_type = CL_MEM_OBJECT_IMAGE2D;
//clImageDesc.image_width = 4096;
//clImageDesc.image_height = 1024;
unsigned size = 4096 * 1024 * 4;
input = malloc(size);
memset(input, 0, size);
input = load_data_with_file_name("***.bin");

ImgMem = clCreateImage2D(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, &format, 4096, 1024, 0, input, NULL);

Thank Alfred

[Sergey_I_Intel1]

Hi Alfred,

Please, try to specify this error with OpenCL official documentation.

It might be wrong usage of clEnqueueNDRangeKernel.

Can you share your full host and kernel code, so we can take a deep look at it.

Best regards,

Sergey

[alfreddong]

Hi Sergey,

I do not think I've misuse the clEnqueueNDRangeKernel. It works well on other platforms actually.

I wondering to know why it returns  error code(-59) whne  running the single CreatRipMapX kernel?

 The clTEnqueueNDRangeKerne API‘s configuration code as following:

global_work_size[0] = 0x100;
global_work_size[1] = 0x100;

local_work_size[0] = 0x40;
local_work_size[1] = 1;

ret |= clEnqueueNDRangeKernel(queue, kernel[0], 3, NULL, global_work_size, local_work_size, 0, NULL, NULL);

Thanks Alfred

[Sergey_I_Intel1]

Hi Alfred,

I assume, that local_work_size in this case is the reason of error.

We had such errors on Windows (on Linux it works), when using local_work_size in clEnqueueNDRangeKernel bring errors.

Try to set "NULL" there and check it out.

Best regards,

Sergey

[alfreddong]

Hi Sergey, 

I think it is unreasonable to set local_work_size to NULL.

Anyway, I tried it according to your suggestion, and it return the same error.

Thanks Alfred

[alfreddong]

Hi Sergey,

Are you working at Intel? If you are an employee of Intel, could you or your team may fix this issue?

Thanks Alfred

[alfreddong]

Hi Sergey,

Are you working at Intel? If you are an employee of Intel, could you or your team may fix this issue?

Thanks Alfred

[Sergey_I_Intel1]

Hi Alfred,

Yes, we are taking a look at your problem.

If you can sent us your host_code part, where you run it, it will be helpful to reproduce it.

By the way, what's your GPU driver version?

Best regards,

Sergey

[alfreddong]

Hi Sergey,

The host codes has little difference with any other common test, and I've send you the key configuration in previous mail.

I post the host code as an attachment now. You can check it if it is helpful for you.

By the way, I tested both driver versions, but neither one worked

Thanks Alfred

[Sergey_I_Intel1]

Hi Alfred,

So, -59 is CL_INVALID_OPERATION error.

CL_​INVALID_​OPERATION if SVM pointers are passed as arguments to a kernel and the device does not support SVM or if system pointers are passed as arguments to a kernel and/or stored inside SVM allocations passed as kernel arguments and the device does not support fine grain system SVM allocations.

Can you check on your machine, if that's the case? It might not work on some devices.

Here you can find sample to check it.

Best regards,

Sergey

[alfreddong]

Hi Sergey,

I've verified the fine grain system grain allocation(SVMBasicFineGrained.exe) can works on my machine(intel UHD630).

I am confused of why the sub-test can not work properly on the device but the whole test of 

LocalToneMapping can works well.

Thanks Alfred