使用 CUDA 并行化四个或更多嵌套循环

Parallelize four and more nested loops with CUDA

本文关键字：四个 嵌套循环 CUDA 并行化 使用      更新时间：2023-10-16

我正在研究一个生成并行C++代码的编译器。我是 CUDA 编程的新手，但我正在尝试使用 CUDA 并行化C++代码。

目前，如果我有以下顺序C++代码：

for(int i = 0; i < a; i++) {
    for(int j = 0; j < b; j++) {
        for(int k = 0; k < c; k++) {
            A[i*y*z + j*z + k*z +l] = 1;
        }
    }
}

这将产生以下 CUDA 代码：

__global__ void kernelExample() {
    int _cu_x = ((blockIdx.x*blockDim.x)+threadIdx.x);
    int _cu_y = ((blockIdx.y*blockDim.y)+threadIdx.y);
    int _cu_z = ((blockIdx.z*blockDim.z)+threadIdx.z);
    A[_cu_x*y*z + _cu_y*z + _cu_z] = 1;
}
因此，每个循环嵌套

都映射到一个维度，但是并行化四个或更多嵌套循环的正确方法是什么：

for(int i = 0; i < a; i++) {
    for(int j = 0; j < b; j++) {
        for(int k = 0; k < c; k++) {
            for(int l = 0; l < d; l++) {
                A[i*x*y*z + j*y*z + k*z +l] = 1;
            }
        }
    }
}

有没有类似的方法？值得注意的是：所有循环维度都是平行的，迭代之间没有依赖关系。

提前感谢！

编辑：目标是将所有迭代映射到 CUDA 线程，因为所有迭代都是独立的并且可以并发执行。

您可以保持外部循环不变。此外，最好将.x用作最内部的循环，以便您可以有效地访问全局内存。

__global__ void kernelExample() {
    int _cu_x = ((blockIdx.x*blockDim.x)+threadIdx.x);
    int _cu_y = ((blockIdx.y*blockDim.y)+threadIdx.y);
    int _cu_z = ((blockIdx.z*blockDim.z)+threadIdx.z);
    for(int i = 0; i < a; i++) {
        A[i*x*y*z + _cu_z*y*z + _cu_y*z + _cu_x] = 1;
    }
}

但是，如果a,b,c,d都非常小，则可能无法获得足够的并行度。在这种情况下，您可以将线性索引转换为 n-D 索引。

__global__ void kernelExample() {
    int tid = ((blockIdx.x*blockDim.x)+threadIdx.x);
    int i = tid / (b*c*d);
    int j = tid / (c*d) % b;
    int k = tid / d % c;
    int l = tid % d;
    A[i*x*y*z + j*y*z + k*z + l] = 1;
}

但请注意，计算i,j,k,l可能会引入大量开销，因为整数除法和 mod 在 GPU 上很慢。作为替代方案，您可以将i,j映射到.z和.y，并以类似的方式仅从.x计算k,l和更多维度。

相关文章：