Sample code for linear memory -> 3d texture

Does anyone have some sample code for

  1. allocating a device linear memory space
  2. writing it with a kernel
  3. converting/copying that space to a 3d texture
  4. using that texture in another kernel

Thanks!

I seem to have answered my own question, so ill post it here in case it can help someone else.

float  *d_density;

CUDA_SAFE_CALL( cudaMalloc((void **)&d_density,sizeof(float)*num_xyz) );

convert_ct_density<<<dimGrid,dimBlock,0>>>(d_ct,d_density); //write to d_density

cudaArray *density_arr;

	cudaChannelFormatDesc floatTex = cudaCreateChannelDesc<float>();;

	const cudaExtent volumeSize = make_cudaExtent(map->x_dim, map->y_dim, map->z_dim);

	CUDA_SAFE_CALL( cudaMalloc3DArray(&density_arr, &floatTex, volumeSize) );

prepareCudaTexture(d_density,density_arr,floatTex,volumeSize);

static void prepareCudaTexture(float* d_source, cudaArray *a_Data,cudaChannelFormatDesc floatTex, const cudaExtent volumeSize)

{

	cudaMemcpy3DParms copyParams = {0};

	copyParams.srcPtr   = make_cudaPitchedPtr((void*)d_source, volumeSize.width*sizeof(float), volumeSize.width, volumeSize.height);

	copyParams.dstArray = a_Data;

    copyParams.extent   = volumeSize;

    copyParams.kind     = cudaMemcpyDeviceToDevice;

    CUDA_SAFE_CALL( cudaMemcpy3D(&copyParams) );

	densityData.normalized = false;                      

    densityData.filterMode = cudaFilterModePoint;      

    densityData.addressMode[0] = cudaAddressModeWrap;   // wrap texture coordinates

    densityData.addressMode[1] = cudaAddressModeWrap;

    densityData.addressMode[2] = cudaAddressModeWrap;

   // bind array to 3D texture

    CUDA_SAFE_CALL(cudaBindTextureToArray(densityData, a_Data, floatTex));

}

Use at your own risk for now, this code is 1h old!

1 Like