3D texture problem

Accelerated Computing CUDA CUDA Programming and Performance
1

Hey

I’m new to cuda and I’m trying to use a 3D texture to manipulate data but I can’t seem to get it to work.

I can compile the code but all the outputs just turn out to be 0.

I tried this on a 2D case and it worked.

I believe it’s something wrong with copying the data.

#include <cuda.h>

#include <cutil.h>

#include <iostream>

#define max_H 3

#define max_W 3

#define max_D 3

using namespace std;

texture<float, 3, cudaReadModeElementType> tex;

__global__ void texTest(float* matrix_p)

{

	int index_x = blockIdx.x*blockDim.x+threadIdx.x;

	int index_y = blockIdx.y*blockDim.y+threadIdx.y;

	int index_z = blockIdx.z*blockDim.z+threadIdx.z;

	matrix_p[index_z*max_W*max_H+index_y*max_W+index_x]=2*tex3D(tex,index_x,index_y,index_z);

}

int main (void)

{

	float matrix [max_H][max_W][max_D];

	float *matrix_ptr;

	float matrix2 [max_H][max_W][max_D];

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

	{

		for (int j=0; j<max_W ; j++)

		{

			for(int k=0; k<max_D;k++)

			{

				matrix[i][j][k]=i*max_W*max_H+j*max_H + k;

				cout<<matrix[i][j][k]<<" ";

			}

			cout<<"\n"<<endl;

		}

		cout << endl;

	}

	cout <<"\n\n\n"<< endl;

	//array init

	cudaArray* cu_array=NULL;

	cudaChannelFormatDesc description = cudaCreateChannelDesc<float>(); 

	const cudaExtent extent = make_cudaExtent(max_W, max_H, max_D);

	cudaMalloc3DArray(&cu_array, &description,extent);

	//copy data

	cudaPitchedPtr pitch_ptr;

	cudaMalloc3D(&pitch_ptr, extent);

	cudaMemset3D(pitch_ptr, 0, extent);

	cudaMemcpy3DParms copyParams = {0};

	copyParams.srcPtr   = pitch_ptr;

	copyParams.dstArray = cu_array;

	copyParams.extent   = extent;

	copyParams.kind      = cudaMemcpyDeviceToDevice;

    CUDA_SAFE_CALL(cudaMemcpy3D(&copyParams));

	//texture parameters

	 

    tex.addressMode[0] = cudaAddressModeClamp;

    tex.addressMode[1] = cudaAddressModeClamp;

	tex.addressMode[2] = cudaAddressModeClamp;

    tex.filterMode = cudaFilterModePoint;

    tex.normalized = false; // do not normalize coordinates

	

	//Binding texture to data

	cudaBindTextureToArray(tex, cu_array,description);

	

	//Kernal call to manipulate data

	cudaMalloc((void**)&matrix_ptr,max_H*max_D*max_W*sizeof(float));

	dim3 block(max_W,max_H,max_D);

	texTest<<<1,block>>>(matrix_ptr);

	//Copying data from device back to host

	cudaMemcpy(matrix2,matrix_ptr,max_D*max_W*max_H*sizeof(float),cudaMemcpyDeviceToHost); 

	CUDA_SAFE_CALL( cudaUnbindTexture(tex) );

	CUDA_SAFE_CALL( cudaFreeArray(cu_array) );

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

	{

		for (int j=0; j<max_W ; j++)

		{

			for (int k=0; k<max_D ; k++)

			{

				cout<<matrix2[i][j][k]<<" ";

			}

			cout << endl;

		}

		cout << endl;

	}

	return 0;

}