topic Re: Computer bench test LocalToneMapping issues in OpenCL* for CPU

Computer bench test LocalToneMapping issues

alfreddong — Tue, 07 Jul 2020 01:39:31 GMT

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.

Re: Computer bench test LocalToneMapping issues

alfreddong — Tue, 07 Jul 2020 08:37:17 GMT

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); });
}
}

Re: Computer bench test LocalToneMapping issues

Sergey_I_Intel1 — Mon, 20 Jul 2020 10:09:32 GMT

Hi,

On what system did you run it?

How did you call createRipMapY and createRipMapX separately?

Best regards,

Sergey

Re: Computer bench test LocalToneMapping issues

alfreddong — Mon, 20 Jul 2020 10:45:30 GMT

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

Re: Computer bench test LocalToneMapping issues

Sergey_I_Intel1 — Thu, 23 Jul 2020 10:41:12 GMT

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

Re: Computer bench test LocalToneMapping issues

alfreddong — Fri, 24 Jul 2020 05:19:21 GMT

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

Re: Computer bench test LocalToneMapping issues

Sergey_I_Intel1 — Fri, 24 Jul 2020 10:23:26 GMT

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

Re: Computer bench test LocalToneMapping issues

alfreddong — Fri, 24 Jul 2020 11:19:39 GMT

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

Re: Computer bench test LocalToneMapping issues

alfreddong — Wed, 29 Jul 2020 06:02:34 GMT

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

Re: Computer bench test LocalToneMapping issues

alfreddong — Wed, 29 Jul 2020 06:03:48 GMT

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

Re: Computer bench test LocalToneMapping issues

Sergey_I_Intel1 — Wed, 29 Jul 2020 09:27:13 GMT

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

Re: Computer bench test Local Tone Mapping issues

alfreddong — Thu, 30 Jul 2020 01:39:47 GMT

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

Re: Computer bench test Local Tone Mapping issues

Sergey_I_Intel1 — Thu, 30 Jul 2020 13:03:30 GMT

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

Re: Computer bench test Local Tone Mapping issues

alfreddong — Fri, 31 Jul 2020 02:22:27 GMT

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