[Mesa-dev] Mixing Pixel Shaders and Compute Shaders

[Dorrington, Albert]

In trying to implement Image support in Clover, I have discovered that the existing CL image related calls result in the generation of Pixel Shader sequences for copies of images to and from the GPU.

I initially thought that this would be fine, and was able to implement image read tests that use clEnqueueWriteImage() to get an image into a kernel.
The clEnqueueWriteImage(), through the routines in clover/api/transfer.cpp generates a Pixel shader which copies the image to the GPU.
The Compute Shader then picks the image up from where the Pixel Shader left it.

I had some issues initially with mixing the Pixel and Compute Shaders, until I added a PS_PARTIAL_FLUSH event along with the CS_PARTIAL_FLUSH event at the start of evergreen_init_atom_start_compute_cs(). I think this helped because it made the Pixel Shader Execute before the Compute Shader (not entirely sure?)

When I try to call clEnqueueReadImage(), after a clEnqueueNDRangeKernel(); the clover/aop/transfer.cpp again generates a Pixel Shader, which gets integrated into the command stream after the Compute Shader entries (so I send up with:

Command Sequence
1  [Initial Configuration]
2  [PixelShader]
3  [ComputeShader]
4  [PixelShader]
5  [Final Configuration/Cleanup/Wait]

The problem is, now I am encountering GPU Lockup CP Stalls at the end of 'section 4' and the start of 'section 5'
I am not sure I entirely understand why this is happening, but I know it has to do with the fact that the Pixel Shader is in the command stream after the Compute Shader commands.
I'm assuming something in how the flushes are configured for the Pixel Shader are not waiting for the Compute Shader to complete before executing, but again, I'm not entirely sure.

I figure there are two possible approaches to resolving this:

1.       Figure out the right way to get the Compute Shader and Pixel Shader to interact properly

2.       Do away with the need for the Pixel Shader by doing the image transfer entirely within the Compute Shader context. (Probably a lot of driver code to replace the existing routines that use the vbo and blitter draw routines?)

>From reviewing the R600/Evergreen register documentation, I see that the CB_COLOR#_INFO registers have a RAT bit (bit 26 in GPU registers 0x28c70-0x28ea4)
I also found that if this flag is set, that the surface is treated as a RAT and can only be manipulated by Compute Shader operations. (Which I suppose is the cause of the conflict between the Pixel Shader and Compute Shader trying to manipulate the same Color buffer/Texture.)

My biggest issue with this, is I have not really found any documentation that describes how you are supposed to transfer buffers/textures within a compute shader, so I feel like I am missing something that might be a very basic foundation for understanding these routines, which is resulting in my overcomplicating the concepts and confusing myself...

If anyone is familiar with this area and is willing to provide some more insight, I would greatly appreciate it.

While our team's goal is to implement OpenCL capability in an alternate operating system, my hope is that once I understand all of this and get it working in that environment, I will be able to contribute back Clover image support to the main Mesa baseline.

Thanks,
Al Dorrington
Software Engineer Sr
Lockheed Martin, Mission Systems and Training

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/mesa-dev/attachments/20140423/553e473b/attachment-0001.html>