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

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

我想知道是否有其他方法解决了这个问题，除了在论文中提出的方法之外？ 似乎荒谬的是不能在内核内malloc/free而不带有这种惩罚。

- RNs_Ghost

---

你的伪代码中的 tmp 是从哪里来的？ - talonmies

那么它是每个内核调用都是常量吗？如果是这样，为什么还要动态内存分配呢？ - talonmies

Answer 1

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

在内核中执行malloc可能意味着你试图在单个内核中做太多的工作。如果每个线程分配不同数量的内存，则每个线程在for循环中运行的次数也不同，这会导致大量的warp分歧。

如果warp中的每个线程运行相同次数的循环，那么就提前分配内存。即使它们运行次数不同，也可以使用常量大小。但我认为你应该考虑如何重构代码，以完全从内核中删除该循环。

- Roger Dahl

---

除非程序员使用 __shared__ 限定符将内核变量定义为共享内存，否则编译器永远不会将它们分配到共享内存中。只能使用寄存器或本地内存。 - talonmies