CUDA Texture Cache似乎有错误的数据？

Question

我编写 CUDA 代码已经有一段时间了，但我现在才开始了解如何使用纹理缓存。

使用 Nvidia SDK 中的 simpleTexture 示例作为灵感，我编写了一个使用纹理缓存的简单示例。主机将 Lena 图像复制到 GPU 并将其绑定为纹理。内核只是将纹理缓存的内容复制到一个输出数组中。

奇怪的是，结果（参见代码下方的全灰色图像）与输入不匹配。 对可能出现的问题有任何想法吗？

代码（看texCache_dummyKernel）：

texture<float, 2, cudaReadModeElementType> tex; //declare texture reference for 2D float texture

//note: tex is global, so no input ptr is needed
__global__ void texCache_dummyKernel(float* out, const int width, const int height){ //copy tex to output
    int x = blockIdx.x*blockDim.x + threadIdx.x; //my index into "big image"
    int y = blockIdx.y*blockDim.y + threadIdx.y;
    int idx = y*width+x;

    if(x < width && y < height)
        out[idx] = tex2D(tex, y, x);
}

int main(int argc, char **argv){        
    cv::Mat img = getRawImage("./Lena.pgm");
    img.convertTo(img, CV_32FC1);
    float* hostImg = (float*)&img.data[0];
    int width = img.cols; int height = img.rows;

    dim3 grid;  dim3 block;
    block.x = 16;  block.y = 16;
    grid.x = width/block.x + 1;          
    grid.y = height/block.y + 1;

    cudaArray *dImg; //cudaArray*, not float*
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);        
    CHECK_CUDART(cudaMallocArray(&dImg, &channelDesc, width, height));
    CHECK_CUDART(cudaMemcpyToArray(dImg, 0, 0, hostImg, width*height*sizeof(float), cudaMemcpyHostToDevice));
    setTexCacheParams(); //defined below
    CHECK_CUDART(cudaBindTextureToArray(tex, dImg, channelDesc)); //Bind the array to the texture

    float* dResult; //device memory for output
    CHECK_CUDART(cudaMalloc((void**)&dResult, sizeof(float)*width*height));

    texCache_dummyKernel<<<grid, block>>>(dResult, width, height); //dImg isn't an input param, since 'tex' is a global variable
    CHECK_CUDART(cudaGetLastError()); //make sure kernel didn't crash

    float* hostResult = (float*)malloc(sizeof(float)*width*height);
    CHECK_CUDART(cudaMemcpy(hostResult, dResult, sizeof(float)*width*height, cudaMemcpyDeviceToHost));
    outputProcessedImage(hostResult, width, height, "result.png"); //defined below
}

我可能应该提供一些我在上面使用过的辅助函数：

void setTexCacheParams(){ //configuration directly pulled from simpleTexture in nvidia sdk
    tex.addressMode[0] = cudaAddressModeWrap;
    tex.addressMode[1] = cudaAddressModeWrap;
    tex.filterMode = cudaFilterModeLinear;
    tex.normalized = true;    // access with normalized texture coordinates
}

void outputProcessedImage(float* processedImg, int width, int height, string out_filename){
    cv::Mat img = cv::Mat::zeros(height, width, CV_32FC1);
    for(int i=0; i<height; i++)
        for(int j=0; j<width; j++)
            img.at<float>(i,j) = processedImg[i*width + j]; //just grab the 1st of the 4 pixel spaces in a uchar4

    img.convertTo(img, CV_8UC1); //float to uchar
    vector<int> compression_params;
    compression_params.push_back(CV_IMWRITE_PNG_COMPRESSION);
    compression_params.push_back(9);
    cv::imwrite(out_filename, img, compression_params);
}

输入：

输出：

---

• 对不起，这篇文章是这样一堵代码墙！如果有任何关于使这样的代码更简洁的建议，我将不胜感激。
• 我在上面的文件 I/O 中使用了 OpenCV...希望这不会造成混淆。
• 当我更改内核以从一维 float* 数组中读取输入图像，并且我几乎保持其他所有内容相同时，我得到了正确的结果。

Answer 1

在您的原始代码中，您已将纹理初始化为使用标准化坐标。这意味着纹理在每个空间维度的 [0,1] 上进行寻址。所以你的内核应该是这样的：

__global__ 
void texCache_dummyKernel(float* out, const int width, const int height)
{
    int x = blockIdx.x*blockDim.x + threadIdx.x; //my index into "big image"
    int y = blockIdx.y*blockDim.y + threadIdx.y;
    int idx = y*width+x;

    if(x < width && y < height) {
        float u = float(x)/float(width), v = float(y)/float(height);
        out[idx] = tex2D(tex, u, v);
    }
}

[标准免责声明：在浏览器中编写，未经编译或测试，使用风险自负]

即。您应该将坐标传递给tex2D，通过除以图像宽度和高度进行归一化。

或者，正如您所发现的，您可以将纹理定义更改为normalized=false，并在绝对而不是相对纹理坐标中使用寻址。即使这样，在您的代码中读取的纹理也应该如下所示：

out[idx] = tex2D(tex, float(x)+0.5f, float(y)+0.5f);

因为纹理寻址总是使用浮点坐标完成，并且纹理数据以体素为中心，因此每个坐标添加 0.5 以确保读取来自每个插值区域或体积内的质心纹理。

您可以在 CUDA C 编程指南的一个附录中找到关于纹理过滤和寻址模式及其对插值的影响的描述。