CUDA内核中类成员赋值为何出现不一致问题？

在CUDA环境中定义了FP_union5类，包含rows和cols两个成员变量，实现了带两个int参数的构造函数用于初始化成员。编写内核函数调用返回该类对象的函数时，通过NSight调试器观察到异常现象：

• 构造函数内部，this指针指向的rows和cols被正确设置为1；
• 构造函数返回后，函数内要返回的变量（myfp3）的rows和cols显示为0；
• 回到内核函数中，接收返回值的变量（myfp5）的rows和cols又变回1。

该异常仅出现在要返回的变量上，函数内其他同类型变量（如myfp2）的成员赋值始终正常。

示例代码 #

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#define mydevice __device__

class FP_union5
{
public:
    // 成员变量
    int rows;
    int cols;

    // 析构函数
    mydevice ~FP_union5();

    // 默认构造函数
    mydevice inline FP_union5(void) : rows(0), cols(0)
    {
    }

    // 参数化构造函数
    mydevice FP_union5(int r, int c);

    // 拷贝构造函数
    mydevice FP_union5(const FP_union5& other); 
    // 赋值运算符
    mydevice FP_union5& operator=(const FP_union5& other); 
};

mydevice FP_union5::~FP_union5() {}

mydevice FP_union5::FP_union5(int r, int c)
{
    rows = r;
    cols = c;

    if (rows * cols <= 0)
    {
        rows = cols = 0;
        return;
    }
}

mydevice FP_union5::FP_union5(const FP_union5& other) {
    rows = other.rows;
    cols = other.cols;
}

mydevice FP_union5& FP_union5::operator=(const FP_union5& other) {
    if (this != &other) {
        rows = other.rows;
        cols = other.cols;
    }
    return *this;
}

mydevice FP_union5 callfpu5(int i)
{
    FP_union5 myfp2(i, i);
    FP_union5 myfp3(i, i);
 
    return(myfp3);
}

__global__ void addKernel(int j)
{
    int i = threadIdx.x;
    FP_union5 myfp5 = callfpu5(1);
}

int main()
{
    const int arraySize = 5;
    addKernel <<<1, arraySize>>> (arraySize);
}

问题原因分析 #

这种现象几乎可以确定是NSight调试器的可视化bug，或是编译器返回值优化（RVO）导致的调试视图异常：

1. 代码逻辑完全合规：拷贝构造函数和赋值操作符都正确实现了成员变量的复制，最终内核中myfp5的值正确，说明程序运行逻辑没有问题；
2. 返回值优化（RVO）是C++编译器的常用优化，会直接在调用方的内存空间构造返回对象，跳过中间临时变量的拷贝，这可能导致调试器无法正确追踪中间变量的状态，出现显示错误；
3. 仅返回变量出现异常，是因为RVO仅作用于返回值相关的对象，其他局部变量不受影响。

解决方案 #

1. 以最终运行结果为准：既然内核中接收的变量值正确，说明程序逻辑没问题，可忽略调试器中间步骤的异常显示；
2. 调整编译选项（可选）：尝试禁用返回值优化（如添加-fno-elide-constructors编译选项，需确认CUDA编译器支持），但这会降低性能，仅用于调试验证；
3. 更新调试器：升级NSight到最新版本，这类调试可视化bug通常会在新版本中修复。

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