Random errors on texture accesses "Uknown error"

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

I’m writing a simple Laplace equation solver using Gauss-Seidel SOR with red/black ordering using texture memory. The code is depicted bellow:

[codebox]

global void kcalcRedBlack(unsigned int offset, float *dataD, unsigned int N, unsigned int width, float omega){

const unsigned int iy = blockIdx.y * BLOCK_SIZENY + threadIdx.y + 1;

const unsigned int ix = blockIdx.x * BLOCK_SIZENX + threadIdx.x;

char line_indicator = (offset+iy)%2;

unsigned int ptrid = ix + __umul24(iy, width);

bool inbounds = ix<N && iy<N-1;

if( inbounds ){

	if( ix>=line_indicator && ix<N-(1-line_indicator) && (line_indicator+ix)%2 ){

		float oldv = tex1Dfetch(texData, ptrid);

		float newv = (1.0f-omega)*oldv+omega*(tex1Dfetch(texData, ptrid-1)+tex1Dfetch(texData, ptrid+1)+tex1Dfetch(texData, ptrid-width)+tex1Dfetch(texData, ptrid+width))/4.0f;

		dataD[ptrid] = newv;

	}

}

}

[/codebox]

texData is bound to the same memory space that dataD points to.

At random execution points a get an “Unknown error”. Sometimes it runs correctly.

Does it have to do with CUDA version or driver? I’m using CUDA 2.0 with driver 185.85 (latest ver).

Thanks in advance.

2

Hello again,

I upgraded to CUDA 2.2 but situation stays the same.

I have attached a simplified source code for anyone who wishes to reproduce it.

Code sometimes terminates with “unknown error” and the screen flashes for a moment:

[codebox]

Matrix dimensions 1024x1024 with h=1/1023

Parameters omega=1.993881

ERROR: unknown error

[/codebox]

and sometimes runs fine

[codebox]

Matrix dimensions 1024x1024 with h=1/1023

Parameters omega=1.993881

Total iterations 5000

38.45 seconds total time

[/codebox]

I still can’t find the problem.
ktxtest.zip (1.36 KB)

3

I suggest you change the tex1Dfetch lines to something like this:

tex1Dfetch( texData, threadIdx.x )

that way you run something consistant and you’re sure you’re not going out of bounds.

If that goes well, and gives you consistant results (even though they are garbage) then you

probably have syncronization or out-of-bounds issues with your texture accesses.

I doubt it is a driver/Cuda issue :)

hope that helps.

eyal

4

[quote name=‘eyalhir74’ post=‘548165’ date=‘Jun 3 2009, 01:44 PM’]

I suggest you change the tex1Dfetch lines to something like this:

[codebox]__global__ void kcalc(unsigned int offset, float *dataD, unsigned int N, unsigned int width, float omega){ 
const unsigned int iy = blockIdx.y * BLOCK_SIZEN + threadIdx.y + 1;

const unsigned int ix = blockIdx.x * BLOCK_SIZEN + threadIdx.x;

char line_indicator = (offset+iy)%2;

unsigned int ptrid = ix + iy*width;

if( (ix<N-1) && (ix>0) && (iy<N-1) && (iy>0) && ((line_indicator+ix)%2) ){

	if( ptrid>width && ptrid<(N-1)*width )

		dataD[ptrid] = (1.0f-omega)*tex1Dfetch(texData, ptrid)+omega*(tex1Dfetch(texData, ptrid-1)+tex1Dfetch(texData, ptrid+1)+tex1Dfetch(texData, ptrid-width)+tex1Dfetch(texData, ptrid+width))/4.0f;

// dataD[threadIdx.x] = (1.0f-omega)tex1Dfetch(texData, ptrid)+omega(tex1Dfetch(texData, ptrid-1)+tex1Dfetch(texData, ptrid+1)+tex1Dfetch(texData, ptrid-width)+tex1Dfetch(texData, ptrid+width))/4.0f;

// dataD[ptrid] = 1.0f;

}

}[/codebox]

Consider this:

If I change the line so all writes are done to dataD[threadIdx.x] instead of dataD[ptrid] then no errors occur. So that can’t be texture accesses out of bounds.

If I change the line so just a constant value is written to dataD[ptrid] without any texture reads then again no errors occur.

So this is confusing and I can’t figure out what is causes it. I’m thinking that this might be a driver issue (bug?).

Is there any other reason that might be the cause of this error (“Unknown error”)? I’m trying to locate it for days without a reasonable explanation.

P.S. I’m using Windows Vista (32bit)
ktxtest.zip (1.36 KB)

5

On Windows XP the error is :

ERROR: unspecified launch failure

6

Now I tried using 2D texture to linear memory (with function cudaBindTexture2D() provided by CUDA 2.2) instead of 1D texture and it seems to work without errors.

Kernel function is now like this :

[codebox]global void kcalc(unsigned int offset, float *dataD, unsigned int N, unsigned int width, float omega){

const unsigned int iy = blockIdx.y * BLOCK_SIZEN + threadIdx.y + 1;

const unsigned int ix = blockIdx.x * BLOCK_SIZEN + threadIdx.x;

char line_indicator = (offset+iy)%2;

unsigned int ptrid = ix + iy*width;

if( (ix<N-1) && (ix>0) && (iy<N-1) && (iy>0) && ((line_indicator+ix)%2) ){

	if( ptrid>width && ptrid<(N-1)*width ){

		dataD[ptrid] = (1.0f-omega)*tex2D(texData, ix, iy)+omega*(tex2D(texData, ix-1, iy)+tex2D(texData, ix+1, iy)+tex2D(texData, ix, iy-1)+tex2D(texData, ix, iy+1))/4.0f;

	}

}

}

[/codebox]

However, I would like to know why the 1D texture caused the random errors.