在内核块执行之间，GPU多处理器的共享内存会如何处理？

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

GPU可以在每个SM上同时运行许多块。给定内核同时运行的块数取决于各种因素。因此，例如，如果共享内存是限制因素，则每个SM将运行与共享内存中一样多的块。因此，如果共享内存为48K，而一个块需要10K，则可以同时使用40K运行4个块。因此，如果您的设备具有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上运行的，该GPU也用于显示Windows 7 with Aero（A 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;
    }
  }
}