【问题标题】:Can I use vector for both device and host class in CUDA我可以在 CUDA 中对设备和主机类使用向量吗
【发布时间】:2013-11-13 14:57:57
【问题描述】:

我正在编写一个 c++ cuda 程序。我有一个非常简单的结构:

struct A
{
int size;
float* tab; 
}

还有一个内核:

__global__ void Kernel(A* res, int n,args*) //
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < n)
{
    res[i] = AGenerator::Generate(args[i]);
}
}

其中 AGenerator::Generate 创建 A 对象并填充选项卡数组。这里发生的情况是,当结果发送到主机时,选项卡指针无效。为了防止这种情况,我需要将 三人法则 应用于此类。由于会有很多这样的类,我想避免编写太多额外的代码。

我进行了研究,发现有一个带有 device_vector 和 host_vector 的推力库,这可能有助于解决我的问题,但问题是我希望 struct A 和类似的结构 可以从主机和主机调用device 所以 device 和 host_vector 不适合这个目的。有什么结构可以用来解决这个问题吗?

编辑 我发现按值传递结构会对我有所帮助,但由于性能非常重要,它似乎不是一个好的解决方案。

【问题讨论】:

  • 您打算如何将结果发送给主机?这是回答您的问题的一个非常重要的因素。我不认为这是三个问题的规则,因为在设备和主机之间共享数据时,具有嵌入式指针的结构本质上更难以在 CUDA 中处理,因为内存空间是分开的。当然,您的结构 definition 可用于设备和主机代码。但是在设备代码中使用该定义创建的对象需要几个步骤才能传输到主机。您将需要一些主机分配的空间,例如cudaMalloc.
  • 现在我有最简单的解决方案。 cudaMalloc((void**)&d_Res, Num * sizeof(CashflowsStruct));内核(d_Res,args)。 cudaMemcpy(h_Res, d_Res, Num* sizeof(CashflowsStruct), cudaMemcpyDeviceToHost)) 但我愿意接受任何建议。关于规则三:你很可能是对的。我试图实现分配运算符,但它仍然不起作用,所以这比我想象的要复杂
  • @RobertCrovella:如何使用主机分配的空间?
  • 一个对象不能由设备代码创建,并且还具有主机可访问的数据区域,除非您使用自定义分配器,该分配器在主机分配的预先分配的设备内存池中工作,例如与cudaMalloc。这很复杂。那是你要的吗?换一种说法,您应该首先定义您希望与A 支持的所有基本交互:在设备代码中创建A 的对象?在主机代码中?在主机和设备之间复制对象?如果您对所有这些回答“是”,那么我能想到的唯一方法就是使用自定义分配器。其他人可能有其他想法。
  • 这可能对所有这些都是肯定的。通过自定义分配器,您的意思是这样的:stackoverflow.com/questions/299761/… ?

标签: c++ vector cuda


【解决方案1】:

这是我对自定义分配器和池的想法的粗略概述,它将隐藏在主机和设备上使用类的一些机制。

我不认为它是卓越编程的典范。这只是我认为可能涉及的步骤的粗略概述。我敢肯定有很多错误。我没有包含它,但我认为你会想要一个公共方法,它也会获得 size。

#include <iostream>
#include <assert.h>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

typedef float mytype;

__device__ unsigned int pool_allocated = 0;
__device__ unsigned int pool_size = 0;
__device__ mytype *pool = 0;

__device__ unsigned int pool_reserve(size_t size){
  assert((pool_allocated+size) < pool_size);
  unsigned int offset = atomicAdd(&pool_allocated, size);
  assert (offset < pool_size);
  return offset;
}

__host__ void init_pool(size_t psize){
  mytype *temp;
  unsigned int my_size = psize;
  cudaMalloc((void **)&temp, psize*sizeof(mytype));
  cudaCheckErrors("init pool cudaMalloc fail");
  cudaMemcpyToSymbol(pool, &temp, sizeof(mytype *));
  cudaCheckErrors("init pool cudaMemcpyToSymbol 1 fail");
  cudaMemcpyToSymbol(pool_size, &my_size, sizeof(unsigned int));
  cudaCheckErrors("init pool cudaMemcpyToSymbol 2 fail");
}


class A{
  public:
  mytype *data;
  __host__ __device__ void pool_allocate_and_copy() {
  assert(d_data == 0);
  assert(size != 0);
#ifdef __CUDA_ARCH__
  unsigned int offset = pool_reserve(size);
  d_data = pool + offset;
  memcpy(d_data, data, size*sizeof(mytype));
#else
  cudaMalloc((void **)&d_data, size*sizeof(mytype));
  cudaCheckErrors("pool_allocate_and_copy cudaMalloc fail");
  cudaMemcpy(d_data, data, size*sizeof(mytype), cudaMemcpyHostToDevice);
  cudaCheckErrors("pool_allocate_and_copy cudaMemcpy fail");
#endif /* __CUDA_ARCH__ */

  }
  __host__ __device__ void update(){
#ifdef __CUDA_ARCH__
  assert(data != 0);
  data = d_data;
  assert(data != 0);
#else
  if (h_data == 0) h_data = (mytype *)malloc(size*sizeof(mytype));
  data = h_data;
  assert(data != 0);
  cudaMemcpy(data, d_data, size*sizeof(mytype), cudaMemcpyDeviceToHost);
  cudaCheckErrors("update cudaMempcy fail");
#endif
  }
  __host__ __device__ void allocate(size_t asize) {
    assert(data == 0);
    data = (mytype *)malloc(asize*sizeof(mytype));
    assert(data != 0);
#ifndef __CUDA_ARCH__
    h_data = data;
#endif
    size = asize;
  }
  __host__ __device__ void copyobj(A *obj){
    assert(obj != 0);
#ifdef __CUDA_ARCH__
    memcpy(this, obj, sizeof(A));
#else
    cudaMemcpy(this, obj, sizeof(A), cudaMemcpyDefault);
    cudaCheckErrors("copy cudaMempcy fail");
#endif
    this->update();
  }
  __host__ __device__ A();
    private:
    unsigned int size;
    mytype *d_data;
    mytype *h_data;
};

__host__ __device__ A::A(){
  data = 0;
  d_data = 0;
  h_data = 0;
  size = 0;
}

__global__ void mykernel(A obj, A *res){
  A mylocal;
  mylocal.copyobj(&obj);
  A mylocal2;
  mylocal2.allocate(24);
  mylocal2.data[0]=45;
  mylocal2.pool_allocate_and_copy();
  res->copyobj(&mylocal2);
  printf("kernel data %f\n", mylocal.data[0]);
}




int main(){
  A my_obj;
  A *d_result, h_result;
  my_obj.allocate(32);
  my_obj.data[0] = 12;
  init_pool(1048576);
  my_obj.pool_allocate_and_copy();
  cudaMalloc((void **)&d_result, sizeof(A));
  cudaCheckErrors("main cudaMalloc fail");
  mykernel<<<1,1>>>(my_obj, d_result);
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel fail");
  h_result.copyobj(d_result);
  printf("host data %f\n", h_result.data[0]);

  return 0;
}

【讨论】:

    【解决方案2】:

    我很确定问题的方向和相关的 cmets 是命运多舛。设备内存和主机内存在概念上和物理上都是完全不同的东西。指针只是不会结转!

    请返回第 1 步,并通过阅读reference manual 和progamming guide 了解有关在主机和设备之间复制值的更多详细信息。

    为了更准确地回答您的问题,请说明这些A 结构是如何在设备上分配的包括这些tab 浮点数的分配。还请展示AGenerator::Generate 如何以某种有意义的方式操纵那些tabs。我最好的选择是您在这里使用未分配的设备内存,并且您可能应该使用预先分配的浮点数组并在数组中插入而不是设备指针。然后这些索引将优雅地转移到主机。

    【讨论】:

    • 你的赌注是正确的,AGenerator::Generate 创建了一个结构,并为标签分配内存(使用 operator new),其大小是从它接收的函数参数中获得的。我认为这就是我与罗伯特·克罗维拉的观点。
    • @rank1 在__device__ 上运行的代码没有operator new。实际上根本没有动态内存分配。您必须预先分配所有内容。请参阅第 1 步!
    • @LumpN:自 2010 年发布 Fermi GPU 以来,CUDA 在 GPU 代码(新的或 malloc)中支持动态内存分配。
    猜你喜欢
    • 1970-01-01
    • 2011-03-19
    • 1970-01-01
    • 2020-08-20
    • 2020-03-30
    • 2010-12-16
    • 2012-03-16
    • 1970-01-01
    • 2020-05-07
    相关资源
    最近更新 更多