cuda统一内存和指针别名

Posted 2023-03-23

【中文标题】cuda统一内存和指针别名

【英文标题】：cuda unified memory and pointer aliasing

【发布时间】：2017-04-30 01:14:43

【问题描述】：

我正在用 cuda 提神醒脑，特别是统一内存（我最后一次真正的 cuda 开发是在 3 年前），我有点生疏了。

pb：

我正在使用统一内存从容器创建任务。但是，经过几天的调查，我遇到了崩溃， 我不能说崩溃在哪里（复制构造函数），但不能说为什么。因为所有的指针都被正确分配了。

我与 Nvidia 的帖子无关 (https://devblogs.nvidia.com/parallelforall/unified-memory-in-cuda-6/) 关于C++和统一内存

#include <cuda.h>
#include <cstdio>

template<class T>
struct container
    container(int size = 1) cudaMallocManaged(&p,size*sizeof(T));
    ~container()cudaFree(p);
    __device__ __host__ T& operator[](int i) return p[i];
    T * p;
;

struct task
    int* a;
;

__global__ void kernel_gpu(task& t, container<task>& v) 
    printf(" gpu value task %i, should be 2 \n", *(t.a)); // this work
    task tmp(v[0]); // BUG
    printf(" gpu value task from vector %i, should be 1 \n", *(tmp.a));


void kernel_cpu(task& t, container<task>& v) 
    printf(" cpu value task %i, should be 2 \n", *(t.a)); // this work
    task tmp(v[0]);
    printf(" cpu value task from vector %i, should be 1 \n", *(tmp.a));


int main(int argc, const char * argv[]) 
    int* p1; 
    int* p2; 
    cudaMallocManaged(&p1,sizeof(int));
    cudaMallocManaged(&p2,sizeof(int));
    *p1 = 1;
    *p2 = 2;

    task t1,t2;
    t1.a=p1;
    t2.a=p2;

    container<task> c(2);

    c[0] = t1; 
    c[1] = t2; 

    //gpu does not work
    kernel_gpu<<<1,1>>>(c[1],c);
    cudaDeviceSynchronize();

    //cpu should work, no concurent access
    kernel_cpu(c[1],c);

    printf("job done !\n");

    cudaFree(p1);
    cudaFree(p2);

    return 0;

客观地说，我可以将一个对象作为参数传递给正确分配内存的地方。但是，看起来无法使用第二学位 间接的（这里是容器）

我犯了一个概念上的错误，但我看不出在哪里。

最好的，

提莫咖啡

我的机器：cuda 7.5、gcc 4.8.2、Tesla K20 m

【问题讨论】：

您的 task 类没有构造函数，无法使其在 GPU 上正常工作。

【参考方案1】：

虽然内存被分配为统一内存，但容器本身在主机代码中声明并分配在主机内存中：container<task> c(2);。您不能将它作为对设备代码的引用传递，并且在内核中取消引用它很可能会导致非法内存访问。

您可能希望使用cuda-memcheck 来识别此类问题。

【讨论】：

好的，但为什么它对任务有效（第一个参数）？任务对象在主机上分配，外部内存使用cuda统一，最糟糕的是我使用向量及其运算符[]

传递任务

您的 task t1 是 [en.cppreference.com/w/cpp/concept/…，所以没关系。但是，您的容器不是因为它具有自定义析构函数，因此它违反了 [en.cppreference.com/w/cpp/language/… destructor) 使您的容器不是 TriviallyCopyable 的要求。

正是因为 TriviallyCopyable，我们可以将指针（在主机代码中声明但填充在设备上分配的内存）传递给内核。

嗯，你认为如果我实现复制/分配/（也许是移动）构造函数，我应该有一些工作吗？

我不确定这一点，因为我从未尝试过。请随意尝试，让我知道结果。谢谢！