CUDA：写入共享内存increses内核时执行了很多

我试图减少65536个元素数组（计算它的元素的总和）与CUDA的帮助。内核看上去像下面（请无视* dev_distanceFloats和索引参数现在）CUDA：写入共享内存increses内核时执行了很多

__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]; 
    } 
}

蚂蚁又推出作为一个块与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来操作）。我认识到，有8个为每个线程冲突写SDATA变量（我上有32个存储能力3.0 GTX670）时。作为一个实验 - 我试图在启动内核时将线程数减少到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; 做了这个。

这是编译器优化（可能只是删除了这一行吗？）或者由于某些原因从本地内存（register？）复制到共享需要很长时间？

来源

2013-03-15 Arsen

这两个内核都不会做任何事情，因为它们不会将内存写出来，内核完成后仍然可以访问内存。

在第一种情况下，编译器足够聪明，可以注意到这一点并优化整个计算。在共享存储器被涉及的第二种情况下，编译器不会注意到此作为信息通过共享存储器中的流动将更难跟踪。因此它将计算放入。

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

来源

2013-03-15 11:06:10 tera

我明白了。所以，据我了解 - 事实上在第一种情况下，cuda运行空内核（因为编译器决定不会在任何情况下使用结果）？在第二种情况下，内核代码不是空的，并且对于内存访问和其他正常的事情，时间增加到258毫秒，对吧？ – Arsen 2013-03-15 11:12:26

是的，没错。 – tera 2013-03-15 11:16:56

谢谢，泰拉！这对我来说看起来是真的。 – Arsen 2013-03-15 11:22:54

共享内存是不能这样做的正确的事情。你需要的是经纱原子操作，总结一个经纱，然后传达经纱之间的中间结果。有一些示例代码演示了与CUDA一起发货。

总结元素是其中大规模并行化不会有太大帮助的任务之一，并且GPU实际上可以超过CPU。

来源

2013-03-15 10:37:11 datenwolf

datenwolf，谢谢你的回答。我意识到减少并不完全揭示并行计算的力量，然而，这是我的algorythm中的一个步骤:)。我检查了Mark Harris的“优化CUDA中的并行还原”中的代码，但在他的情况下使用共享内存并没有引入这样的效果。 – Arsen 2013-03-15 10:52:34

CUDA：写入共享内存increses内核时执行了很多

回答

相关问题