如何在结构中嵌入 CUDA 纹理对象？答案

【问题标题】：How to embed CUDA Texture Objects in structs?如何在结构中嵌入 CUDA 纹理对象？
【发布时间】：2016-03-03 10:14:11
【问题描述】：

我们已成功使用以下帖子来帮助创建包含基本类型（如 int *）的结构。纹理为只读数组提供了很好的性能提升。我们使用了很多，这使得内核和内核子函数的参数列表又长又复杂。我们希望将纹理嵌入结构中以减少参数长度和复杂性。

Copying a struct containing pointers to CUDA device

这是一个代表我们使用的代码方法的 sn-p。它编译，但在运行时崩溃。

// Initialize texture description
memset(&textureDescription, 0, sizeof(textureDescription));
textureDescription.readMode = cudaReadModeElementType;

// Create Texture from variable
cudaTextureObject_t texture = 0;
cudaResourceDesc resource;
memset(&resource, 0, sizeof(resource));
resource.resType = cudaResourceTypeLinear;
resource.res.linear.devPtr = intArray;
resource.res.linear.desc.f = cudaChannelFormatKindSigned;
resource.res.linear.desc.x = 32; // bits per channel
resource.res.linear.sizeInBytes = count*sizeof(int);
cudaCreateTextureObject(&texture, resource, &textureDescription, NULL);

// These declarations are in the .h file
typedef struct SampleStructure {
   cudaTextureObject_t texture;
} SampleStructure;
SampleStructure *structureHost;
SampleStructure *structureDevice;

// Create host and device structures
structureHost = (SampleStructure *)malloc(sizeof(SampleStructure));
cudaMalloc(&structureDevice, sizeof(SampleStructure));

// Assign the texture object to the host structure
structureHost->texture = texture;

// Copy the host structure to Global Memory
cudaMemcpy(structureDevice, structureHost, sizeof(SampleStructure), cudaMemcpyHostToDevice));

// Pass Texture and Texture-embedded-in-structure to kernel
kenerl<<<1,1>>>(texture, structureDevice);

...
__global__ void
kernel(cudaTextureObject_t texture, SampleStructure *structureDevice) {
    value = tex1Dfetch<int>(texture, index); // Runs successfully at runtime
    value = tex1Dfetch<int>(structureDevice->texture, index); // Crashes at runtime
}

当在内核代码（或子函数）中使用“纹理”变量时，它正确运行。改用“structureDevice->texture”时，它会在运行时崩溃。

谁能展示一个简单的代码来展示如何成功地将纹理对象嵌入到传递给内核并在不崩溃的情况下运行的结构中？或者有人可以指出我们提供的代码中的错误可能在哪里？

【问题讨论】：

为什么不直接按值而不是按引用传递结构？
@talonmies 仅通过进行内核调用，按值传递就会导致内核中的运行时错误。我修改了参数列表以按值接受结构，但是以这种方式传递结构会导致运行时失败。内核>(texture, *structureDevice);
这不是按值传递结构。那是取消引用主机上的设备指针，这显然是非法的。只需在主机内存中创建一个结构并按值传递即可。支持的架构可以使用大小最大为 4kb 的参数列表，因此在使用按值传递时没有实际的大小限制
@talonmies 按价值传递有效！谢谢！
请随意添加您所做的作为答案。稍后您将能够接受答案，这会将问题从未回答列表中删除，并使其在搜索中更加可见（我也会投票赞成它......）

标签： c++ c struct cuda

【解决方案1】：

按值传递结构得到了一个可行的解决方案。这是使它工作的等效代码。感谢@talonmies 的建议。

虽然结构可以简化参数列表，但它会减慢执行速度，因为系统必须对全局内存进行 2 次调用，而不是 1:1 调用来获取结构和 1 次调用来获取纹理。为了提高性能，可以将结构复制到共享内存。在共享内存中使用该结构可以提高性能。

// Create the Texture Object
cudaResourceDesc resource;
memset(&resource, 0, sizeof(resource));
resource.resType = cudaResourceTypeLinear;
resource.res.linear.devPtr = intArray;
resource.res.linear.desc.f = cudaChannelFormatKindSigned;
resource.res.linear.desc.x = 32; // bits per channel
resource.res.linear.sizeInBytes = count*sizeof(int);
cudaCreateTextureObject(&texture, resource, &textureDescription, NULL);

// These structure declarations are in the .h file
typedef struct SampleStructure {
   cudaTextureObject_t texture;
} SampleStructure;
SampleStructure structureHost;

// Assign the texture object to the host structure
structureHost.texture = texture;

// Pass Texture and Texture-object-embedded-in-structure to kernel
kenerl<<<1,1>>>(texture, structureHost);

...
__global__ void
kernel(cudaTextureObject_t texture, SampleStructure structureDevice) {
    __shared__ SampleStructure structureSharedMemory;

    // Copy the structure to shared memory for faster access
    if (threadIdx.x == 0)
       structureSharedMemory = structureDevice;
    __threadfence_block();

    value = tex1Dfetch<int>(texture, index); // Runs successfully at runtime
    value = tex1Dfetch<int>(structureSharedMemory.texture, index); // Runs successfully at runtime
}

【讨论】：