[Mesa-dev] Mixing Pixel Shaders and Compute Shaders

[Tom Stellard]

On Wed, Apr 23, 2014 at 01:27:11PM +0000, Dorrington, Albert wrote:
> 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?)
>

PS_PARTIAL_FLUSH guarantees all pixel shader have completed before the
next packet is processed, so we should be using this any place where we
may need the result of a pixel shader.
 
> 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]
Need PS_PARTIAL_FLUSH here
> 3  [ComputeShader]
Need CS_PARTIAL_FLUSH here
> 4  [PixelShader]
Need PS_PARTIAL_FLUSH here
> 5  [Final Configuration/Cleanup/Wait]
> 

You also may need to flush the various caches after the pixel shader
and compute shaders have completed.  See r600_flush_emit in r600_hw_context.c

> 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.)
>

Does the documentation say that setting the RAT bit means it can only be
used by compute shaders, or have you discovered this from your testing?
If this is the case, you may have to use a different CB_COLOR# for the
image when coying it with a pixel shader.

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

Looking forward to your contribution.  As always the sooner you can post the code the better
as that will make it easier to review and may help uncover some of your issues.

-Tom

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

> _______________________________________________
> mesa-dev mailing list
> mesa-dev at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/mesa-dev