CUDA上的2D数组

Question

我想动态分配全局2D数组CUDA.我怎样才能做到这一点？

在我的主要部分,我Kernel在循环中呼唤我.但在我调用内核之前,我需要在内核上分配一些内存GPU.在内核调用之后,从GPU向CPU发送单个整数以通知问题是否已解决.
如果问题没有解决,我将不会释放旧内存,因为还需要它,我应该GPU为内核分配新内存并再次调用内核.

显示一个sudocode:

int n=0,i=0;
while(n==0)
{
    //allocate 2d memory for MEM[i++] 
    //call kernel(MEM,i)
    // get n from kernel       
}


__global__ void kernerl(Mem,int i)
{
    Mem[0][5]=1;
    Mem[1][0]=Mem[0][5]+23;//can use this when MEM[1] is allocated before kernel call
}

有什么建议？谢谢.

Answer 1

两个开放注释 - 使用动态分配的2D数组在CUDA中是一个坏主意,并且在循环中进行重复的内存分配也不是一个好主意.两者都会产生不必要的性能损失.

对于主机代码,如下所示:

size_t allocsize = 16000 * sizeof(float);
int n_allocations = 16;
float * dpointer
cudaMalloc((void **)&dpointer, n_allocations * size_t(allocsize));

float * dcurrent = dpointer;
int n = 0;
for(int i=0; ((n==0) && (i<n_allocations)); i++, dcurrent+=allocsize) {

    // whatever you do before the kernel

    kernel <<< gridsize,blocksize >>> (dcurrent,.....);

    // whatever you do after the kernel

}

是优选的.在这里,您只需调用一次cudaMalloc,并将偏移量传递给分配,这样可以在循环内部释放内存分配和管理.循环结构也意味着你无法无休止地运行并耗尽所有GPU内存.

在二维阵列问题本身,有两个原因,这是一个坏主意.首先,分配要求具有N行的2D阵列需要(N + 1)cudaMalloc调用和主机设备存储器副本,这是缓慢且丑陋的.其次在内核代码中,为了获取数据,GPU必须执行两次全局内存读取,一次用于指针间接获取行地址,然后一次从行中获取数据.这比这个替代方案要慢得多:

#define idx(i,j,lda) ( (j) + ((i)*(lda)) )
__global__ void kernerl(float * Mem, int lda, ....)
{
    Mem[idx(0,5,lda)]=1; // MemMem[0][5]=1;
}

它使用索引到1D分配.在GPU内存中,交易非常昂贵,但FLOPS和IOPS都很便宜.单个整数乘法 - 加法是执行此操作的最有效方法.如果需要访问先前内核调用的结果,只需将偏移量传递给先前的结果并在内核中使用两个指针,如下所示:

__global__ void kernel(float *Mem, int lda, int this, int previous)
{
   float * Mem0 = Mem + this;
   float * Mem1 = Mem + previous;

}

高效的分布式内存程序(以及CUDA实际上是一种分布式内存编程)在一段时间后开始看起来像Fortran,但这是您为便携性,透明度和效率付出的代价.

希望这有帮助.