OpenCL* for CPU
Ask questions and share information on Intel® SDK for OpenCL™ Applications and OpenCL™ implementations for Intel® CPU
Announcements
This forum covers OpenCL* for CPU only. OpenCL* for GPU questions can be asked in the GPU Compute Software forum. Intel® FPGA SDK for OpenCL™ questions can be ask in the FPGA Intel® High Level Design forum.
1663 Discussions

Computer bench test LocalToneMapping issues

alfreddong
Novice
1,983 Views

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.

Labels (1)
0 Kudos
13 Replies
alfreddong
Novice
1,974 Views

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
Employee
1,941 Views

Hi,

 

On what system did you run it?

How did you call createRipMapY and createRipMapX separately?

 

Best regards,

Sergey

alfreddong
Novice
1,938 Views

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
Employee
1,921 Views

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
Novice
1,910 Views

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
Employee
1,890 Views

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
Novice
1,883 Views

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
Novice
1,859 Views

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
Novice
1,859 Views

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
Employee
1,844 Views

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
Novice
1,834 Views

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
Employee
1,821 Views

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
Novice
1,816 Views

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

Reply