cancel
Showing results for 
Search instead for 
Did you mean: 

OpenCL

mcphatty
Journeyman III

OpenCL: Wrong access to image2d_array_t from kernel, Radeon Pro W5500

 

Hello,

I've got a little problem with my OpenCL code, maybe someone can help me out with that.

The code works flawlessly on dedicated NVidia and integrated Intel graphics, but with my AMD Radion Pro W5500 I've encountered problems.

My code uses boost compute, and I reduced the problem to a minimal example. I'm pushing several images to a image2d_array_t, which seems to work fine, as after clEnqueueWriteImage I immediately call clEnqueueReadImage to verify the content of the array. The images are identical, so I guess in principle the pushed data should be ok.

Now, when I execute my minimal kernel, I'm using a linear sampler and image_readf to read from the image2d_array_t and write the content of a certain image in the array to a image2d_t. Then I get the conent of this result image back to the host and compare the conent to what I stuffed in into the array.

The kernel looks as follows:

__constant sampler_t samplerLinear  = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_LINEAR;

__kernel void arrayTest_kernel(__write_only image2d_t img_out, __read_only image2d_array_t img1, int w, int h, int arrImgIdx)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if (x >= w || y >= h)
        return;

    float4 resf = read_imagef(img1, samplerLinear, (float4)(x + 0.5f, y + 0.5f, arrImgIdx, 0.0f));

    if (x == 89 && y == 84)
        printf("linear, f: %f, %f, %f, %f\n", resf.x, resf.y, resf.z, resf.w);

    unsigned int out = 255 * resf.x;
    
    write_imageui(img_out, (int2)(x, y), (uint4)(out, 0, 0, 0));
}

 

There is a total of 5 images in the array. What I experience on my AMD card is that

  • the access at arrImgIdx = 0 fetches the 5th image, and not the 1st as expected
  • All other images of the array (arrImgIdx > 0) are just blank, meaning that the printf() statement in these cases also gives 0, 0, 0, 1

As the code works on all other platforms I assume either a driver related bug or some specific handling needed on AMD platforms I'm missing.

So the scenario is:

  • Copy multiple host images to image2d_array_t
  • In a kernel read these via read_imagef using (x, y, imgIdx, 0)
  • Is this working as expected on the W5500?

My driver version is: 20.45.40.04-210405a-366334c

Thank you in advance,

Phil

0 Likes
Reply
4 Replies
dipak
Community Manager
Community Manager

Re: OpenCL: Wrong access to image2d_array_t from kernel, Radeon Pro W5500

Hi @mcphatty ,

Thank you for reporting it. I've moved the post to the OpenCL developer forum. Also I've whitelisted you for the AMD developer community.

Please provide a minimal test-case (host code + kernel) that reproduces the issue and share the clinfo output and OS details.

Thanks. 

0 Likes
Reply
mcphatty
Journeyman III

Re: OpenCL: Wrong access to image2d_array_t from kernel, Radeon Pro W5500

Thank you, I had a little feeling I could be in the wrong place of the forums. 🙂

I'll provide a minmal example asap, could take me a moment.

0 Likes
Reply
mcphatty
Journeyman III

Re: OpenCL: Wrong access to image2d_array_t from kernel, Radeon Pro W5500

I've prepared a minimal example.

What is the preferred way of providing a code example in this community?

When I post a code example with rounded parentheses the site complains about that, even when putting it in a code snippet. It then prompts me to wait for 10 minutes to do my next post. Is there a possibility to post a file, or do I need to host my file somewhere and provide the link?

Just as a feedback: The 600 second rule is extremely annoying, is this limit lowered with time when earning status by posting?

Regards,

Phil

0 Likes
Reply
mcphatty
Journeyman III

Re: OpenCL: Wrong access to image2d_array_t from kernel, Radeon Pro W5500

clinfo 

main.cpp 

It seems there is a problem with parentheses which are quoted, so I've hosted the minmal example. I also hosted the output of clinfo.

The minimal example is just started in a command line, all devices are listed and you are prompted to enter the index of a listed device.

The test example then

  • creates two grayvalue images, one with vertical and one with horizontal white stripes on a black background
  • sends those images to an image array of size 2
  • uses a kernel to access each subimage of the array and write it to an output image
  • compares this output image element-wise to the respective input image

The result should in the optimal case be

"output image 1 matches input!"

"output image 2 matches input!"

On my ATI card both images do not match their inputs, as the behaviour equals the one in my original code.

  • Output image 1 equals the last array image (input image 2)
  • Output image 2 is just black

If opencv is present it is possible to write the images to a disk by commenting in WITH_OPENCV and specifying 'debugPath'. 

My operating system is Windows 10 Version 2009 10.0.19042.

 

 

0 Likes
Reply