cudaMalloc与cudaMemcpy调用后__device__变量仍为nullptr求助

你的核心问题在于主机端错误地直接操作__device__变量的地址：

• __device__ int* var;是GPU全局内存中的符号变量，主机端无法直接通过&gpu::var获取它的GPU内存地址并传递给cudaMalloc。你当前的cudaMalloc调用，实际是把分配的设备内存地址写到了主机端的临时占位地址，而非GPU上的gpu::var变量本身。
• 后续的cudaMemcpy虽然没有报错，但因为gpu::var在GPU上仍是nullptr，这步拷贝的目标地址无效，只是CUDA异步执行模式下不会立刻触发错误，核函数运行时才会暴露gpu::var未被正确赋值的问题。

需要通过CUDA提供的符号操作函数完成主机端对__device__变量的赋值，具体步骤：

1. 调用cudaMalloc分配设备内存，得到主机可见的设备指针；
2. 使用cudaMemcpyToSymbol将该设备指针的值写入GPU上的__device__ int* var变量；
3. 用cudaMemcpy把主机数据拷贝到设备指针指向的内存；
4. 核函数即可正确访问gpu::var指向的内存。

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

#define cugo(ans) execute_cuda_runtime((ans), __FILE__, __LINE__)

namespace calcu {
    namespace cpu {
        int var; 
    }

    namespace gpu {
        __device__ int* var;
    }

    void execute_cuda_runtime(cudaError_t code, const char* fname, int line) {
        if (code != cudaSuccess) {
            printf("CUDA execution error at %s:%d\n", fname, line);
            printf("Error is %s\n", cudaGetErrorString(code));
            exit(886);
        }
    }

    void cuda_test_print(int useless);
}

namespace calcu {
    __global__ void cuda_test_print_kernel(int useless) {
        if (gpu::var == nullptr) {
            printf("wtf... gpu::var is nullptr\n");
        } else {
            printf("ftw... gpu::var is %p, value is %d\n", gpu::var, *gpu::var);
        }
    }

    void cuda_test_print(int useless) {
        cuda_test_print_kernel<<<1, 1>>>(useless);
        cugo(cudaDeviceSynchronize());
    }
}

int main() {
    using namespace calcu;

    cpu::var = 1;
    int* d_temp_ptr;

    // 1. 分配设备内存
    cugo(cudaMalloc((void**)&d_temp_ptr, sizeof(int)));
    // 2. 将设备指针写入GPU上的__device__变量
    cugo(cudaMemcpyToSymbol(gpu::var, &d_temp_ptr, sizeof(int*)));
    // 3. 拷贝主机数据到设备内存
    cugo(cudaMemcpy(d_temp_ptr, &cpu::var, sizeof(int), cudaMemcpyHostToDevice));

    cuda_test_print(0);

    // 释放设备内存
    cugo(cudaFree(d_temp_ptr));

    return 0;
}

• 新增临时指针d_temp_ptr存储cudaMalloc的结果；
• 用cudaMemcpyToSymbol替代直接赋值，把设备指针写入GPU的__device__变量；
• 给cudaDeviceSynchronize加上错误检查，避免漏掉异步执行的错误；
• 核函数新增*gpu::var的打印，验证数据是否正确拷贝。

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