Extending sobel example to 3D volume raster

Accelerated Computing CUDA CUDA Programming and Performance
1

So I have a 3D texture. I’ve used it to perform a 3D gaussian, and now I need to do a 3D sobel filter.

The given sdk example is quite obscured compared to the sobel theory. I’d implement naively but having X,Y, and Z temporaries (3 times that if I separate the convolutions), wouldn’t fit in device memory, so I need something efficient.

Example 3D sobel kernels are here

http://www.aravind.ca/cs788h_Final_Project…mators.htm#idsb

Essentially I want to take these two routines from the sobel sdk example and make them perform a 3D sobel using tex3D, and the like

__device__ unsigned char

ComputeSobel(unsigned char ul, // upper left

			 unsigned char um, // upper middle

			 unsigned char ur, // upper right

			 unsigned char ml, // middle left

			 unsigned char mm, // middle (unused)

			 unsigned char mr, // middle right

			 unsigned char ll, // lower left

			 unsigned char lm, // lower middle

			 unsigned char lr, // lower right

			 float fScale )

{

	short Horz = ur + 2*mr + lr - ul - 2*ml - ll;

	short Vert = ul + 2*um + ur - ll - 2*lm - lr;

	short Sum = (short) (fScale*(abs(Horz)+abs(Vert)));

	if ( Sum < 0 ) return 0; else if ( Sum > 0xff ) return 0xff;

	return (unsigned char) Sum;

}

__global__ void 

SobelTex( Pixel *pSobelOriginal, unsigned int Pitch, 

		  int w, int h, float fScale )

{ 

	unsigned char *pSobel = 

	  (unsigned char *) (((char *) pSobelOriginal)+blockIdx.x*Pitch);

	for ( int i = threadIdx.x; i < w; i += blockDim.x ) {

		unsigned char pix00 = tex2D( tex, (float) i-1, (float) blockIdx.x-1 );

		unsigned char pix01 = tex2D( tex, (float) i+0, (float) blockIdx.x-1 );

		unsigned char pix02 = tex2D( tex, (float) i+1, (float) blockIdx.x-1 );

		unsigned char pix10 = tex2D( tex, (float) i-1, (float) blockIdx.x+0 );

		unsigned char pix11 = tex2D( tex, (float) i+0, (float) blockIdx.x+0 );

		unsigned char pix12 = tex2D( tex, (float) i+1, (float) blockIdx.x+0 );

		unsigned char pix20 = tex2D( tex, (float) i-1, (float) blockIdx.x+1 );

		unsigned char pix21 = tex2D( tex, (float) i+0, (float) blockIdx.x+1 );

		unsigned char pix22 = tex2D( tex, (float) i+1, (float) blockIdx.x+1 );

		pSobel[i] = ComputeSobel(pix00, pix01, pix02, 

								 pix10, pix11, pix12,

								 pix20, pix21, pix22, fScale );

	}

}
2

no ideas?

3

Why don’t you make a general 3D convolver? It’s not that hard.