0

I want to use texture memory of GPU for a 1D array and use it in a kernal. I create a texture object and and assigned resDesc.res.linear.devPtr to a device array to bind my desired array to texture object. The problem is that when I fetch the data, it is different from the data I bounded to the texture memory. The output of this code should be 1, 4, 9, ... while it is 0,0,0,....

#include<cuda.h>
#include<cuda_runtime.h>
#include<iostream>
#include<stdio.h>
using namespace std;
// Simple transformation kernel
__global__ void squareKernel(float* output,float *dh, int size,  cudaTextureObject_t texObj) {
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    if(x>size)return;
    float y=tex1D<float>(texObj, x);
    printf("%d, %f, %f\n", x, y, dh[x]);
    output[x]=y*y;
}
#define width 10
int main() {
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(width, 0, 0, 0,
            cudaChannelFormatKindFloat);
    float *hA=(float *)malloc(width*sizeof(float));
    for(int i=0;i<width;i++){
        hA[i]=i;
    }
    float *dA;
    cudaMalloc((void **)&dA, width*sizeof(float));
    cudaMemcpy(dA, hA, width*sizeof(float), cudaMemcpyHostToDevice);
    
    cudaResourceDesc resDesc;
    memset(&resDesc, 0, sizeof(resDesc));

    resDesc.resType = cudaResourceTypeLinear;
    resDesc.res.linear.devPtr=dA;
    resDesc.res.linear.sizeInBytes=width*sizeof(float);
    
    // Specify texture object parameters
    cudaTextureDesc texDesc;
    memset(&texDesc, 0, sizeof(texDesc));

    texDesc.filterMode = cudaFilterModeLinear;
    texDesc.readMode = cudaReadModeElementType;
    texDesc.normalizedCoords = 0;
    // Create texture object
    cudaTextureObject_t texObj = 0;
    cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);



    float* output;
    cudaMalloc(&output, width * sizeof(float));
    squareKernel<<<1, width>>>(output, dA, width,texObj);
    cudaMemcpy(hA, output, width*sizeof(float), cudaMemcpyDeviceToHost);
    for(int i=0; i<width;i++){
        cout<<i<<","<<hA[i]<<endl;
    }
    
    // Destroy texture object
    cudaDestroyTextureObject(texObj);
    // Free device memory

    cudaFree(dA);
    free(hA);
    return 0;
}
4
  • "The problem is that when I fetch the data" -- what problem would that be? Could you explain in a bit more detail about what isn't working? Commented Sep 18, 2023 at 8:46
  • @talonmies, I already edited the question. Commented Sep 18, 2023 at 10:23
  • 2
    OK. So it doesn't seem you have defined a channel description anywhere in your code. I would guess that cudaCreateTextureObject is failing with an error, but you don't check for runtime errors, so no-one could say for sure Commented Sep 18, 2023 at 10:45
  • 2
    # compute-sanitizer ./t47 ========= COMPUTE-SANITIZER ========= Program hit cudaErrorInvalidChannelDescriptor (error 20) due to "invalid channel descriptor" on CUDA API call to cudaCreateTextureObject. Commented Sep 19, 2023 at 23:34

1 Answer 1

-1

I think you can simply replace the tex1D<float>(texObj, x); by tex1Dfetch<float>(texObj, x);

One use tex1D when the resource data is stored in a cuArray.

Sign up to request clarification or add additional context in comments.

3 Comments

How will that fix a texture definition without a valid channel description?
@talonmies Sorry, I did not see the details about the channel description in this code, but I think the cause of the zero output is the usage of tex1D, and I found a very similar question here.
The texture isn’t even created. Did you try running your answer? I guess not..

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Start asking to get answers

Find the answer to your question by asking.

Ask question

Explore related questions

See similar questions with these tags.