这是一个与性能相关的问题。我基于“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 社区的问题很简单:为什么?