<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40"> <head> <meta http-equiv="Content-Type" content="text/html; charset=us-ascii"> <meta name="Generator" content="Microsoft Word 14 (filtered medium)"> <style></style> </head> <body lang="EN-US" link="blue" vlink="purple"> <div class="WordSection1"> 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. <o:p></o:p> <o:p> </o:p> 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.<o:p></o:p> The clEnqueueWriteImage(), through the routines in clover/api/transfer.cpp generates a Pixel shader which copies the image to the GPU.<o:p></o:p> The Compute Shader then picks the image up from where the Pixel Shader left it.<o:p></o:p> <o:p> </o:p> 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?)<o:p></o:p> <o:p> </o:p> 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: <o:p></o:p> <o:p> </o:p> Command Sequence<o:p></o:p> 1 [Initial Configuration]<o:p></o:p> 2 [PixelShader]<o:p></o:p> 3 [ComputeShader]<o:p></o:p> 4 [PixelShader]<o:p></o:p> 5 [Final Configuration/Cleanup/Wait]<o:p></o:p> <o:p> </o:p> The problem is, now I am encountering GPU Lockup CP Stalls at the end of ‘section 4’ and the start of ‘section 5’<o:p></o:p> 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.<o:p></o:p> 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.<o:p></o:p> <o:p> </o:p> I figure there are two possible approaches to resolving this:<o:p></o:p> <![if !supportLists]>1. <![endif]>Figure out the right way to get the Compute Shader and Pixel Shader to interact properly<o:p></o:p> <![if !supportLists]>2. <![endif]>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?)<o:p></o:p> <o:p> </o:p> 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)<o:p></o:p> 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.)<o:p></o:p> <o:p> </o:p> 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...<o:p></o:p> <o:p> </o:p> If anyone is familiar with this area and is willing to provide some more insight, I would greatly appreciate it.<o:p></o:p> <o:p> </o:p> 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.<o:p></o:p> <o:p> </o:p> Thanks,<o:p></o:p> Al Dorrington<o:p></o:p> Software Engineer Sr<o:p></o:p> Lockheed Martin, Mission Systems and Training<o:p></o:p> <o:p> </o:p> </div> </body> </html>