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.
1718 Discussions

Computer bench test LocalToneMapping issues

alfreddong
Novice
3,418 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
3,409 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); });
}
}

0 Kudos
Sergey_I_Intel1
Employee
3,376 Views

Hi,

 

On what system did you run it?

How did you call createRipMapY and createRipMapX separately?

 

Best regards,

Sergey

0 Kudos
alfreddong
Novice
3,373 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

0 Kudos
Sergey_I_Intel1
Employee
3,356 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

0 Kudos
alfreddong
Novice
3,345 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

0 Kudos
Sergey_I_Intel1
Employee
3,325 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

0 Kudos
alfreddong
Novice
3,318 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

0 Kudos
alfreddong
Novice
3,294 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

0 Kudos
alfreddong
Novice
3,294 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

0 Kudos
Sergey_I_Intel1
Employee
3,279 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

0 Kudos
alfreddong
Novice
3,269 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

0 Kudos
Sergey_I_Intel1
Employee
3,256 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
0 Kudos
alfreddong
Novice
3,251 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

0 Kudos
Reply