0

我对以下任务有疑问：

“给定一个二维数组“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;
        }
    }
}
4

1 回答 1

1

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

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

CUDA 设备代码可以使用内核或创建新分配，但是来自此类分配的数据不能直接传输回主机。要将其传输回主机需要主机端分配，这样分配中的数据可以复制到其中，这使您回到原始问题。newmalloc

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

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

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

  1. 只运行一次算法，让内核在内核中分配以提供所需的额外空间。主机端操作无法访问此空间。该内核还将返回此类分配的大小和/或排列。

  2. 基于设备大小分配的返回大小/排列，主机将分配所需大小的新内存。

  3. 然后主机将启动一个新的“复制内核”，它将数据从步骤 1 中的设备端分配复制到步骤 2 中提供的主机端分配。

  4. 然后，主机将在步骤 2 中将数据从主机端分配复制到主机内存。

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

于 2019-12-01T17:04:50.137 回答