tex1Dfetch 意外返回 0

Question

我认为这与此处报告的问题不同：

Bound CUDA texture reads zero

CUDA 1D texture fetch always return 0

在我的 CUDA 应用程序中，我注意到 tex1Dfetch 没有返回预期值，超过了缓冲区中的某个索引。应用程序中的初步观察是可以正确读取索引 0 处的值，但在 12705625 处读取的值为 0。我制作了一个小测试程序来调查这个问题，如下所示。结果让我有点莫名其妙。我正在尝试探查哪些索引值不再被正确读取。但是随着值 arraySize 的更改，"firstBadIndex" 也会更改。即使 arraySize =2，第二个值也读错了！随着 arraySize 变大，firstBadIndex 变大。当绑定到 float、float2 或 float4 数组时会发生这种情况。如果数据是从设备缓冲区读取的（切换 FetchTextureData 中的注释行），那么一切都很好。这是在 Tesla c2075 上使用 CUDA 6.5。感谢您的任何见解或建议。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

#define FLOATTYPE float4
texture<FLOATTYPE,cudaTextureType1D,cudaReadModeElementType> texture1D;

const unsigned int arraySize = 1000;
FLOATTYPE* host;
FLOATTYPE* device;
FLOATTYPE* dTemp;
FLOATTYPE hTemp[1];

__global__ void FetchTextureData(FLOATTYPE* data,FLOATTYPE* arr,int idx)
{
    data[0] = tex1Dfetch(texture1D, idx);
    //data[0] = arr[idx];
}

bool GetTextureValues(int idx){

    FetchTextureData<<<1,1>>>(dTemp,device,idx);

    // copy to the host
    cudaError_t err = cudaMemcpy(hTemp,dTemp,sizeof(FLOATTYPE),cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        throw "cudaMemcpy failed!";
    }

    if (cudaDeviceSynchronize() != cudaSuccess) {
        throw "cudaDeviceSynchronize failed!";
    }

    return hTemp[0].x == 1.0f;
}

int main()
{

    try{

        host = new FLOATTYPE[arraySize];
        cudaError_t err = cudaMalloc((void**)&device,sizeof(FLOATTYPE) * arraySize);
        cudaError_t err1 = cudaMalloc((void**)&dTemp,sizeof(FLOATTYPE));
        if (err != cudaSuccess || err1 != cudaSuccess) {
            throw "cudaMalloc failed!";
        }

        // make some host data
        for(unsigned int i=0; i<arraySize; i++){
            FLOATTYPE data = {1.0f, 0.0f, 0.0f, 0.0f};
            host[i] = data;
        }

        // and copy it to the device
        err = cudaMemcpy(device,host,sizeof(FLOATTYPE) * arraySize,cudaMemcpyHostToDevice);
        if (err != cudaSuccess){
            throw "cudaMemcpy failed!";
        }

        // set up the textures
        cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<FLOATTYPE>();
        texture1D.addressMode[0] = cudaAddressModeClamp;
        texture1D.filterMode = cudaFilterModePoint;
        texture1D.normalized = false;
        cudaBindTexture(NULL, texture1D, device, channelDesc, arraySize);

        // do a texture fetch and find where the fetches stop working
        int lastGoodValue = -1, firstBadValue = -1;
        float4 badValue = {-1.0f,0.0f,0.0f,0.0f};

        for(unsigned int i=0; i<arraySize; i++){

            if(i % 100000 == 0) printf("%d\n",i);
            bool isGood = GetTextureValues(i);

            if(firstBadValue == -1 && !isGood)
                firstBadValue = i;

            if(isGood)
                lastGoodValue = i;
            else
                badValue = hTemp[0];
        }

        printf("lastGoodValue %d, firstBadValue %d\n",lastGoodValue,firstBadValue);
        printf("Bad value is (%.2f)\n",badValue.x);

    }catch(const char* err){
        printf("\nCaught an error : %s\n",err);
    }
    return 0;
}

Answer 1

问题出在纹理设置上。这个：

cudaBindTexture(NULL, texture1D, device, channelDesc, arraySize);

应该是：

    cudaBindTexture(NULL, texture1D, device, channelDesc, 
                                   arraySize * sizeof(FLOATTYPE));

根据documentation，size 参数是以字节为单位的内存区域的大小，而不是元素的数量。我原以为在钳位寻址模式下，代码仍能按预期工作。使用边界模式，您应该得到一个零值，这看起来会触发您的错误值检测。我实际上没有运行你的代码，所以也许我在某处遗漏了一个微妙的地方。对于这样一个简单的重现案例，您的代码结构相当复杂且难以理解（至少在我正在阅读的移动 phone 屏幕上）。

编辑补充说，在我开始写这篇文章和写完这段时间之间，@njuffa 在评论中指出了同样的错误

tex1Dfetch 意外返回 0

tex1Dfetch unexpectedly returning 0

c++

textures

cuda

gpu