cuda - GPU 多处理器的内核 block 执行之间的共享内存会发生什么变化？

Question

假设我有一个带有一堆 block 的 CUDA 内核，并且假设某个 block 紧接在同一个对称多处理器上的另一个 block 之后(即，所有线程束具有相同共享内存区域的单元)。 NVIDIA 目前并未在 API 或每个 GPU 文档中指定执行之间的共享内存会发生什么情况。但实际上，关于 block 的共享内存内容，以下哪项成立？ :

• 它处于上一个调度 block 离开它时的相同状态。
• 它是空白的。
• 它包含不可预见的垃圾。

为了缩小可能出现的情况的变化，请具体引用每个 block 使用最大可能共享内存量的情况 - 在 Kepler GPU 上为 48 KB。

Answer 1

NVIDIA 不会发布此级别的硬件行为，因此您应该将其视为未定义(如@datenwolf 所说)。当然，给定 block 看到的共享内存的内容不会是随机的。硬件没有必要花时间清理内存。

GPU 可以在每个 SM 上同时运行多个 block 。给定内核同时运行的 block 数取决于各种因素。因此，例如，如果共享内存是限制因素，则每个 SM 将运行与共享内存中适合的 block 数。所以，如果有 48K 的共享内存，一个 block 需要 10K，那么 4 个 block 可能同时运行，使用 40K。因此，如果您的设备有 8 个 SM，我的猜测是给定 block 的共享内存将有 32 (4 * 8) 个可能的固定位置。因此，当安排一个新 block 时，它将被分配到这些位置之一，并查看共享内存，因为它是由在该位置运行的前一个 block 留下的。

API 无法让 block 检测它在哪个位置运行。 block 的调度是动态确定的，可能很难预测。

如果 GPU 用于显示，它可能同时运行其他内核(着色器)，可能以奇怪而奇妙的方式覆盖 CUDA 内核中 block 之间的共享内存。甚至 CUDA 也可能在幕后运行其他内核。

编辑:

我写了一个小程序来测试(包括在下面)。该程序将一个 block 应存储在共享内存中的整数个数作为参数。然后它启动 100,000 个 block ，每个 block 有一个线程。每个 block 检查其共享内存是否已初始化。如果它被初始化，该 block 什么都不做。如果未初始化，该 block 会初始化内存并增加全局计数。初始化模式是递增的数字序列，以避免部分重叠的初始化共享内存缓冲区看起来有效。

在 GTX660(开普勒、CC 3.0、5 个 SM)上，配置 48K 共享内存，CC 3.0 发布版本，我得到了以下结果:

C:\rd\projects\cpp\test_cuda\Release>test_cuda.exe 10000
Shared memory initializations: 5

我运行了好几次，每次都得到相同的结果。这符合我最初的猜测，因为 10000 个整数占用 ~40K，因此每个 SM 有一个并发 block 的空间，并且这个设备有 5 个 SM。

但是，当我将共享内存减少到 2500 个整数(~10K)，期望进行 20 次初始化并运行几次时，我得到了不同的高数字:

Shared memory initializations: 32,822
Shared memory initializations: 99,996
Shared memory initializations: 35,281
Shared memory initializations: 30,748

因此，在这种情况下，我对固定位置的猜测是完全无效的。

然后我尝试将共享内存减少到 100 个整数(在 48K 中将有 122 个 block 的空间)并始终如一地得到:

Shared memory initializations: 480

因此，同样，这不是预期的数量，而且令人惊讶的是，尽管每个 block 使用的共享内存量较小，但可能的变化显然较少。

看起来，如果你决心搬起石头砸自己的脚，你可以使用一个大的共享内存块来保持一致 :) 此外，这是在也用于显示的 GPU 上运行的，Windows 7 Aero(一种 GPU 加速主题)并且渲染似乎不会受到干扰，因为在内核运行时桌面会卡住。

程序:

#include "cuda_runtime.h"

#include <iostream>
#include <sstream>
using namespace std;

#define assertCudaSuccess(ans) { _assertCudaSuccess((ans), __FILE__, __LINE__); }
inline void _assertCudaSuccess(cudaError_t code, char *file, int line)
{
  if (code != cudaSuccess) {
    fprintf(stderr,"CUDA Error: %s %s %d\n", cudaGetErrorString(code), file, line);
    exit(code);
  }
}

__global__ void shared_memory_persistence_test(int n_shared_ints);
__device__ int init_cnt_d(0);

int main(int argc, char* argv[])
{
  cout.imbue(locale(""));
  int n_shared_ints;
  stringstream(string(argv[1])) >> n_shared_ints;
  shared_memory_persistence_test<<<dim3(100, 1000), 1, n_shared_ints * sizeof(int)>>>(n_shared_ints);
  assertCudaSuccess(cudaPeekAtLastError());
  assertCudaSuccess(cudaDeviceSynchronize());
  int init_cnt_h;
  assertCudaSuccess(cudaMemcpyFromSymbol(&init_cnt_h, init_cnt_d, sizeof(int), 0, cudaMemcpyDeviceToHost));
  cout << "Shared memory initializations: " << init_cnt_h << endl;
  return 0;
}

__global__ void shared_memory_persistence_test(int n_shared_ints)
{
  extern __shared__ int shared[];

  for (int i(0); i < n_shared_ints; ++i) {
    if (shared[i] != i) {
      for (int i(0); i < n_shared_ints; ++i) {
        shared[i] = i;
      }
      atomicAdd(&init_cnt_d, 1);
      break;
    }
  }
}