[Mesa-dev] EXTERNAL: Re: Mixing Pixel Shaders and Compute Shaders

[Dorrington, Albert]

> -----Original Message-----
> From: Dorrington, Albert
> > -----Original Message-----
> > From: Tom Stellard>
> > On Wed, Apr 23, 2014 at 01:27:11PM +0000, Dorrington, Albert wrote:
> <snip>
> > > 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 <-- this one I think I have with what I added
> to evergren_init_atom_start_compute_cs()
> 3  [ComputeShader]
> Need CS_PARTIAL_FLUSH here
> 4  [PixelShader]
> Need PS_PARTIAL_FLUSH here
> 5  [Final Configuration/Cleanup/Wait]
> 
> The other two, I've been experimenting trying to add them in various places,
> but I haevn't seen a change in behavior yet (probably not putting them in the
> right place yet...)
> 

I spent most of my day experimenting with the CS and PS Partial Flushes, at the end of the Compute Shader Dispatch and the start and end of the r600_draw_vbo(), and I have had no luck.\
I only seem to be moving the stall point around the command sequence.

I wrote a program to take the data words from a command stream and decode them. I have attached a file which includes my latest attempt.
In this attempt, the trace_cs indicates the lockup occured at/after dw 0x000005be
This is in the Pixel Shader Code generated by the clEnqueueReadImage() call that goes through the transfer_map routines with r600_draw_vbo()...

I don't know who might understand these command sequences well enough suggest what might be causing the stall.

The program that generates this is fairly simple. It creates an Image and a Buffer, it populates the buffer and enqueues it, it enqueues the kernel, then does an clEnqueueReadImage to retrieve the image. All the kernel does is set some pixels within the image to known values.

The clSetKernelArg() for the image argument generates the first Pixel Shader, the clEnqueueNDRangeKernel() generates the Compute Shader, and then the clEnqueueReadImage() generates the last PixelShader.

I am starting to think that this approach of using the existing Texture/Pixel Shader code to move the images in and out of the kernel just won't work, and that the compute shader should be able to do this through the configured RAT instead.

Thanks
-Al

-------------- next part --------------
An embedded and charset-unspecified text was scrubbed...
Name: stall10.txt
URL: <http://lists.freedesktop.org/archives/mesa-dev/attachments/20140423/72972696/attachment-0001.txt>