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

Question

假设我有一个包含许多块的CUDA内核，假设某个块在同一对称多处理器上紧随另一个块的安排（也就是说，所有warp的共享内存区域相同）。目前，NVIDIA没有在API或每个GPU文档中指定共享内存在执行之间发生了什么。实际上，以下哪种情况适用于块的共享内存内容？

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

为了缩小可能出现的情况范围，请具体参考每个块使用最大共享内存量的情况-Kepler GPU上的48 KB。

- einpoklum

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版本，我得到以下结果：

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

我多次运行了这个程序，每一次都得到了相同的结果。这符合我最初的猜测：10000个整数占用约40K的空间，因此每个SM只有一个并发块的空间，而这个设备有5个SMs。

但是，当我将共享内存减少到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;
    }
  }
}

- Roger Dahl

---

从语言学的角度来看，这是未定义的行为，这就是你需要了解的全部内容。 - talonmies

重点是执行顺序和硬件亲和性未定义。访问在当前范围内被定义但未初始化的内存是未定义行为。编写故意依赖于任何未定义行为的代码只是纯粹的愚蠢。 - talonmies

@einpoklum：通常有些事情是故意留下未定义的，这样实现就可以在不必关心明确将某些事物放入特定状态的情况下进行优化和重新排列。实际上，许多编译器的优化都基于某些事情是未定义的，因此不能在程序中执行，这为编译器提供了一些关于程序员可能不希望从他所在的环境中得到的知识的余地。这是关于未定义行为的一个事实，你只能接受它，没有其他办法。 - datenwolf

@datenwolf：也许你是对的。但是-我们不知道。也许这仅仅是因为NVidia不想承诺未来GPU可能如何改变。也许没有人费心将其放入规范中。也许这种行为很丑陋和复杂。这就是我问这个问题的原因。 - einpoklum

@einpoklum：是的，这些都是离开它未定义的非常有效的理由。但这并不意味着在某个地方有一个“秘密”或“隐含”的定义可以依赖。某些事情为什么未定义并不重要。唯一重要的是它们是未定义的，这本质上在上面贴了一张明亮的、红色的胶带，上面用非常粗的字母写着“不要以任何方式使用或依赖它”。 - datenwolf

Answer 2

状态未定义。这意味着它可以是任何东西，包括您猜测的三个中的任何一个。但从未初始化的内存中读取可能会导致您的GPU出现人工智能。

- datenwolf

---

@einpoklum：那个问题无法回答，因为行为可能会随着每个驱动程序版本的变化而改变，甚至可能取决于将GPU放入的系统硬件配置。 - datenwolf

请求某人给你一个未定义行为的定义，这对我来说是一个非常奇怪的问题。向那些愿意尝试回答它的人致敬。 - Robert Crovella