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

Question

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

目前，如果我有以下顺序 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 线程，因为所有迭代都是独立的并且可以同时执行。

Answer 1

你可以保持外循环不变。此外，最好使用.x 作为最内层循环，这样您就可以使用access the global memory efficiently。

__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 维索引。

__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，并以类似的方式仅计算k,l 和.x 的更多维度。