CUDA 12.6读取1D浮点纹理内存异常：仅读取到0的问题排查

你的代码读取到全0的核心问题是cudaMemcpy3D的参数配置错误，另外缺少CUDA错误检查会导致你无法及时定位问题。以下是具体分析和修正方案：

问题分析 #

1. cudaMemcpy3D参数误用：
   你尝试用cudaMemcpy3D复制一维数据到cudaArray，但make_cudaPitchedPtr的参数设置不符合要求。主机端的内存是连续的一维数组，用cudaMemcpyToArray会更直接，避免pitch参数的配置错误。

2. 无错误检查：
   所有CUDA API调用（如cudaMalloc、cudaMemcpy3D）和核函数启动后都没有错误检查，这会导致你无法确认哪一步操作失败，直接表现为读取到初始的0值。

修正后的代码 #

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

#define CHECK_CUDA_ERROR(err) \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    }

// 从纹理内存读取的核函数
__global__ void readFromTexture(cudaTextureObject_t texObj, float *output, int size) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size) {
        output[idx] = tex1Dfetch<float>(texObj, idx); // 从纹理对象获取数据
    }
}

int main() {
    const int size = 16;
    float h_input[size], h_output[size];
    float *d_output;
    cudaError_t err;

    // 初始化主机端输入数组
    for (int i = 0; i < size; i++) {
        h_input[i] = static_cast<float>(i * 2);
    }

    // 为输出分配设备端内存
    err = cudaMalloc((void**)&d_output, size * sizeof(float));
    CHECK_CUDA_ERROR(err);

    // 在设备端内存中分配CUDA数组
    cudaArray *d_array;
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
    err = cudaMallocArray(&d_array, &channelDesc, size);
    CHECK_CUDA_ERROR(err);

    // 替换cudaMemcpy3D，用cudaMemcpyToArray直接复制一维数据
    err = cudaMemcpyToArray(d_array, 0, 0, h_input, size * sizeof(float), cudaMemcpyHostToDevice);
    CHECK_CUDA_ERROR(err);

    // 创建纹理对象
    cudaResourceDesc resDesc = {};
    resDesc.resType = cudaResourceTypeArray;
    resDesc.res.array.array = d_array;

    cudaTextureDesc texDesc = {};
    texDesc.addressMode[0] = cudaAddressModeClamp;
    texDesc.filterMode = cudaFilterModePoint;
    texDesc.readMode = cudaReadModeElementType;
    texDesc.normalizedCoords = 0;

    cudaTextureObject_t texObj = 0;
    err = cudaCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
    CHECK_CUDA_ERROR(err);

    // 启动核函数
    int threadsPerBlock = 8;
    int blocksPerGrid = (size + threadsPerBlock - 1) / threadsPerBlock;
    readFromTexture<<<blocksPerGrid, threadsPerBlock>>>(texObj, d_output, size);
    err = cudaGetLastError();
    CHECK_CUDA_ERROR(err);
    err = cudaDeviceSynchronize();
    CHECK_CUDA_ERROR(err);

    // 将结果复制回主机端
    err = cudaMemcpy(h_output, d_output, size * sizeof(float), cudaMemcpyDeviceToHost);
    CHECK_CUDA_ERROR(err);

    // 打印结果
    printf("Input:  \tOutput:\n");
    for (int i = 0; i < size; i++) {
        printf("%6.1f\t%6.1f\n", h_input[i], h_output[i]);
    }

    // 清理资源
    err = cudaDestroyTextureObject(texObj);
    CHECK_CUDA_ERROR(err);
    err = cudaFree(d_output);
    CHECK_CUDA_ERROR(err);
    err = cudaFreeArray(d_array);
    CHECK_CUDA_ERROR(err);

    return 0;
}

关键修正点 #

• 用cudaMemcpyToArray替代cudaMemcpy3D，简化一维数据到cudaArray的复制操作，避免pitch参数配置错误。
• 添加CHECK_CUDA_ERROR宏，对所有CUDA操作进行错误检查，方便定位问题。
• 核函数启动后添加cudaGetLastError和cudaDeviceSynchronize，检查核函数是否启动成功并等待执行完成。

内容的提问来源于stack exchange，提问作者Chopp