cuda - double 的纹理对象

Question

我想将纹理对象(而不是引用)与 double 值一起使用。下面的代码在使用 float 时有效，但 double 不是受支持的数据类型。

我可以使用 2d 纹理来解决这个问题吗？如果可以，我该如何设置这样的纹理？

纹理引用有类似的问题，但纹理对象没有。 Support for double type in texture memory in CUDA

__global__ void my_print(cudaTextureObject_t texObject)
{
    printf("%f\n",tex1Dfetch<double>(texObject,0));

    return;
}

int main()
{

    double i = 0.35;
    int numel = 50;

    double* d_data;
    cudaMalloc(&d_data,numel*sizeof(double));
    cudaMemcpy((void*)d_data,&i,1*sizeof(double), cudaMemcpyHostToDevice);


    cudaTextureDesc td;
    memset(&td, 0, sizeof(td));

    td.normalizedCoords = 0;
    td.addressMode[0] = cudaAddressModeClamp;
    td.readMode = cudaReadModeElementType;


    struct cudaResourceDesc resDesc;
    memset(&resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeLinear;
    resDesc.res.linear.devPtr = d_data;
    resDesc.res.linear.sizeInBytes = numel*sizeof(double);
    resDesc.res.linear.desc.f = cudaChannelFormatKindFloat;
    resDesc.res.linear.desc.x = 32;

    cudaTextureObject_t texObject = 0;
    gpuErrchk(cudaCreateTextureObject(&texObject, &resDesc, &td, NULL));

    my_print<<<1,1>>>(texObject);

    gpuErrchk(cudaDeviceSynchronize());
    return 0;
}

Answer 1

这个想法与纹理引用完全相同。您可以通过将数据绑定(bind)到支持的 64 位类型并将结果读取转换为 double 来访问 double 。如果您像这样修改代码:

#include <vector>
#include <cstdio>

static __inline__ __device__ double fetch_double(uint2 p){
    return __hiloint2double(p.y, p.x);
}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}
__global__ void my_print(cudaTextureObject_t texObject)
{
    uint2 rval = tex1Dfetch<uint2>(texObject, 0);
    double dval = fetch_double(rval);
    printf("%f\n", dval);
}

int main()
{

    double i = 0.35;
    int numel = 50;

    std::vector<double> h_data(numel, i);
    double* d_data;
    cudaMalloc(&d_data,numel*sizeof(double));
    cudaMemcpy((void*)d_data, &h_data[0], numel*sizeof(double), cudaMemcpyHostToDevice);


    cudaTextureDesc td;
    memset(&td, 0, sizeof(td));
    td.normalizedCoords = 0;
    td.addressMode[0] = cudaAddressModeClamp;
    td.readMode = cudaReadModeElementType;


    struct cudaResourceDesc resDesc;
    memset(&resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeLinear;
    resDesc.res.linear.devPtr = d_data;
    resDesc.res.linear.sizeInBytes = numel*sizeof(double);
    resDesc.res.linear.desc.f = cudaChannelFormatKindUnsigned;
    resDesc.res.linear.desc.x = 32;
    resDesc.res.linear.desc.y = 32;

    cudaTextureObject_t texObject;
    gpuErrchk(cudaCreateTextureObject(&texObject, &resDesc, &td, NULL));

    my_print<<<1,1>>>(texObject);

    gpuErrchk(cudaDeviceSynchronize());
    return 0;
}

即将 channel 描述修改为 64 位，从纹理对象读取 uint2，然后将其转换为 double，它应该按照您的要求工作。