0

我正在尝试在 CUDA 的帮助下减少 65536 个元素数组（计算其中元素的总和）。内核如下所示（请暂时忽略 *dev_distanceFloats 和索引参数）

__global__ void kernel_calcSum(float *d, float *dev_distanceFloats, int index) {
  int tid = threadIdx.x;
  float mySum = 0;
  for (int e = 0; e < 256; e++) {
    mySum += d[tid + e];
  }
}

ant 它作为一个具有 256 个线程的块启动：

kernel_calcSum <<<1,  256 >>>(dev_spFloats1, dev_distanceFloats, index);

到目前为止，一切顺利，256 个线程中的每一个都从全局内存中获取 256 个元素，并在局部变量 mySum 中计算它的总和。内核执行时间约为 45 毫秒。下一步是在块中的这 256 个线程之间引入共享内存（以计算 mySum 的总和），因此内核变为如下：

__global__ void kernel_calcSum(float *d, float *dev_distanceFloats, int index) {
  __shared__ float sdata[256];
  int tid = threadIdx.x;
  float mySum = 0;
  for (int e = 0; e < 256; e++) {
    mySum += d[tid + e];
  }
  sdata[tid] = mySum;
}

我刚刚添加了对共享内存的写入，但执行时间从 45 毫秒增加到 258 毫秒（我在 NVidia Visual Profiler 5.0.0 的帮助下对此进行了检查）。我意识到在写入 sdata 变量时每个线程有 8 个存储库冲突（我在 GTX670 上，它具有 32 个存储库的能力 3.0）。作为一个实验——我尝试在启动内核时将线程数减少到 32 个——但时间仍然是 258 毫秒。

问题 1：为什么在我的情况下写入共享内存需要这么长时间？问题 2：是否有任何工具可以详细显示“执行计划”（内存访问时间、冲突等）？

感谢您的建议。

更新： 使用内核 - 我将 sdata 设置为每个线程的某个常量：

__global__ void kernel_calcSum(float *d, float *dev_distanceFloats, int index) {
  __shared__ float sdata[256];
  int tid = threadIdx.x;
  float mySum = 0;
  for (int e = 0; e < 256; e++) {
    mySum += d[tid + e];
  }
  sdata[tid] = 111;
}

时间回到 48 毫秒。所以，改变 sdata[tid] = mySum; 到 sdata[tid] = 111; 做了这个。

这个编译器优化（可能只是删除了这一行？）还是出于某种原因从本地内存（寄存器？）复制到共享需要很长时间？

4

2 回答 2

0

您的两个内核都不做任何事情，因为它们不会将结果写入内核完成后仍可访问的内存。

在第一种情况下，编译器足够聪明地注意到这一点并优化整个计算。在涉及共享内存的第二种情况下，编译器不会注意到这一点，因为通过共享内存的信息流将更难跟踪。因此，它留下了计算。

传入一个指向全局内存的指针（就像你已经做的那样）并通过这个指针写出结果。

于 2013-03-15T11:06:10.543 回答
0

共享内存不适合这个。你需要的是warp原子操作，在warp中总结，然后在warp之间传达中间结果。有示例代码演示了这种与 CUDA 一起发布的内容。

总结元素是大规模并行化无济于事的任务之一，实际上 GPU 可以被 CPU 超越。

于 2013-03-15T10:37:11.897 回答