[Mesa-dev] EXTERNAL: Re: Radeon r600_ring_test/evergreen_resume errors

[Tom Stellard]

On Fri, Feb 21, 2014 at 09:48:50PM +0000, Dorrington, Albert wrote:
> > -----Original Message-----
> > From: Tom Stellard [mailto:tom at stellard.net]
> > Subject: Re: [Mesa-dev] EXTERNAL: Re: Radeon
> > r600_ring_test/evergreen_resume errors
> > 
> > On Fri, Feb 21, 2014 at 05:53:02PM +0000, Dorrington, Albert wrote:
> > > > -----Original Message-----
> > > > From: Alex Deucher [mailto:alexdeucher at gmail.com] You are seeing a
> > > > GPU hang and the driver attempts to reset it which doesn't always
> > > > work.  Probably a problem in the OpenGL or OpenCL driver in mesa.
> > > >
> > > > Alex
> > >
> > > I assume there is some sort of watchdog timer monitoring the GPU, is
> > there a way to increase this timer duration?
> > > The results I am seeing are somewhat inconsistent, as the same code
> > doesn't appear to cause the crash in a repeatable manner.
> > >
> > 
> > Can you share the opencl kernel you are trying to compile?
> > 
> > -Tom
> 
> I have finally been able to capture the kernel that is being generated, and when I compiled on the command line I received an assert error
> 
> llc: /home/aldorr/opensrc/llvm_test/include/llvm/MC/MCRegisterInfo.h:65: unsigned int llvm::MCRegisterClass::getRegister(unsigned int) const: Assertion `i < getNumRegs() && "Register number out of range!"' failed.
> 
> I am assuming that the stalling is a result of the program i'm using not recognizing the fact that the compile was bad, and attempting to push a buffer full of something other than the program into the video card.
> 
> With being able to see the error from the command line build, I understand that something in the code being compiled is resulting in running out of registers. But I am not sure what that is.
> Here is the kernel code, which I believe is trying to verify that what goes in comes out the same. (Disclaimer: I didn't write this kernel code) ;-)
> 
> __kernel void test_kernel( __global long4 *in, __global long4 *out)
> {
>   __private long4 internal[ 32 ];

The problem is this array.  Private arrays are stored in registers, so it requires
2 * 4 * 32 = 256 registers.

We only store arrays in the X component of the vector registers, so this gives
us only 127 registers to use for arrays.

Private memory handling is broken for larger arrays like this, and I think we
should look into moving larger private arrays into either LDS or scratch memory.

-Tom

>   int tid = get_global_id( 0 );
> 
>   for( int i = 0; i < 32; i++ ) {
>     internal[ i ] = in[ i ];
>   }
>   out[ tid ] = internal[tid];
> }  
>