On Fri, Feb 21, 2014 at 09:48:50PM +0000, Dorrington, Albert wrote:
> > -----Original Message-----
> > From: Tom Stellard [mailto:t...@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:alexdeuc...@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];
> }
>
_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/mesa-dev
