OpenCL

[mcphatty]

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

[dipak]

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. 

[mcphatty]

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.

[mcphatty]

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

[mcphatty]

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.

[dipak]

Hi @mcphatty ,

Thank you for providing the reproducible code and other information. We will look into this.

Regarding the source code sharing, the community does support file attachment option (see attached screenshot) which can be used to share source files and other information. However, due to some restriction, the "attachment" option might be temporarily disabled for new users. Hence you are not seeing it.

Thanks. 

[mcphatty]

I see, thanks.

[dipak]

Hi @mcphatty,

I was able to reproduce the issue. I've reported it to the OpenCL team. I'll let you know once I get any update on this.

Thanks.

[mcphatty]

Thank you, sounds good!

[mcphatty]

Hey @dipak,

are there news on the issue yet, or some info on if it will be possible to fix this in the future?

This is really a problem for us, as at the moment we can only recommend NVidia cards to our customers, which is not what I would prefer.

Kind regards

[dipak]

The concerned team is looking into this issue. Can't provide any timeline about the fix. I'll notify you if I get any news on this.

Thanks.