小编Cri*_*rro的帖子

编辑:似乎,至少在这种情况下,转置网格对L2缓存带宽有负面影响.这是从视觉分析器获得的.原因尚不清楚.

我已经到了需要转置CUDA网格的GPU计算情况.因此,如果block_ {x,y}最初作用于数据区域d_ {x,y},则它现在作用于数据区域d_ {y,x},因此block_ {y,x}将作用于数据区域d_ {x, y}.下图显示了一个示例.

值得一提的是,线程不会在每个块内部进行转置,也就是说,一旦定位了块,threadIdx.x和threadIdx.y值分别以正常方式用于它们的x和y偏移.

据我所知,理论上这个设计应该不会对性能产生任何影响,因为内存合并模式仍然保留,即块内的线程没有转置,只是网格重新排列其块.但是我发现在转置网格时,内核运行大约.比正常情况慢2倍.我做了一个玩具示例来说明情况.

?  transpose-grid ./prog 10000 10000 100 0
init data.....................done: zero matrix of 10000 x 10000
copy data to GPU..............done
preparing grid................done: block(32, 32, 1), grid(313, 313, 1)
normal_kernel (100 rep).......done: 0.935132 ms
verifying correctness.........ok
?  transpose-grid ./prog 10000 10000 100 1
init data.....................done: zero matrix of 10000 x 10000
copy data to GPU..............done
preparing grid................done: …

使用CUDA Dynamic Parallelism时,我们遇到了性能问题.目前,CDP的表现比传统方法慢至少3倍.我们制作了最简单的可重现代码来显示这个问题,即将数组的所有元素的值增加+1.即

a[0,0,0,0,0,0,0,.....,0] --> kernel +1 --> a[1,1,1,1,1,1,1,1,1]

这个简单示例的目的只是为了查看CDP是否可以像其他CDP一样执行,或者是否存在严重的开销.

代码在这里:

#include <stdio.h>
#include <cuda.h>
#define BLOCKSIZE 512

__global__ void kernel_parent(int *a, int n, int N);
__global__ void kernel_simple(int *a, int n, int N, int offset);


// N is the total array size
// n is the worksize for a kernel (one third of N)
__global__ void kernel_parent(int *a, int n, int N){
    cudaStream_t s1, s2;
    cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
    cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);

    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if(tid == …