CUDA处理CV_16UC1格式OpenCV图像拷贝输出全黑问题求助

基于CUDA实现CV_16UC1格式N张图像平均值计算时，输出始终为全黑图像。简化逻辑为单张输入图像经CUDA核函数直接拷贝至输出、再回传主机端后，全黑问题仍然复现。
测试代码如下：

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <opencv2/opencv.hpp>
#include <opencv2/core/cuda/common.hpp>

using namespace std;

__global__ void vectorAdd10(unsigned short* input, unsigned short* output, int width, int height, int inStep, int outStep) {
    // Index of current thread
    const int x = blockIdx.x * blockDim.x + threadIdx.x;
    const int y = blockIdx.y * blockDim.y + threadIdx.y;

    // Number of channels
    const int in_c = inStep / width;
    const int out_c = outStep / width;

    // Only valid threads perform memory I/O
    if ((x < width) && (y < height)) {

        // Location of pixel
        const int in_tid = y * inStep + (in_c * x);
        const int out_tid = y * outStep + (out_c * x);

        // Invert
        for (int i = 0; i < in_c; ++i) {
            output[out_tid + i] = static_cast<unsigned short>(input[in_tid + i]);
        }
    }
}

int main() { //Umesh: replace this function with other name

    cv::Mat& result = cv::Mat(1024, 1024, CV_16UC1);;
    cv::Mat input = cv::imread("Frame1.png");
    const int inBytes = input.step * input.rows;
    const int outBytes = result.step * result.rows;
    
    unsigned short* b_input, * b_output;
    std::cout << "step 1" << std::endl;
    b_input = (unsigned short*)malloc(inBytes);
    b_output = (unsigned short*)malloc(outBytes);
    cout << "step 21" << endl;
    cudaMalloc<unsigned short>(&b_input, inBytes);
    cudaMalloc<unsigned short>(&b_output, outBytes);
    cout << "step 2" << endl;

    cudaMemcpy(b_input, input.ptr(), inBytes, cudaMemcpyHostToDevice);
    cout << "step 3" << endl;

    // Threadblock size
    const dim3 block(16, 16);

    const dim3 grid(cv::cuda::device::divUp(input.cols, block.x), cv::cuda::device::divUp(input.rows, block.y));

    // Grid size

    vectorAdd10 << <grid, block >> > (b_input, b_output, input.cols, input.rows, input.step, result.step);
    cout << "step 4" << endl;

    cudaSafeCall(cudaDeviceSynchronize());
    // Copy sum vector from device to host
    cudaSafeCall(cudaMemcpy(result.ptr(), b_output, outBytes, cudaMemcpyDeviceToHost));
    cout << "step 5" << endl;

    cv::imshow("Result", result);
    cv::waitKey(3000);
    cout << "step 6" << endl;


    printf("COMPLETED SUCCESFULLY\n");
}

预期单张图直接拷贝后输出与输入完全一致，实际输出为全黑图像。

代码存在4处核心错误，任意一处都可能导致全黑输出：

• 图像读取参数错误：cv::imread默认以cv::IMREAD_COLOR模式加载图像，强制转成8位3通道CV_8UC3格式，与预期的16位单通道CV_16UC1格式完全不匹配，输入数据的位深、通道数从读取阶段就错了。
• 无效的主机内存分配：提前用malloc给设备内存指针分配主机堆内存，后续cudaMalloc会直接覆盖指针值造成主机内存泄漏，且代码全程未检查cudaMalloc、cudaMemcpy、核函数启动的返回状态，CUDA接口调用失败时无任何提示，极易出现数据未正常拷贝、核函数未正常执行的情况。
• Mat对象引用绑定非法：cv::Mat& result = cv::Mat(1024, 1024, CV_16UC1)将非常量引用绑定到临时对象，属于C++未定义行为，多数编译环境下临时Mat对象在该行执行完就会析构，result成为野引用，后续向result.ptr拷贝数据属于野内存写入，显示时无法读取到正确像素值。
• 核函数寻址逻辑完全错误：传入核函数的inStep/outStep是OpenCV定义的每行字节步长，但核函数内用unsigned short*类型指针做寻址（C++指针偏移单位为指向类型的长度，此处为2字节），直接用字节步长计算元素偏移会把每行起始地址算到实际位置的2倍内存外，访问的全是值为0的未初始化显存；另外自行用inStep/width计算通道数不可靠，OpenCV的step可能包含行尾对齐填充字节，计算结果不准。

1. 图像读取时传入cv::IMREAD_UNCHANGED参数，保留原图像的位深和通道数，加载后必须校验图像是否为空、格式和尺寸是否匹配预期。
2. 删除多余的malloc调用，所有CUDA接口调用后都要检查错误码，核函数启动后立刻调用cudaGetLastError检查启动状态，设备内存使用完后调用cudaFree释放。
3. 去掉result定义的引用符号，直接实例化Mat对象，避免绑定临时对象的未定义行为。
4. 核函数内不要用字节步长直接计算元素偏移，提前把字节步长除以sizeof(unsigned short)转成元素步长，同时直接传入Mat的通道数，不要自行通过step计算。

修正后的核心代码：

// 核函数修正：传入元素步长、通道数，不要用字节步长直接寻址
__global__ void vectorAdd10(unsigned short* input, unsigned short* output, int width, int height, int inStepElems, int outStepElems, int channels) {
    const int x = blockIdx.x * blockDim.x + threadIdx.x;
    const int y = blockIdx.y * blockDim.y + threadIdx.y;

    if ((x < width) && (y < height)) {
        const int in_tid = y * inStepElems + (channels * x);
        const int out_tid = y * outStepElems + (channels * x);
        for (int i = 0; i < channels; ++i) {
            output[out_tid + i] = input[in_tid + i];
        }
    }
}

int main() {
    // 修正：去掉引用，直接实例化Mat
    cv::Mat result(1024, 1024, CV_16UC1);
    // 修正：传入IMREAD_UNCHANGED保留原图格式
    cv::Mat input = cv::imread("Frame1.png", cv::IMREAD_UNCHANGED);
    // 输入合法性校验
    if (input.empty() || input.type() != CV_16UC1) {
        printf("Input load failed or type mismatch, require CV_16UC1\n");
        return -1;
    }
    if (input.cols != 1024 || input.rows != 1024) {
        printf("Input size mismatch, require 1024x1024\n");
        return -1;
    }

    const int inBytes = input.step * input.rows;
    const int outBytes = result.step * result.rows;
    
    unsigned short* b_input, * b_output;
    cudaError_t cudaStatus;
    // 修正：删除多余malloc，直接分配设备内存并检查错误
    cudaStatus = cudaMalloc<unsigned short>(&b_input, inBytes);
    if (cudaStatus != cudaSuccess) {
        printf("cudaMalloc input failed: %s\n", cudaGetErrorString(cudaStatus));
        return -1;
    }
    cudaStatus = cudaMalloc<unsigned short>(&b_output, outBytes);
    if (cudaStatus != cudaSuccess) {
        printf("cudaMalloc output failed: %s\n", cudaGetErrorString(cudaStatus));
        cudaFree(b_input);
        return -1;
    }

    cudaStatus = cudaMemcpy(b_input, input.ptr(), inBytes, cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        printf("Memcpy H2D failed: %s\n", cudaGetErrorString(cudaStatus));
        cudaFree(b_input);
        cudaFree(b_output);
        return -1;
    }

    const dim3 block(16, 16);
    const dim3 grid(cv::cuda::device::divUp(input.cols, block.x), cv::cuda::device::divUp(input.rows, block.y));

    // 修正：传入元素步长、通道数，检查核函数启动错误
    vectorAdd10 << <grid, block >> > (
        b_input, b_output, input.cols, input.rows,
        input.step / sizeof(unsigned short),
        result.step / sizeof(unsigned short),
        input.channels()
    );
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        printf("Kernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        cudaFree(b_input);
        cudaFree(b_output);
        return -1;
    }

    cudaSafeCall(cudaDeviceSynchronize());
    cudaSafeCall(cudaMemcpy(result.ptr(), b_output, outBytes, cudaMemcpyDeviceToHost));

    // 释放设备内存
    cudaFree(b_input);
    cudaFree(b_output);

    cv::imshow("Result", result);
    cv::waitKey(3000);

    printf("COMPLETED SUCCESFULLY\n");
    return 0;
}

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