GPU 多处理器的内核块执行之间的共享内存会发生什么情况？

Question

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

• 它与最后一个预定块离开时的状态相同。
• 它是空白的
• 它包含不可预见的垃圾。

为了缩小可能出现的情况的变化，请具体参考每个块使用最大可能共享内存量的情况 - 在开普勒 GPU 上为 48 KB。

Answer 1

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

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

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

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

编辑：

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

在 GTX660（Kepler、CC 3.0、5 个 SM）上，配置 48K 共享内存，CC 3.0 Release build，我得到以下结果：

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

我跑了好几次，每次都得到相同的结果。这符合我最初的猜测，因为 10000 个整数占用 ~40K，所以每个 SM 有一个并发块的空间，并且这个设备有 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 个块的空间）并始终如一地得到：

Shared memory initializations: 480

所以，再次，不是预期的数量，令人惊讶的是，即使每个块使用的共享内存量较小，但可能的变化显然更少。

看起来，如果你决心搬起石头砸自己的脚，你可以使用一个大的共享内存块来保持一致:)此外，这是在也用于显示的 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;
    }
  }
}

Answer 2

状态未定义。这意味着它可以是任何东西，包括你猜到的三件事中的任何东西。但是从未初始化的内存中读取数据也可能导致你的 GPU 出现人工智能。