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

因此,再次,不是预期的数字,并且令人惊讶的是,即使每个块使用的共享内存量较小,但可能的变化明显较少.

看起来,如果你决定用脚射击自己,你可以使用一个大的共享内存块来保持一致:)此外,这是在一个也用于显示的GPU上运行,Windows 7带有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;
    }
  }
}