如果为整个网格分配的共享内存量超过 48kB，则内核启动失败

Question

我正在处理需要大量 共享内存 的 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, ...);

对于小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 随心所欲）。

Answer 1

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

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

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

不过我们可以从中得出两个结论：

========= Invalid __global__ write of size 8

内核确实启动了。
报告的失败与越界索引到全局内存有关，在写入到全局内存时，不是共享内存。（因此，正如您的猜想所暗示的那样，它不可能发生在从全局内存读取以填充共享内存时。）

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

如果您想追踪对内核中特定代码行的无效全局写入的来源，this question/answer 可能会感兴趣。

Answer 2

我认为这里的问题是块内的所有线程必须运行在同一个 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;

}

Answer 3

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

如果为整个网格分配的共享内存量超过 48kB，则内核启动失败

Kernel launch failure if the amount of shared memory allocated for the whole grid exceeds 48kB

cuda

gpu-shared-memory