Take Garbage Value wrong output how to use shared memory in a program

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,

I read a program to find the Multiplication of SquareMatrix Using shared Memory.But It gives the wrong result(outputsnapshot)and my expected result is other(shapshot expect).

Tell me?

How Can I get a exacted output?

Thanks in Advance

#include<stdio.h> #include<cuda.h> __global__ void Shar(float *a,float *b,float *c,int n) {	__shared__ float aTile[4][4],bTile[4][4];	int row=blockIdx.y*blockDim.y+threadIdx.y;	int col=blockIdx.x*blockDim.x+threadIdx.x;	float sum=0.0;	aTile[threadIdx.y][threadIdx.x]=a[row*n+threadIdx.x];	bTile[threadIdx.y][threadIdx.x]=b[threadIdx.y*n+col];	__syncthreads();	for(int i=0;i<n;i++)	{	sum+=aTile[threadIdx.y][i]*bTile[i][threadIdx.x];	}	c[row*n+col]=sum; } int main() {	float *a_h,*b_h,*c_h,*a_d,*b_d,*c_d;	int i,n;	n=4;	size_t size=sizeof(float)*(n*n);	a_h=(float*)malloc(size);	b_h=(float*)malloc(size);	c_h=(float*)malloc(size);	cudaMalloc((void**)&a_d,size);	cudaMalloc((void**)&b_d,size);	cudaMalloc((void**)&c_d,size);	for(i=0;i<(n*n);i++)	{	a_h[i]=1;	}	cudaMemcpy(a_d,a_h,size,cudaMemcpyHostToDevice);	for(i=0;i<(n*n);i++)	{ b_h[i]=1;	}	cudaMemcpy(b_d,b_h,size,cudaMemcpyHostToDevice);	int blocksize=4;	int nblock=n/blocksize+(n%blocksize==0?0:1);	int TILE_DIM=4;	Shar<<<nblock,blocksize,TILE_DIM>>>(a_d,b_d,c_d,n);	cudaMemcpy(c_h,c_d,size,cudaMemcpyDeviceToHost);	printf("\nMultiplication Of Matrix");	for(i=0;i<(n*n);i++)	{	printf("\n%f",c_h[i]);	}	free(a_h);	free(b_h);	free(c_h);	cudaFree(a_d);	cudaFree(b_d);	cudaFree(c_d);	return 0; }

expect.png
expect.png1366×768 3.25 KB

output.png
output.png1366×768 5.57 KB

2

You forgot to add:

cudaThreadSynchronize();

after kernels execution. Therefore you might copy arrays before kernels stop operation on them.

Additionally If someone use shared memory, extra care should be taken to prevent collisions. I have many time received unstable results because one of the threads was reading from shared memory befor other have put something there.

3

No, CUDA implicitly synchronizes between operations on the same stream. (If you don’t specify a stream, you are on stream 0.) You do not need cudaThreadSynchronize() for anything but benchmarking kernel runtime (and perhaps one obscure use of zero-copy memory).

You should check the return codes on your CUDA calls to see if your kernel is being run at all.