Does anyone have some sample code for
- allocating a device linear memory space
- writing it with a kernel
- converting/copying that space to a 3d texture
- using that texture in another kernel
Thanks!
Does anyone have some sample code for
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(©Params) );
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!