CUDA - 在内核中动态重新分配更多全局内存答案

【问题标题】：CUDA - dynamically reallocate more global memory in KernelCUDA - 在内核中动态重新分配更多全局内存
【发布时间】：2019-12-01 16:28:25
【问题描述】：

我对以下任务有疑问：

“给定一个二维数组“a[N][M]”，所以 N 行长度为 M。数组的每个元素都包含一个介于 0 和 16 之间的随机整数值。编写一个内核“compact(int *a, int *listM, int *listN)”，它只包含一个块的N个线程，每个线程计算数组的一行有多少元素值为 16。

线程将这些数字写入共享内存中长度为 N 的数组“num”，然后（在屏障之后）其中一个线程执行下面列出的前缀代码“PrefixSum(int *num, int N)”（在下面的代码中，我解释了这段代码的作用）。最后（再次障碍），每个线程“Idx”将其行中值为 16 的元素的 N 值和 M 值，分别位置（或“x 和 y 坐标”）写入两个数组中“全局内存中的 listM" 和 "listN"，从这些数组中的位置 "num[Idx]" 开始。为了更容易实现这最后一个任务，有上面提到的前缀代码。”

我已经编写了一个内核和一个合适的 main 来测试它。但是，我仍然有一个我无法解决的问题。

在两个数组“listeM”和“listeN”中，应该存储数组“a[M][N]”中出现的每16个的单独位置。因此，它们的大小必须等于 16 的出现总数，这可能会有所不同。

由于您不知道值为 16 的元素的确切数量，因此您只能在内核运行时知道两个数组“listeM”和“listeN”需要多少内存。当然，您可以在程序启动时为最大可能数量释放足够的内存，即 N 乘以 M，但这将非常低效。是否可以编写内核，以便每个线程在计算其行中值为 16 的元素数量（仅此数字）后动态扩大两个数组“listeM”和“listeN”？

这是我的内核：

__global__ void compact(int* a, int* listM, int* listN)
{
    int Idx = threadIdx.x;
    int elements, i;

    i = elements = 0;

    __shared__ int num[N];

    for (i = 0; i < M; i++)
    {
        if (a[Idx][i] == 16)
        {
            elements++;
        }
    }
    num[Idx] = elements;

        //Here at this point, the thread knows the number of elements with the value 16 of its line and would 
        //need to allocate just as much extra memory in "listeM" and "listeN". Is that possible ?

    __syncthreads();

    if (Idx == 0)
    {
                //This function sets the value of each element in the array "num" to the total value of the 
                //elements previously counted in all lines with the value 16.
                //Example: Input: num{2,4,3,1} Output: num{0,2,6,9}
        PrefixSum(num, N);
    }

    __syncthreads();

        // The output of PrefixSum(num, N) can now be used to realize the last task (put the "coordinates" of 
        //each 16 in the two arrays ("listM" and "listN") and each thread starts at the position equal the 
        //number of counted 16s).
    for (i = 0; i < M; i++)
    {
        if (a[Idx][i] == 16)
        {
            listM[num[Idx] + i] = Idx;
            listN[num[Idx] + i] = i;
        }
    }
}

【问题讨论】：

标签： cuda kernel nvidia dynamic-memory-allocation memory-reallocation

【解决方案1】：

是否可以编写内核，以便每个线程在计算其行中值为 16 的元素数量（仅此数字）后动态扩大两个数组“listeM”和“listeN”？

CUDA 设备代码无法扩大使用主机端 cudaMalloc、cudaMallocManaged、cudaHostAlloc 或类似名称创建的现有分配。

可以使用内核中的new 或malloc 将CUDA 设备代码转换为create new allocations，但是来自此类分配的数据不能直接传输回主机。要将其传输回主机将需要主机端分配，这样分配中的数据可以复制到其中，这使您回到最初的问题。

因此，确实没有方便的方法来做到这一点。您的选择是：

（过度）根据可能返回的最大大小分配所需的大小。
创建一个算法，运行内核一次以确定所需的大小，然后将该大小返回给主机。然后，主机分配该大小并将其传递给内核以供在第二次调用算法时使用，该算法执行实际所需的工作。

“可能的”第三种方法是：

只运行一次算法，让内核在内核中分配以提供所需的额外空间。主机端操作无法访问此空间。该内核还将返回此类分配的大小和/或排列。
根据返回的设备大小分配的大小/排列，主机将分配所需大小的新内存。
然后主机将启动一个新的“复制内核”，它将数据从步骤 1 中的设备端分配复制到步骤 2 中提供的主机端分配。
然后主机会将数据从第 2 步中的主机端分配复制到主机内存。

对于您所概述的这样一个微不足道的问题，这是一个极端复杂的问题，显而易见的解决方案就是过度分配所需的空间并完成它。

【讨论】：