Number of threads affecting answer; this should not happen a VERY strange error..

Accelerated Computing CUDA CUDA Programming and Performance
#1

Hi all

I have been having a very strange error. My kernel gives my “right” answer only if my number of threads per block is <=8 . Number of blocks don’t matter.

Generally, this would mean I have some problem in shared memory access (some type of conflict). But in my case there is NO inter-thread data transfer via shared memory, each thread just reuses it own bunch of shared memory in the kernel.

Hence as there is no co-orporation either at global memory or shared memory level , this shouldn’t happen rite ? :no: …!

I do launch multiple kernels with each kernel using the results from the previous kernels. But kernels cannot run concurrently, i guess ?

What I mean is something like this:

gpu_ford_phia<<<dimGrid,dimBlock>>>(gpu_space,ff,tempvar);		

	   for(int i=0;i<12;i++)

	   {	

	   gpu_ford_phib<<<dimGrid,dimBlock>>>(gpu_space,ff,ypass,tempvar,i);	

	   gpu_ford_phic<<<dimGrid2,dimBlock2>>>(gpu_space,ff,ypass,tempvar,i);	  

	   }

Here I am assuming kernels run one after the another.

The code runs fine and gives rite vales in device-Emu mode for any number of threads per block. :unsure:

For number of threads <=8 i get answer correct up-to 13 digit in device mode (am using double precision) but if my number of threads per block is say 64 my answers are TOTALLY wrong.

I don’t understand what is going on :( … I think it should work correctly.

I have attached the kernel here along with the main file.

Please is someone has time take a quick look if you can find any errors in the code :unsure: … am really desperate as my advisor wants some result in 2 days.

Also anyone has any clues why such a thing can happen ?

Thank you all… really appreciate any help

Nittin Arora
testcode.zip (4.69 KB)

#2

Check your code for races.

#3

I just see this code in your phic kernel…

for(k=0;k<3;k++)

			 a += ff[k+mul].x[indx]*sh[k][ind];

You are using “k” as FOR loop index here…

In the code above this, you are calculating the value of “k” in a loop.

And that value is completely destroyed by using “k” in the FOR loop…

Is that a bug?

#4

No its not a bug , the previous k value was just there for calculating the rite index to copy a column (last three values) to the sh variable.

The next use of k just loops over the 3 values stored and multiplies them with a another matrix in global memory. I am doing this to minimize register usage.

and I am sorry, but I don’t understand how can there be a race condition if there is no inter thread co-orporation going on… ?

… why does it only give correct answer run for <=8 threads per block…

thanks for looking at my code :)

NA

#5

how come im struggling to download the zip file… i feel seriously computer illiterate… keeps redirecting me. could you maybe pastebin it, or explain how to download it haha
aren’t threads per block supposed to be >=32 seeing that warpsize is 32. (not that it helps your problem - just a question)

#6

Hmm that’s strange… anyway…

Here is the paste-bin link to the code:

the kernels: link kernels

the main calling program : link main

and yes my actual block size is 320, I was just testing to find for which value of block size the problem occurs.

thanks

NA

#7

Try calling cudaThreadSynchronize() between your kernel calls.

#8

Nope that does not help… the kernels are executed serially anyway

Also I found one more thing, if I print results (after memcpy back to cpu) which were evaluated by threadidx.x = 0 ; they are right only when block size < = 8.

I think thread 0 should behave the same way irrespective of the block size ? :( … this is getting stranger

thanks all…

#9

I found the part of the code which is not getting executed properly

in kernel gpu_for_phib

switch(i)

										{

									   case 6:  ypass[i].x[indx] = 1.0 + h*cpy2;break;

										case 13: ypass[i].x[indx] = 1.0 + h*cpy2;break;				

										case 20: ypass[i].x[indx] = 1.0 + h*cpy2;break;

										case 27: ypass[i].x[indx] = 1.0 + h*cpy2;break;

										case 34: ypass[i].x[indx] = 1.0 + h*cpy2;break;

										case 41:ypass[i].x[indx]  = 1.0 + h*cpy2;break;

										default:ypass[i].x[indx]  = h*cpy2;break;

										}

This switch block is in an if else statement as shown in the kernel.cu file.

Investigating deeper I found out… both the switch statements in my code are not executing properly. If I convert them to if loops the code works fine for blocksize > 8. But why is this the case ?

I don’t understand what is wrong with this :mellow: , maybe the compiler is generating some wrong look up table or something or am I doing some stupid mistake here in these switch statements?