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

[Dorrington, Albert]

> -----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 ];
  int tid = get_global_id( 0 );

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