CUDA 内核中的内存分配

Question

我有以下 (sn-p) 的内核。

__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 allocate memory in __device__ function

这是一个相关的问题：Efficiency of Malloc function in CUDA

我想知道除了论文中提出的方法之外，是否还有其他方法可以解决这个问题？ 如果没有这种惩罚，就不能在内核中进行 malloc/free，这似乎很荒谬。

Answer 1

我认为引入 malloc() 会减慢代码速度的原因是它在全局内存中分配内存。当你使用一个固定大小的数组时，编译器很可能会把它放到寄存器文件中，这样会快很多。

必须在内核中执行 malloc 可能意味着您尝试使用单个内核做太多工作。如果每个线程分配不同数量的内存，那么每个线程在 for 循环中运行的次数也不同，你会得到很多 warp 发散。

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