topic Hi ScottJ, in OpenCL* for CPU

Very basic beginner question using CodeBuilder and a sample shader

johnson__scott1 — Mon, 13 May 2019 21:23:41 GMT

I am trying a sample shader that intends to simply copy the pixels form a input texture to an output texture. My goal is to see it run in Intel CodeBuilder without host code. I have setup a session to run this small kernel. (The kernel is from the textbook "OpenCL In Action").

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

__kernel void simple_image(read_only image2d_t src_image, write_only image2d_t dst_image)
{
  int2 coord = (int2) (get_global_id(0), get_global_id(1));
  uint4 pixel = read_imageui(src_image, sampler, coord);

  write_imageui(dst_image, coord, pixel);
}

I've created two image variables. The input image is 642x482 RGBA. So I made the output image have the same properties as the input image except it is labeled as an output image. I assigned the image variables to the input and outputs in the "Code Builder Analysis Input" window. I selected a global size of X:642 Y:482 Z:0. The local size is (1, 1, 0).

When I run the kernel in CodeBuilder it builds fine and runs but the output image is always black. The simple kernel code should copy the image. I've tried setting "pixel" to a nonzero constant value but still the output image shows up black in the output report.

Hi ScottJ,

Michael_C_Intel1 — Tue, 14 May 2019 00:22:25 GMT

Hi ScottJ,

Thanks for the interest... the feedback is useful ... I did get walking through the process just now... but I had trouble hardcoding a pixel value as well... But for me it was because I didn't set my channel data type correctly.

For some pointers to debug...

A formatted printf(...) can be dropped in the kernel... it may be fruitful to add some code like this to check the data the kernel sees emitted by say work-item 0,0...

  if(get_global_id(0) == 0 && get_global_id(1) == 0)
      printf("%u %u %u %u\n", pixel[0], pixel[1], pixel[2], pixel[3]);

If my input variable usez CL_UNSIGNED_INT8 as the channel data type... when I dump the uint4 pixel I see for a toy image @ 0,0 I get: 237 27 36 255...

If I leave the image object as CL_UNORM_INT8 channel type I see 1064168942 1037621465 1041272977 1065353216 for the same pixel in a successfully validated kernel write... Per spec... using read_imageui is undefined if the the channel data type isn't friendly:

Reference:

read_imageui can only be used with image objects created with image_channel_data_type set to one of the following values: CL_UNSIGNED_INT8, CL_UNSIGNED_INT16, andCL_UNSIGNED_INT32. If the image_channel_data_type is not one of the above values, the values returned by read_imageui are undefined.

This may have affected your attempts at hard coding... With a real host side program... this would be defined in host code. Similar spec limitations are documented for the unsigned write version of the function as well.

The Code Builder Variable pane can help edit variable types in the Codebuilder environment.

On ISS2019u3 w/ MSVS2017 on Intel® Core™ i5-6300U I was able to assign a 600x446 mspaint toy jpg to src_image and have the session validate successfully. I'll attach screen shots to give a general reference... I'm hoping some old fashioned printf(...) and checking the buffer types can help.

Let me know if you're still stuck.

-MichaelC

How were you able to use CL

johnson__scott1 — Tue, 14 May 2019 15:12:51 GMT

How were you able to use CL_UNSIGNED_INT8 as the channel data type? I get the error "Failed to create an image for this argument (dst_image): -10 (CL_IMAGE_FORMAT_NOT_SPECIFIED).

simple_out

Data Type : image2d_t

Width : 4

Height : 4

Depth : 1

Array Size : 1

Row Pitch : 0

Slice Pitch : 0

Source Format : RGB - RGBA

Channel Data Type : CL_UNSIGNED_INT8

Channel Order : CL_BGRA

IO Mode : Output

Memory Flags : CL_MEM_USE_HOST_PTR

Hi ScottJ,

Michael_C_Intel1 — Tue, 21 May 2019 19:26:38 GMT

Hi ScottJ,

Thanks again for the feedback. This feedback encourages me to file clarifications for multiple issues with the product.

Difficult to tell where there is immediate divergence between your configuration and mine...

A few comments:

I think there's a sighting here for an error code issue... per OCL docs it should be CL_IMAGE_FORMAT_NOT_SUPPORTED... CL_IMAGE_FORMAT_NOT_SPECIFIED looks like a typo in the support tool backend. I will file this to the dev team.
On image formats related to that error... it's important to match the OCL manual for the cl_image_format description data structure. "These formats can only be used if channel data type = "
What is simple_out? I only use the same variable names in the Assign Variable pane as I use in the Code Builder Configuration Assign Input pane... I've attached a screen shot of a successful configuration with dst_image reduced to a 4x4.
Also notice the channel order is different than mine.

I was successful running only 4x4 workitems and size for the dst_image memory object. Of course that didn't validate against the much larger original image....

Can you post a screen shot of the error to this thread if you still can't get it to operate? Can you archive the codebuilder project directory and post it here?

Thanks,

-MichaelC

For reference for Intel

Michael_C_Intel1 — Tue, 21 May 2019 22:55:32 GMT

For reference for Intel employee viewers. I've filed OPENCLSDK-4908,4909,and 4910 related to this posting... thanks.

-MichaelC