求和 3D 矩阵 cuda答案

【问题标题】：Sum 3D matrix cuda求和 3D 矩阵 cuda
【发布时间】：2012-03-30 10:09:28
【问题描述】：

我需要进行如下计算：A[x][y] = sum{从 z=0 到 z=n}{B[x][y][z]+C[x][y][z ]}，其中矩阵 A 的维度为 [height][width]，矩阵 B,C 的维度为 [height][width][n]。

值通过以下方式映射到内存：

index = 0;
for (z = 0; z<n; ++z)
    for(y = 0; y<width; ++y)
        for(x = 0; x<height; ++x) {
            matrix[index] = value;
            index++;
        }

Q1：这个 Cuda 内核可以吗？

idx = blockIdx.x*blockDim.x + threadIdx.x;
idy = blockIdx.y*blockDim.y + threadIdx.y;

for(z=0; z<n; z++){
    A[idx*width+idy] += B[idx*width+idy+z*width*height] + C[idx*width+idy+z*width*height];
}

Q2：这种计算方式更快吗？

idx = blockIdx.x*blockDim.x + threadIdx.x;
idy = blockIdx.y*blockDim.y + threadIdx.y;
idz = blockIdx.z*blockDim.z + threadIdx.z;

int  stride_x = blockDim.x * gridDim.x;
int  stride_y = blockDim.y * gridDim.y;
int  stride_z = blockDim.z * gridDim.z;

while ( idx < height && idy < width && idz < n ) {
    atomicAdd( &(A[idx*width+idy]), B[idx*width+idy+idz*width*height] + C[idx*width+idy+idz*width*height] );
    idx += stride_x;
    idy += stride_y;
    idz += stride_z;
}

【问题讨论】：

标签： c cuda gpu

【解决方案1】：

第一个内核没问题。但是我们还没有合并对矩阵B 和C 的访问。

至于第二个核函数。您有数据竞赛，因为不仅一个线程能够写入A[idx*width+idy] 地址。你需要额外的同步，比如AttomicAdd

关于一般问题：我认为实验表明它更好。这取决于您拥有的典型矩阵大小。请记住，Fermi

【讨论】：

数据竞赛的好处，谢谢 :) 当我添加原子函数时，代码形式 Q2 将与 Q1 中的代码完全相同？我的主要目标是速度，所以我正在寻找最快的解决方案。
“实验表明它更好”，使用原子函数更好吗？还是不使用？
通常同步点会显着降低性能。它与一段时间内并行执行的扭曲数密切相关。最好使用共享内存来存储中间结果。在共享内存中同步是可以的，但我更喜欢在全局内存中避免它。

【解决方案2】：

ArrayFire 中的真正简单：

array A = randu(nx,ny,nz);
array B = sum(A,2); // sum along 3rd dimension
print(B);

【讨论】：

【解决方案3】：

Q1：用你知道答案的矩阵来测试它

备注：使用非常大的矩阵时可能会遇到问题。使用具有适当增量的 while 循环。 Cuda by Example 和往常一样是参考书。

可以在此处找到实现嵌套循环的示例：For nested loops with CUDA。实现了一个while循环。

marina.k 关于比赛条件是正确的。这将有利于方法一，因为原子操作往往会减慢代码速度。

【讨论】：

线程块的最大维数 = 3，线程块网格的最大维数 = 2（对于 CC = 2）。所以我认为没有问题。我有大型矩阵，但没有发现问题。
你提醒：if ( idx >= height || idy >= width || idz >= n ) return; ?
是的，但你不需要 if 而是一段时间。假设您有 1024 个内核，但矩阵的大小为 2048，那么您将无法达到所有索引。