works in one dimension, but not two dimensions?!

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi All,

I’m pretty new to programming and to CUDA, and I’m trying to write a program for image dilation (just brute force).

And sorry for the title, I don’t know how else to describe my problem.

My CPU version of the code works perfectly:

void dilateOnHost(matrix *a, matrix *b, matrix *c,int *result){

	int i,j,down,across,max,value,resultIndex, sampleIndex, aheight, bheight;

	

	for(j=0;j<c->height;j++){	//moves filter down

		for (i=0;i<c->width;i++){	//moves filter across

			resultIndex = i+j*c->width;

			sampleIndex = i+j*a->width;

			result[resultIndex]=0;

			max = 0;

			for(down=0;down<b->height;down++){		//multiplies each column

				aheight = down*a->width;

				bheight = down*b->width;

				for(across=0;across<b->width;across++){		//multiplies each row

					value = b->arrayPtr[across+bheight] * a->arrayPtr[sampleIndex + across+aheight];

					if(value>max)

						max = value;

				}

			}

			result[resultIndex]=max;

		}

	}

}

where I’m dilating sample, ‘a’, by structural element ‘b’, and the result is ‘c’ but i’m storing the result in ‘result’

But my CUDA code does not produce correct results, even though everything is pretty much the same:

I’m using a 6x6 grid and 11x11 threadblock because that’s the size of the matrix i’m dealing with.

//CPU code

		int numThreadsPerBlock = 11;

	dim3 dimGrid(66/numThreadsPerBlock, 66/numThreadsPerBlock); //dimensions of the grid in terms of number of blocks

	dim3 dimBlock(numThreadsPerBlock,numThreadsPerBlock); //dimension of block in terms of numebr of threads

	//launch gpu kernel

	dilateOnDevice<<< dimGrid, dimBlock >>>(d_sample, d_f, d_result, d_widths, d_heights);

	   ....

//GPU code

__global__ void dilateOnDevice(int *d_sample, int *d_f, int *d_result, int *d_widths, int *d_heights)

{

	//cuPrintf("4 numbers are: %d %d %d %d", d_widths[0], d_widths[1], d_widths[2], d_heights[0]);

	int sampleWidth, kernelWidth, resultWidth, kernelHeight;

	sampleWidth = d_widths[0];

	kernelWidth = d_widths[1];

	resultWidth = d_widths[2];

	kernelHeight = d_heights[0];

	

	int down, max, across, value, i,j, resultIndex, sampleIndex, aheight, bheight;

	j = blockIdx.y * blockDim.y + threadIdx.y;

	i = blockIdx.x * blockDim.x + threadIdx.x;

	resultIndex = i+j*resultWidth;

	sampleIndex = i+j*sampleWidth;

	

	d_result[resultIndex]=0;

	max = 0;

	

	for(down=0;down<kernelHeight;down++){		//multiplies each column

		aheight = down*sampleWidth;

		bheight = down*kernelWidth;

		for(across=0;across<kernelWidth;across++){		//multiplies each row

			value = d_f[across+bheight] * d_sample[sampleIndex + across+aheight];

			if(value>max)

				max = value;

		}

	}

	d_result[resultIndex]=max;

}

when I replaced

i = blockIdx.x * blockDim.x + threadIdx.x;

with

for (i=0;i<resultWidth;i++){

then the code works fine.

Any ideas why?

2

I can’t immediately see anything wrong with it, which seems to suggest that resultWidth≠66.

3

I can’t immediately see anything wrong with it, which seems to suggest that resultWidth≠66.

4

yea you’re right.

resultWidth = 64

kernelWidth = 3

sampleWidth = 66

kernelHeight = 3

sample is already zero-padded

5

yea you’re right.

resultWidth = 64

kernelWidth = 3

sampleWidth = 66

kernelHeight = 3

sample is already zero-padded

6

[quote name=‘KChou’ post=‘1124665’ date=‘Sep 30 2010, 11:02 PM’]

Hi All,

I’m pretty new to programming and to CUDA, and I’m trying to write a program for image dilation (just brute force).

And sorry for the title, I don’t know how else to describe my problem.

My CPU version of the code works perfectly:

[codebox][font="Courier New"]#include

#include <cutil_inline.h>

#define IMW 11

#define IMH 66

#define FILTW 3

#define FILTH 3

class matrix

{

public:

int width,height;

int *h_arrayPtr,*d_arrayPtr;

matrix *d_matrix;

matrix() { h_arrayPtr=d_arrayPtr=NULL; d_matrix=NULL; width=0; height=0; }

matrix(int w, int h)

{

	width=w;height=h;

	d_arrayPtr=NULL; d_matrix=NULL;

	h_arrayPtr=(int *)malloc(width*height*sizeof(*h_arrayPtr));

}

~matrix()

{

	if (h_arrayPtr) free(h_arrayPtr);

	if (d_arrayPtr) cudaFree(d_arrayPtr);

	if (d_matrix) cudaFree(d_matrix);

}

int matrixCPUelement(int i, int j) { if ((unsigned)i<(unsigned)height && (unsigned)j<(unsigned)width) return h_arrayPtr[i*width+j]; else return 0; }

int matrixCPUelement(int i) { if ((unsigned)i<(unsigned)(width*height)) return h_arrayPtr[i]; else return 0; }

__device__ int matrixGPUelement(int i, int j) { if ((unsigned)i<(unsigned)height && (unsigned)j<(unsigned)width) return d_arrayPtr[i*width+j]; else return 0; }

__device__ int matrixGPUelement(int i) { if ((unsigned)i<(unsigned)(width*height)) return d_arrayPtr[i]; else return 0; }

void fillRandom(int seed);

int copytodevice();

};

int matrix::copytodevice()

{

cudaMalloc((void**)&d_matrix,sizeof(*d_matrix));

cudaMalloc((void**)&d_arrayPtr,width*height*sizeof(*h_arrayPtr));

cudaMemcpy(d_arrayPtr,h_arrayPtr,width*height*sizeof(*h_arra

yPtr),cudaMemcpyHostToDevice);

cudaMemcpy(d_matrix,this,sizeof(*d_matrix),cudaMemcpyHostToD

evice);

return 1;

}

void matrix::fillRandom(int s)

{

srand(s);

for (int h=0;h<height;h++)

	for (int w=0;w<width;w++)

		h_arrayPtr[h*width+w]=rand();

}

global void dilateOnGPU(matrix *dest, matrix *image, matrix *kernel)

{

int h=blockIdx.x,w=threadIdx.x;

int resultIndex = w+h*dest->width;

int sampleIndex = w+h*image->width;

dest->d_arrayPtr[resultIndex]=0;

int maxv = 0;

for(int down=0;down<kernel->height;down++)					//multiplies each column

{

	int a_rowstart = down*image->width;

	int b_rowstart = down*kernel->width;

	for(int across=0;across<kernel->width;across++)        //multiplies each row

	{

		int value = kernel->matrixGPUelement(b_rowstart+across) * image->matrixGPUelement(sampleIndex + a_rowstart+across);

		if(value>maxv) maxv = value;

	}

}

dest->d_arrayPtr[resultIndex]=maxv;

}

void dilateOnHost(matrix &dest, matrix &image, matrix &kernel)

{

for(int h=0;h<dest.height;h++)    //moves filter down

{

    for (int w=0;w<dest.width;w++)    //moves filter across

	{

        int resultIndex = w+h*dest.width;

        int sampleIndex = w+h*image.width;

        dest.h_arrayPtr[resultIndex]=0;

        int maxv = 0;

        for(int down=0;down<kernel.height;down++)					//multiplies each column

		{

            int a_rowstart = down*image.width;

            int b_rowstart = down*kernel.width;

            for(int across=0;across<kernel.width;across++)        //multiplies each row

			{

                int value = kernel.matrixCPUelement(b_rowstart+across) * image.matrixCPUelement(sampleIndex + a_rowstart+across);

                if(value>maxv) maxv = value;

            }

        }

		dest.h_arrayPtr[resultIndex]=maxv;

    }

}

}

int main(void)

{

int gpu_dev;

cudaSetDevice(gpu_dev=cutGetMaxGflopsDeviceId());

matrix img(IMW,IMH),filt(FILTW,FILTH),result(IMW,IMH);

img.fillRandom(2010);

filt.fillRandom(2011);

dilateOnHost(result,img,filt);

img.copytodevice();filt.copytodevice();result.copytodevice()

;

cudaMemset(result.d_arrayPtr,0,result.width*result.height*si

zeof(*result.d_arrayPtr)); // not really necessary…

dilateOnGPU<<<IMH,IMW>>>(result.d_matrix,img.d_matrix,filt.d_matrix);

cutilCheckMsg( "Kernel execution failed" );

cudaThreadSynchronize();

// check gpu results

int *arraygpucpy=(int*)malloc(result.width*result.height*sizeof(

*result.d_arrayPtr));

cudaMemcpy(arraygpucpy,result.d_arrayPtr,result.width*result

.height*sizeof(*result.d_arrayPtr),cudaMemcpyDeviceToHost);

int errors=0;

for (int n=0;n<result.height*result.width;n++)

	errors+=arraygpucpy[n]!=result.h_arrayPtr[n];

printf("%d errors, %s\n",errors,errors==0?"PASSED":"FAILED");



free(arraygpucpy);

return 0;[/font]

}[/codebox]

When I compile this as .cu with VS2008 toolkit 3.2 it runs fine. For a more efficient implementation, the convolutionSeparable example in the SDK might help.

7

[quote name=‘KChou’ post=‘1124665’ date=‘Sep 30 2010, 11:02 PM’]

Hi All,

I’m pretty new to programming and to CUDA, and I’m trying to write a program for image dilation (just brute force).

And sorry for the title, I don’t know how else to describe my problem.

My CPU version of the code works perfectly:

[codebox][font="Courier New"]#include

#include <cutil_inline.h>

#define IMW 11

#define IMH 66

#define FILTW 3

#define FILTH 3

class matrix

{

public:

int width,height;

int *h_arrayPtr,*d_arrayPtr;

matrix *d_matrix;

matrix() { h_arrayPtr=d_arrayPtr=NULL; d_matrix=NULL; width=0; height=0; }

matrix(int w, int h)

{

	width=w;height=h;

	d_arrayPtr=NULL; d_matrix=NULL;

	h_arrayPtr=(int *)malloc(width*height*sizeof(*h_arrayPtr));

}

~matrix()

{

	if (h_arrayPtr) free(h_arrayPtr);

	if (d_arrayPtr) cudaFree(d_arrayPtr);

	if (d_matrix) cudaFree(d_matrix);

}

int matrixCPUelement(int i, int j) { if ((unsigned)i<(unsigned)height && (unsigned)j<(unsigned)width) return h_arrayPtr[i*width+j]; else return 0; }

int matrixCPUelement(int i) { if ((unsigned)i<(unsigned)(width*height)) return h_arrayPtr[i]; else return 0; }

__device__ int matrixGPUelement(int i, int j) { if ((unsigned)i<(unsigned)height && (unsigned)j<(unsigned)width) return d_arrayPtr[i*width+j]; else return 0; }

__device__ int matrixGPUelement(int i) { if ((unsigned)i<(unsigned)(width*height)) return d_arrayPtr[i]; else return 0; }

void fillRandom(int seed);

int copytodevice();

};

int matrix::copytodevice()

{

cudaMalloc((void**)&d_matrix,sizeof(*d_matrix));

cudaMalloc((void**)&d_arrayPtr,width*height*sizeof(*h_arrayPtr));

cudaMemcpy(d_arrayPtr,h_arrayPtr,width*height*sizeof(*h_arra

yPtr),cudaMemcpyHostToDevice);

cudaMemcpy(d_matrix,this,sizeof(*d_matrix),cudaMemcpyHostToD

evice);

return 1;

}

void matrix::fillRandom(int s)

{

srand(s);

for (int h=0;h<height;h++)

	for (int w=0;w<width;w++)

		h_arrayPtr[h*width+w]=rand();

}

global void dilateOnGPU(matrix *dest, matrix *image, matrix *kernel)

{

int h=blockIdx.x,w=threadIdx.x;

int resultIndex = w+h*dest->width;

int sampleIndex = w+h*image->width;

dest->d_arrayPtr[resultIndex]=0;

int maxv = 0;

for(int down=0;down<kernel->height;down++)					//multiplies each column

{

	int a_rowstart = down*image->width;

	int b_rowstart = down*kernel->width;

	for(int across=0;across<kernel->width;across++)        //multiplies each row

	{

		int value = kernel->matrixGPUelement(b_rowstart+across) * image->matrixGPUelement(sampleIndex + a_rowstart+across);

		if(value>maxv) maxv = value;

	}

}

dest->d_arrayPtr[resultIndex]=maxv;

}

void dilateOnHost(matrix &dest, matrix &image, matrix &kernel)

{

for(int h=0;h<dest.height;h++)    //moves filter down

{

    for (int w=0;w<dest.width;w++)    //moves filter across

	{

        int resultIndex = w+h*dest.width;

        int sampleIndex = w+h*image.width;

        dest.h_arrayPtr[resultIndex]=0;

        int maxv = 0;

        for(int down=0;down<kernel.height;down++)					//multiplies each column

		{

            int a_rowstart = down*image.width;

            int b_rowstart = down*kernel.width;

            for(int across=0;across<kernel.width;across++)        //multiplies each row

			{

                int value = kernel.matrixCPUelement(b_rowstart+across) * image.matrixCPUelement(sampleIndex + a_rowstart+across);

                if(value>maxv) maxv = value;

            }

        }

		dest.h_arrayPtr[resultIndex]=maxv;

    }

}

}

int main(void)

{

int gpu_dev;

cudaSetDevice(gpu_dev=cutGetMaxGflopsDeviceId());

matrix img(IMW,IMH),filt(FILTW,FILTH),result(IMW,IMH);

img.fillRandom(2010);

filt.fillRandom(2011);

dilateOnHost(result,img,filt);

img.copytodevice();filt.copytodevice();result.copytodevice()

;

cudaMemset(result.d_arrayPtr,0,result.width*result.height*si

zeof(*result.d_arrayPtr)); // not really necessary…

dilateOnGPU<<<IMH,IMW>>>(result.d_matrix,img.d_matrix,filt.d_matrix);

cutilCheckMsg( "Kernel execution failed" );

cudaThreadSynchronize();

// check gpu results

int *arraygpucpy=(int*)malloc(result.width*result.height*sizeof(

*result.d_arrayPtr));

cudaMemcpy(arraygpucpy,result.d_arrayPtr,result.width*result

.height*sizeof(*result.d_arrayPtr),cudaMemcpyDeviceToHost);

int errors=0;

for (int n=0;n<result.height*result.width;n++)

	errors+=arraygpucpy[n]!=result.h_arrayPtr[n];

printf("%d errors, %s\n",errors,errors==0?"PASSED":"FAILED");



free(arraygpucpy);

return 0;[/font]

}[/codebox]

When I compile this as .cu with VS2008 toolkit 3.2 it runs fine. For a more efficient implementation, the convolutionSeparable example in the SDK might help.

8

nope, I believe you’re looking at my serial code. It runs fine.

Let me try to rephrase, I’m getting the correct output with the serial code, which is this:

[codebox]

111111111111111111111100111111111111111111111111111111111111

0111

111111111111111111111111011111111111111110111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111101111111111111111

1111

011101111111111111111111111111111111111111111111111111111111

0111

011111111111111111111111111111111111111111111111111111110011

1111

111111111111111111111111111111111111111110111111111111111111

1111

111100111111111111111111111111111111111111111111111111111111

1101

111111111110111111111111111111111111111111111111111111111111

1110

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111011111111111000111111111111111111111111111111111111

1011

111111111111111111110101111111111111111011111111111111111111

1111

111111111111111111111111111110111101111011111111111111111111

1111

111111111111111111111111011100111111010111111111101111111011

1111

111111111111111111111111111110011111100111111110011110111101

1111

111111111111111111111111111111111111101111111111111111111111

1111

111111111111111111111111111111111011111111111110111111111111

1111

111111111111111111111111011111011111111111111111111101111111

1111

111111111111111111111110111111111111111111111111111111111111

1111

111111111111111111110101111110111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111011

1111

111111111111111100111111111111111011111111111111111111111111

1111

111111111111111101111111111111011111111111111110111111111101

1110

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111011111111111111111111111111111111111111

1111

111111111111111111111001111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111101111111111110111

1111

111111111111111110111111111111111111111111111111111111111111

1111

111111110111111111111111111111111111011111111111111101011111

1111

111111111111111111111111111111111110111111111110111011111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

011111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111100111111111111111111111111111111111011111111111111

1111

111111001101111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111110101111111111111111111111110111111

1111

111111111111110111111111110101111111111111111111111111111111

1111

111111111111101111110111111111111111111111111111111111111111

1111

111111111111111111101111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111011111111

1111

111111111111111111111111111111111111111111111111111101111111

1111

111111111111111111111111111111001111111111111111111111111111

1111

111111111111111111111111111111110111111111111111111111111111

1111

110111111111111111111111111111111111111111111111011111111111

1011

111111111111111111111111111111111111111111110111111111111101

1111

111111001111111111111111111111111110111111111111111111111111

1111

111111001111111111011111111111111111111111111111111111111111

1111

111111100111111111111101111101111111111111111111111111111111

1111

111111000111111111111111111111111111111111110111111111111111

1111

111111100111111111111111111111111111111111111111111111111111

1111

111111101111111110111111101111111111111111111111111111111111

1111

111111011111111111111111111111111111111111111111111111111111

1111

111111101111111111111111111111111111111111110110111111011111

1111

011111101111111111110111111111111011111111111111011111111111

1111

111111111111111111111111111111111111111111111111011111011111

1110

111111111111111111111111001111111111111111111111111111111111

1100

111111111111111111111111111111111111111111111111111111111111

1111

100111111111111111111111111111111111111111111111111111111111

1111

[/codebox]

I also get the same correct result with this version of my parallel code:

[codebox]

111111111111111111111100111111111111111111111111111111111111

0111

111111111111111111111111011111111111111110111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111101111111111111111

1111

011101111111111111111111111111111111111111111111111111111111

0111

011111111111111111111111111111111111111111111111111111110011

1111

111111111111111111111111111111111111111110111111111111111111

1111

111100111111111111111111111111111111111111111111111111111111

1101

111111111110111111111111111111111111111111111111111111111111

1110

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111011111111111000111111111111111111111111111111111111

1011

111111111111111111110101111111111111111011111111111111111111

1111

111111111111111111111111111110111101111011111111111111111111

1111

111111111111111111111111011100111111010111111111101111111011

1111

111111111111111111111111111110011111100111111110011110111101

1111

111111111111111111111111111111111111101111111111111111111111

1111

111111111111111111111111111111111011111111111110111111111111

1111

111111111111111111111111011111011111111111111111111101111111

1111

111111111111111111111110111111111111111111111111111111111111

1111

111111111111111111110101111110111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

101111111111111111111111111111111111111111111111111111111011

1111

111111111111111100111111111111111011111111111111111111111111

1111

011111111111111101111111111111011111111111111110111111111101

1110

001111111111111111111111111111111111111111111111111111111111

1111

011111111111111111111011111111111111111111111111111111111111

1111

011111111111111111111001111111111111111111111111111111111111

1111

011111111111111111111111111111111111111111101111111111110111

1111

101111111111111110111111111111111111111111111111111111111111

1111

111111110111111111111111111111111111011111111111111101011111

1111

011111111111111111111111111111111110111111111110111011111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

011111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111100111111111111111111111111111111111011111111111111

1111

111111001101111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111110101111111111111111111111110111111

1111

111111111111110111111111110101111111111111111111111111111111

1111

111111111111101111110111111111111111111111111111111111111111

1111

111111111111111111101111111111111111111111111111111111111111

1111

011111111111111111111111111111111111111111111111111111111111

1111

101111111111111111111111111111111111111111111111111011111111

1111

001111111111111111111111111111111111111111111111111101111111

1111

011111111111111111111111111111001111111111111111111111111111

1111

111111111111111111111111111111110111111111111111111111111111

1111

010111111111111111111111111111111111111111111111011111111111

1011

111111111111111111111111111111111111111111110111111111111101

1111

001111001111111111111111111111111110111111111111111111111111

1111

101111001111111111011111111111111111111111111111111111111111

1111

101111100111111111111101111101111111111111111111111111111111

1111

111111000111111111111111111111111111111111110111111111111111

1111

111111100111111111111111111111111111111111111111111111111111

1111

011111101111111110111111101111111111111111111111111111111111

1111

011111011111111111111111111111111111111111111111111111111111

1111

111111101111111111111111111111111111111111110110111111000000

0000

011111101111111111110111111111111011111111111111011111100000

0000

111111111111111111111111111111111111111111111111011111000000

0000

111111111111111111111111001111111111111111111111111111100000

0000

111111111111111111111111111111111111111111111111111111100000

0000

100111111111111111111111111111111111111111111111111111100000

0000

[/codebox]

I get extra zeros at random places, and all I did was replace

for (i=0;i<resultWidth;i++)

with

i = blockIdx.x * blockDim.x + threadIdx.x;

sorry for how the outputs don’t line up exactly.

would this be caused by rounding errors??? but comparing 0’s and 1’s shouldn’t result in rounding errors, right?

9

nope, I believe you’re looking at my serial code. It runs fine.

Let me try to rephrase, I’m getting the correct output with the serial code, which is this:

[codebox]

111111111111111111111100111111111111111111111111111111111111

0111

111111111111111111111111011111111111111110111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111101111111111111111

1111

011101111111111111111111111111111111111111111111111111111111

0111

011111111111111111111111111111111111111111111111111111110011

1111

111111111111111111111111111111111111111110111111111111111111

1111

111100111111111111111111111111111111111111111111111111111111

1101

111111111110111111111111111111111111111111111111111111111111

1110

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111011111111111000111111111111111111111111111111111111

1011

111111111111111111110101111111111111111011111111111111111111

1111

111111111111111111111111111110111101111011111111111111111111

1111

111111111111111111111111011100111111010111111111101111111011

1111

111111111111111111111111111110011111100111111110011110111101

1111

111111111111111111111111111111111111101111111111111111111111

1111

111111111111111111111111111111111011111111111110111111111111

1111

111111111111111111111111011111011111111111111111111101111111

1111

111111111111111111111110111111111111111111111111111111111111

1111

111111111111111111110101111110111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111011

1111

111111111111111100111111111111111011111111111111111111111111

1111

111111111111111101111111111111011111111111111110111111111101

1110

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111011111111111111111111111111111111111111

1111

111111111111111111111001111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111101111111111110111

1111

111111111111111110111111111111111111111111111111111111111111

1111

111111110111111111111111111111111111011111111111111101011111

1111

111111111111111111111111111111111110111111111110111011111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

011111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111100111111111111111111111111111111111011111111111111

1111

111111001101111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111110101111111111111111111111110111111

1111

111111111111110111111111110101111111111111111111111111111111

1111

111111111111101111110111111111111111111111111111111111111111

1111

111111111111111111101111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111011111111

1111

111111111111111111111111111111111111111111111111111101111111

1111

111111111111111111111111111111001111111111111111111111111111

1111

111111111111111111111111111111110111111111111111111111111111

1111

110111111111111111111111111111111111111111111111011111111111

1011

111111111111111111111111111111111111111111110111111111111101

1111

111111001111111111111111111111111110111111111111111111111111

1111

111111001111111111011111111111111111111111111111111111111111

1111

111111100111111111111101111101111111111111111111111111111111

1111

111111000111111111111111111111111111111111110111111111111111

1111

111111100111111111111111111111111111111111111111111111111111

1111

111111101111111110111111101111111111111111111111111111111111

1111

111111011111111111111111111111111111111111111111111111111111

1111

111111101111111111111111111111111111111111110110111111011111

1111

011111101111111111110111111111111011111111111111011111111111

1111

111111111111111111111111111111111111111111111111011111011111

1110

111111111111111111111111001111111111111111111111111111111111

1100

111111111111111111111111111111111111111111111111111111111111

1111

100111111111111111111111111111111111111111111111111111111111

1111

[/codebox]

I also get the same correct result with this version of my parallel code:

[codebox]

111111111111111111111100111111111111111111111111111111111111

0111

111111111111111111111111011111111111111110111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111101111111111111111

1111

011101111111111111111111111111111111111111111111111111111111

0111

011111111111111111111111111111111111111111111111111111110011

1111

111111111111111111111111111111111111111110111111111111111111

1111

111100111111111111111111111111111111111111111111111111111111

1101

111111111110111111111111111111111111111111111111111111111111

1110

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111011111111111000111111111111111111111111111111111111

1011

111111111111111111110101111111111111111011111111111111111111

1111

111111111111111111111111111110111101111011111111111111111111

1111

111111111111111111111111011100111111010111111111101111111011

1111

111111111111111111111111111110011111100111111110011110111101

1111

111111111111111111111111111111111111101111111111111111111111

1111

111111111111111111111111111111111011111111111110111111111111

1111

111111111111111111111111011111011111111111111111111101111111

1111

111111111111111111111110111111111111111111111111111111111111

1111

111111111111111111110101111110111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

101111111111111111111111111111111111111111111111111111111011

1111

111111111111111100111111111111111011111111111111111111111111

1111

011111111111111101111111111111011111111111111110111111111101

1110

001111111111111111111111111111111111111111111111111111111111

1111

011111111111111111111011111111111111111111111111111111111111

1111

011111111111111111111001111111111111111111111111111111111111

1111

011111111111111111111111111111111111111111101111111111110111

1111

101111111111111110111111111111111111111111111111111111111111

1111

111111110111111111111111111111111111011111111111111101011111

1111

011111111111111111111111111111111110111111111110111011111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

011111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111100111111111111111111111111111111111011111111111111

1111

111111001101111111111111111111111111111111111111111111111111

1111

111111111111111111111111111111111111111111111111111111111111

1111

111111111111111111111111110101111111111111111111111110111111

1111

111111111111110111111111110101111111111111111111111111111111

1111

111111111111101111110111111111111111111111111111111111111111

1111

111111111111111111101111111111111111111111111111111111111111

1111

011111111111111111111111111111111111111111111111111111111111

1111

101111111111111111111111111111111111111111111111111011111111

1111

001111111111111111111111111111111111111111111111111101111111

1111

011111111111111111111111111111001111111111111111111111111111

1111

111111111111111111111111111111110111111111111111111111111111

1111

010111111111111111111111111111111111111111111111011111111111

1011

111111111111111111111111111111111111111111110111111111111101

1111

001111001111111111111111111111111110111111111111111111111111

1111

101111001111111111011111111111111111111111111111111111111111

1111

101111100111111111111101111101111111111111111111111111111111

1111

111111000111111111111111111111111111111111110111111111111111

1111

111111100111111111111111111111111111111111111111111111111111

1111

011111101111111110111111101111111111111111111111111111111111

1111

011111011111111111111111111111111111111111111111111111111111

1111

111111101111111111111111111111111111111111110110111111000000

0000

011111101111111111110111111111111011111111111111011111100000

0000

111111111111111111111111111111111111111111111111011111000000

0000

111111111111111111111111001111111111111111111111111111100000

0000

111111111111111111111111111111111111111111111111111111100000

0000

100111111111111111111111111111111111111111111111111111100000

0000

[/codebox]

I get extra zeros at random places, and all I did was replace

for (i=0;i<resultWidth;i++)

with

i = blockIdx.x * blockDim.x + threadIdx.x;

sorry for how the outputs don’t line up exactly.

would this be caused by rounding errors??? but comparing 0’s and 1’s shouldn’t result in rounding errors, right?

10

Rounding errors is out of the question (ints). But as I said, I think you read beyond the (input)array and therefore get random errors.

Please try my version of your code and see if it produces correct results. It does on my machine.

11

Rounding errors is out of the question (ints). But as I said, I think you read beyond the (input)array and therefore get random errors.

Please try my version of your code and see if it produces correct results. It does on my machine.

12

I tried running your code. Well, I incorporated it into my code. It still does not work.

The result I’m getting is still having extra zeros at random places.

Could it be how my CUDA is set up?

I’ll keep working on this and restructure my code to be more like yours, and see if anything changes.

13

I tried running your code. Well, I incorporated it into my code. It still does not work.

The result I’m getting is still having extra zeros at random places.

Could it be how my CUDA is set up?

I’ll keep working on this and restructure my code to be more like yours, and see if anything changes.