CUDA内核中的内存分配

Question

我有一个内核的以下(片段).

__global__ void plain(int* geneVec, float* probs, int* nComponents, float* randomNumbers,int *nGenes)
{

    int xid = threadIdx.x + (blockDim.x * blockIdx.x);

    float* currentProbs= (float*)malloc(sizeof(float)*tmp);

         .....
         .....

    currentProbs[0] = probs[start];
    for (k=1;k<nComponents[0]; k++)
    {
        currentProbs[k] = currentProbs[k-1] + prob;
    }

       ...
       ...
      free(currentProbs);

}

当它是静态的(即使是相同的大小)时它非常快,但是当动态分配CurrentProbs时(如上所述),性能很糟糕.

这个问题说我可以在内核中执行此操作:CUDA在__device__函数中分配内存

这是一个相关的问题:CUDA中Malloc功能的效率

我想知道除了论文中提出的方法之外,是否还有其他任何方法解决了这个问题？如果没有这种惩罚,在内核中不能malloc/free是一件很荒谬的事.

Answer 1

我认为引入malloc()减慢代码速度的原因是它在全局内存中分配内存.当您使用固定大小的数组时,编译器可能会将它放在寄存器文件中,这要快得多.

必须在内核中执行malloc可能意味着您尝试使用单个内核执行太多工作.如果每个线程分配不同的内存量,那么每个线程在for循环中运行的次数不同,并且会产生大量的warp散度.

如果warp中的每个线程循环次数相同,则只需预先分配.即使它们运行的​​次数不同,也可以使用恒定的大小.但相反,我认为您应该看看如何重构代码以完全从内核中删除该循环.