CUDA共享内存速度

Question

这是一个与性能相关的问题。我已经根据“CUDA By Example”示例代码编写了以下简单的 CUDA 内核：

#define N 37426 /* the (arbitrary) number of hashes we want to calculate */
#define THREAD_COUNT 128

__device__ const unsigned char *m = "Goodbye, cruel world!";

__global__ void kernel_sha1(unsigned char *hval) {
  sha1_ctx ctx[1];
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;

  while(tid < N) {
    sha1_begin(ctx);
    sha1_hash(m, 21UL, ctx);
    sha1_end(hval+tid*SHA1_DIGEST_SIZE, ctx);
    tid += blockDim.x * gridDim.x;
  }
}

在我看来，代码是正确的，并且确实吐出了相同哈希的 37,426 个副本（正如预期的那样。根据我对第 5 章第 5.3 节的阅读，我假设每个写入全局内存的线程都以“ hval" 会非常低效。

然后我实现了我认为使用共享内存的性能提升缓存。代码修改如下：

#define N 37426 /* the (arbitrary) number of hashes we want to calculate */
#define THREAD_COUNT 128

__device__ const unsigned char *m = "Goodbye, cruel world!";

__global__ void kernel_sha1(unsigned char *hval) {
  sha1_ctx ctx[1];
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  __shared__ unsigned char cache[THREAD_COUNT*SHA1_DIGEST_SIZE];

  while(tid < N) {
    sha1_begin(ctx);
    sha1_hash(m, 21UL, ctx);
    sha1_end(cache+threadIdx.x*SHA1_DIGEST_SIZE, ctx);

    __syncthreads();
    if( threadIdx.x == 0) {
      memcpy(hval+tid*SHA1_DIGEST_SIZE, cache, sizeof(cache));
    }
    __syncthreads();
    tid += blockDim.x * gridDim.x;
  }
}

第二个版本似乎也可以正常运行，但比初始版本慢几倍。后者代码在大约 8.95 毫秒内完成，而前者在大约 1.64 毫秒内运行。我对 Stack Overflow 社区的问题很简单：为什么？

Answer 1

我通过示例查看了 CUDA，但找不到任何类似的东西。是的，附录中有一些关于 GPU 哈希表的讨论，但它看起来不像这样。所以我真的不知道你的函数是做什么的，尤其是 sha1_end。如果此代码与那本书中的内容相似，请指出，我错过了。

但是，如果 sha1_end 写入一次全局内存（每个线程）并以合并的方式执行，那么它没有理由不能非常高效。大概每个线程都在写入不同的位置，因此如果它们或多或少相邻，则肯定有合并的机会。无需详细讨论合并的细节，只需说它允许多个线程在单个事务中将数据写入内存。而且，如果您要将数据写入全局内存，您将不得不在某个地方至少支付一次这种惩罚。

对于您的修改，您已经完全扼杀了这个概念。您现在已经从单个线程执行了所有数据复制，并且 memcpy 意味着后续数据写入（整数或字符等）发生在单独的事务中。是的，有一个缓存可以帮助解决这个问题，但在 GPU 上这样做是完全错误的。让每个线程更新全局内存，并利用机会并行进行。但是，当您强制在单个线程上进行所有更新时，该线程必须按顺序复制数据。这可能是时间差异中最大的单一成本因素。

使用 __syncthreads() 也会产生额外的成本。

CUDA by Examples 一书的第 12.2.7 节提到了视觉分析器（并提到它可以收集有关合并访问的信息）。视觉分析器是帮助尝试回答此类问题的好工具。

如果您想了解有关高效内存技术和合并的更多信息，我会推荐 NVIDIA GPU 计算webinar，标题为“GPU Computing using CUDA C – Advanced 1 (2010)”。它的直接链接是here 和slides。