Bind CUDA Texture to a float Image

cuda, textures

Solution

When you use the normalized coordinates the texture is accessed via the coordinates from 0 to 1 (exclusive). You forgot translate your integer threadIdx-based coordinates into normalized.

unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; 
float u = x / (float)width;  
float v = y / (float)height;

Problem

I have a 1 channel, float image in C side like the following: ``` int width, height; float* img; ``` I want to pass this image to a CUDA texture. I'm reading the NVIDIA CUDA C Programming Guide (page 42-43) and using the tutorial, wrote a code like the following: main.cpp: ``` int main() { int width, height; float* h_Input; ReadImage(&h_Input, &width, &height); // My function which reads the image. WriteImage(h_Input, width, height); // works perfectly... float* h_Output = (float*) malloc(sizeof(float) * width * height); CalculateWithCuda(h_Input, h_Output, width,height); WriteImage(h_Output, width, height); // writes an empty-gray colored image.... *WHY???* } ``` kernel.cu: ``` texture<float, cudaTextureType2D, cudaReadModeElementType> texRef; // 2D float texture __global__ void Kernel(float* output, int width, int height) { int i = blockIdx.y * blockDim.y + threadIdx.y; // row number int j = blockIdx.x * blockDim.x + threadIdx.x; // col number if(i < height && j < width) { float temp = tex2D(texRef, i + 0.5f, j + 0.5f); output[i * width + j] = temp ; } } void CalculateWithCuda(const float* h_input, float* h_output, int width, int height) { float* d_output; // Allocate CUDA array in device memory cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindFloat); cudaArray* cuArray; cudaMallocArray(&cuArray, &channelDesc, width, height); // Copy to device memory some data located at address h_data in host memory cudaMemcpyToArray(cuArray, 0, 0, h_input, width * height * sizeof(float) , cudaMemcpyHostToDevice); // Set texture parameters texRef.addressMode[0] = cudaAddressModeWrap; texRef.addressMode[1] = cudaAddressModeWrap; texRef.filterMode = cudaFilterModeLinear; texRef.normalized = true; // Bind the array to the texture reference cudaBindTextureToArray(texRef, cuArray, channelDesc); // Allocate GPU buffers for the output image .. cudaMalloc(&d_output, sizeof(float) * width * height); dim3 threadsPerBlock(16,16); dim3 numBlocks((width/threadsPerBlock.x) + 1, (height/threadsPerBlock.y) + 1); Kernel<<<numBlocks, threadsPerBlock>>>(d_output, width,height); cudaDeviceSynchronize(); // Copy output vector from GPU buffer to host memory. cudaMemcpy(h_output, d_output, sizeof(float) * width * height, cudaMemcpyDeviceToHost); // Free GPU memory ... } ``` As I said in the code; this Kernel has to read from the texture and give me the same image as output. However, I'm taking an empty (grays colored) image for the output. I just implemented the same way in the tutorial, why does not this texture work? I will be appreciated if somebody show me a way to fix this issue ... PS: Sure, it's not the all of the code. I just copied the necessary parts. If you need other details, I will support as well. Thanks in advance.

Original source