cuda统一内存和指针别名

我用cuda刷新了我的思想，特别是统一内存（我最后一个真正的cuda dev是3年前），我有点生锈。cuda统一内存和指针别名

铅：

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

我不与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; 
}

客观我可以通过一个对象作为在存储器已被正确分配的参数。然而，它看起来像不可能使用间接性（这里是容器）的第二级

我在做一个概念错误，但我不知道在哪里。

最佳，

Timocafe

我的机器：CUDA 7.5，GCC 4.8.2，特斯拉K20米

来源

2016-12-14 Timocafé

你的'任务'类没有一个构造函数可以让它在GPU上正常工作。 – talonmies

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

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

来源

2016-12-14 16:35:46 yhf8377

好的，但为什么它的任务（第一个参数）？任务对象在主机上分配，而使用cuda的外部存储器统一，最差的是我使用向量及其运算符传递任务[] –

你的'任务t1'是[http://en.cppreference.com/w/cpp/concept/TriviallyCopyable]（TriviallyCopyable），所以很好。然而你的容器并不像它有一个自定义的析构函数，所以它违反了[http://en.cppreference.com/w/cpp/language/destructor#Trivial_destructor](Trivial destructor）使你的容器不是TriviallyCopyable的要求。 – yhf8377

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

cuda统一内存和指针别名

回答

相关问题