1

我正在研究需要大量共享内存的 N 体问题。

基本上，有N独立的任务，每个任务使用 4 个双精度变量，即 32 个字节。单个任务由一个线程执行。

为了快速起见，我一直在为这些变量使用共享内存（假设寄存器也被线程使用）。由于N在编译时不知道任务的数量，因此共享内存是动态分配的。

  • 网格的维度和共享内存的计算取决于N块大小：

    const size_t BLOCK_SIZE = 512;
const size_t GRID_SIZE = (N % BLOCK_SIZE) ? (int) N/BLOCK_SIZE : (int) N/BLOCK_SIZE +1;
const size_t SHARED_MEM_SIZE = BLOCK_SIZE * 4 * sizeof(double);

  • 然后使用这 3 个变量启动内核。

    kernel_function<<<GRID_SIZE, BLOCK_SIZE, SHARED_MEM_SIZE>>>(N, ...);

对于 small N，这可以正常工作，并且内核执行时不会出错。

但是如果超过N = 1500，内核启动失败（以下消息多次出现）：

========= Invalid __global__ write of size 8
=========
========= Program hit cudaErrorLaunchFailure (error 4) due to "unspecified launch failure" on CUDA API call to cudaLaunch.

据我了解，这是由于试图写入超出已分配共享内存的界限。当在内核中将全局内存复制到共享内存中时，就会发生这种情况：

__global__ void kernel_function(const size_t N, double *pN, ...)
{
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;

    if(idx<N)
    {
        extern __shared__ double pN_shared[];
        for(int i=0; i < 4; i++)
        {
            pN_shared[4*idx + i] = pN[4*idx + i];
        }
        ...
    }
}

此错误仅在 时发生，因此当共享内存N > 1500总量超过48kB ( ) 时。 无论网格和块大小如何，此限制都是相同的。1500 * 4 * sizeof(double) = 1500 * 32 = 48000

如果我正确理解了 CUDA 的工作原理，那么网格使用的共享内存的累积量不限于48kB，这只是单个线程块可以使用的共享内存的限制。

这个错误对我来说毫无意义，因为累积的共享内存量只会影响流式多处理器之间的网格调度方式（此外，GPU 设备有 15 个 SM 可供使用）。

4

3 回答 3

3

您在此处动态分配的共享内存量：

kernel_function<<<GRID_SIZE, BLOCK_SIZE, SHARED_MEM_SIZE>>>(N, ...);
                                         ^^^^^^^^^^^^^^^

是每个线程块的数量，并且该数量限制为 48KB（即 49152，而不是 48000）。因此，如果您尝试在那里分配超过 48KB 的空间，则在检查时应该会收到错误消息。

但是，我们可以从中得出两个结论：

========= Invalid __global__ write of size 8
  1. 内核确实启动了。
  2. 报告的失败与全局内存的越界索引、写入全局内存而不是共享内存有关。（因此，正如您的猜想所暗示的，它不会发生在从全局内存中读取以填充共享内存的情况下。）

所以总的来说我认为你的结论是不正确的，你可能需要做更多的调试，而不是得出关于共享内存的结论。

如果您想追踪对内核中特定代码行的无效全局写入的来源，这个问题/答案可能很有趣。

于 2016-05-20T14:00:03.333 回答
1

您正在访问索引 idx*4+0:3 处的共享数组。程序从 N > BLOCK_SIZE 开始不正确。幸运的是，它似乎可以达到 1500。但是使用 cuda mem-check 应该可以指出这个问题。在相关主题上，请注意在另一个位置静态分配的共享内存可能会使用共享内存。打印出指针的值将有助于弄清楚。

于 2016-05-20T15:59:25.603 回答
0

我认为这里的问题是块内的所有线程都必须在同一个 SM 中运行。因此，每个块仍然具有 48kB 共享内存的硬限制。在该块中运行多少线程并不重要。调度无关紧要，因为 GPU 不能跨多个 SM 拆分块中的线程。如果可以的话，我会尝试减少 BLOCK_SIZE，因为这将直接决定每个块的共享内存量。但是，如果您将其减少得太远，您可能会遇到无法充分利用 SM 中的计算资源的问题。这是一种平衡行为，根据我的经验，CUDA 架构提出了许多类似这样的有趣权衡。

同样在您的情况下，我什至不确定您是否需要共享内存。我只会使用局部变量。我认为局部变量存储在全局内存中，但是对它们的访问是对齐的，所以它非常快。如果你想用共享内存做一些简洁的事情来提高性能，这里是我的 N-Body 模拟器的 OpenCL 内核。使用共享内存为块中的每个线程创建缓存可以让我的速度提高大约 10 倍。

在这个模型中，每个线程都负责计算一个物体上的加速度，这是由于其他物体受到引力的影响。这需要每个线程循环遍历所有 N 个主体。这通过共享内存缓存得到了增强，因为块中的每个线程都可以将不同的主体加载到共享内存中并且它们可以共享它们。 

__kernel void acceleration_kernel
(
    __global const double* masses, 
    __global const double3* positions,
    __global double3* accelerations,
    const double G,
    const int N,
    __local double4* cache //shared memory cache (local means shared memory in OpenCL)
)
{
    int idx = get_global_id(0);
    int lid = get_local_id(0);
    int lsz = get_local_size(0);

    if(idx >= N)
        return;

    double3 pos = positions[idx];
    double3 a = { };

    //number of loads required to compute accelerating on Body(idx) from all other bodies
    int loads = (N + (lsz - 1)) / lsz;

    for(int load = 0; load < loads; load++)
    {
        barrier(CLK_LOCAL_MEM_FENCE);

        //compute which body this thread is responsible for loading into the cache
        int load_index = load * lsz + lid;
        if(load_index < N)
            cache[lid] = (double4)(positions[load_index], masses[load_index]);

        barrier(CLK_LOCAL_MEM_FENCE);

        //now compute the acceleration from every body added to the cache
        for(int i = load * lsz, j = 0; i < N && j < lsz; i++, j++)
        {
            if(i == idx)
                continue;
            double3 r_hat = cache[j].xyz - pos; 
            double over_r = rsqrt(0.0001 + r_hat.x * r_hat.x + r_hat.y * r_hat.y + r_hat.z * r_hat.z);
            a += r_hat * G * cache[j].w * over_r * over_r * over_r;
        }
    }

    accelerations[idx] = a;
}
double3 pos = positions[idx];
double3 a = { };

int loads = (N + (lsz - 1)) / lsz;

for(int load = 0; load < loads; load++)
{
    barrier(CLK_LOCAL_MEM_FENCE);
    int load_index = load * lsz + lid;
    if(load_index < N)
        cache[lid] = (double4)(positions[load_index], masses[load_index]);
    barrier(CLK_LOCAL_MEM_FENCE);

    for(int i = load * lsz, j = 0; i < N && j < lsz; i++, j++)
    {
        if(i == idx)
            continue;
        double3 r_hat = cache[j].xyz - pos; 
        double over_r = rsqrt(0.0001 + r_hat.x * r_hat.x + r_hat.y * r_hat.y + r_hat.z * r_hat.z);
        a += r_hat * G * cache[j].w * over_r * over_r * over_r;
    }
}

accelerations[idx] = a;

}

于 2016-05-20T14:01:34.667 回答