如何在CUDA中嵌套for循环？

Question

我想提供一个完整的 CUDA 代码示例，其中包含有人可能想要包含的所有内容，以便尝试编写此类代码的人（例如我自己）引用它。

我主要关心的是是否可以在同一块中的不同线程上同时处理多个 for 循环。这是在示例代码中的案例 3 上运行（对于一个清晰的示例）总共 2016 个线程分为 32 个块和在理论上在每个 for 循环上运行 1024 个线程之间的区别，我们可以运行更少的代码通过在同一块下运行其他案例的 for 循环来另外 2 个块。否则，单独的案例将主要用于处理单独的任务，例如 for 循环。目前看来，CUDA 代码只知道何时并行运行。

// 注意：很少被引用，您可以看似按块并行处理 if 语句，我想说这是使用更多块而不是在调用期间增加每个块的线程数的主要目的，而不是需要多个SM（流式多处理器），上限为 2048 个线程（也是一个块的上限）//

如果我们有以下代码，包括 for 循环和 if 语句，那么优化并行化的代码会是什么？

public void main(string[] args) {

    doMath(3); // we want to process each statement in parallel. For this we use different blocks.
}

void doMath(int question) {
    int[] x = new int{0,1,2,3,4,5,6,7,8,9};
    int[] y = new int{0,1,2,3,4,5,6,7,8,10};
    int[] z = new int{0,1,2,3,4,5,6,7,8,11};
    int[] w = new int{0,1,2,3,4,5,6,7,8,12};
    int[] q = new int[1000];
    int[] r = new int[1000];
    int[] v = new int[1000];
    int[] t = new int[1000];


    switch(question) {
        case 1: 
            for (int a = 0; a < x.length; a++) {
                for (int b = 0; b < y.length; b++) {
                    for (int c = 0; c < z.length; c++) {
                        q[(a*100)+(b*10)+(c)] = x[a] + y[b] + z[c];
                    }
                }
            }
        break;
        case 2: 
            for (int a = 0; a < x.length; a++) {
                for (int b = 0; b < y.length; b++) {
                    for (int c = 0; c < w.length; c++) {
                        r[(a*100)+(b*10)+(c)] = x[a] + y[b] + w[c];
                    }
                }
            }
        break;
        case 3:
            for (int a = 0; a < x.length; a++) {
                for (int b = 0; b < z.length; b++) {
                    for (int c = 0; c < w.length; c++) {
                        v[(a*100)+(b*10)+(c)] = x[a] + z[b] + w[c];
                    }
                }
            }
            for (int a = 0; a < x.length; a++) {
                for (int b = 0; b < y.length; b++) {
                    for (int c = 0; c < w.length; c++) {
                        t[(a*100)+(b*10)+(c)] = x[a] + y[b] + w[c];
                    }
                }
            }
        break;
    }
}

从我看到的示例中，CUDA 代码如下：

// 3 个块用于 3 个 switch case 第三种情况需要 2000 个线程以完全并行完成，而前两种只需要 1000 个。块由 32 的倍数（线程）操作。诀窍是采用所有案例的最大公分母，或者 if/else 语句作为...案例...可能是，并适当每个案例所需的块数。 （在这个例子中，我们需要 127 个 32 个线程的块 (1024 * 2 + 2048 - 32)//

//旁注：每个 Streaming Multiprocessor 或 SM 只能支持 2048 个线程和 2048 / (# of blocks * # of threads/block)//

public void main(string[] args) {

    int *x, *y *z, *w, *q, *r, *t;

    int[] x = new int{0,1,2,3,4,5,6,7,8,9};
    int[] y = new int{0,1,2,3,4,5,6,7,8,10};
    int[] z = new int{0,1,2,3,4,5,6,7,8,11};
    int[] w = new int{0,1,2,3,4,5,6,7,8,12};
    int[] q = new int[1000];
    int[] r = new int[1000];
    int[] t = new int[1000];

    cudaMallocManaged(&x, x.length*sizeof(int));
    cudaMallocManaged(&y, y.length*sizeof(int));
    cudaMallocManaged(&z, z.length*sizeof(int));
    cudaMallocManaged(&w, w.length*sizeof(int));
    cudaMallocManaged(&q, q.length*sizeof(int));
    cudaMallocManaged(&r, r.length*sizeof(int));
    cudaMallocManaged(&t, t.length*sizeof(int));

    doMath<<<127,32>>>(x, y, z, w, q, r, t); 

    cudaDeviceSynchronize();

    cudaFree(x);
    cudaFree(y);
    cudaFree(z);
    cudaFree(w);
    cudaFree(q);
    cudaFree(r);
    cudaFree(t);
}

__global__
void doMath(int *x, int *y, int *z, int *w, int *q, int *r, int *t) {

    int index = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    switch(question) {
        case 1: 
            for (int a = index; a < x.length; a+=stride ) {
                for (int b = index; b < y.length; b+=stride) {
                    for (int c = index; c < z.length; c+=stride) {
                        q[(a*100)+(b*10)+(c)] = x[a] + y[b] + w[c];
                    }
                }
            }
        break;
        case 2: 
            for (int a = index; a < x.length; a+=stride) {
                for (int b = index; b < y.length; b+=stride) {
                    for (int c = index; c < w.length; c+=stride) {
                        r[(a*100)+(b*10)+(c)] = x[a] + y[b] + w[c];
                    }
                }
            }
        break;
        case 3:
            for (int a = index; a < x.length; a+=stride) {
                for (int b = index; b < y.length; b+=stride) {
                    for (int c = index; c < z.length; c+=stride) {
                        q[(a*100)+(b*10)+(c)] = x[a] + y[b] + w[c];
                    }
                }
            }
            for (int a = index; a < x.length; a+=stride) {
                for (int b = index; b < y.length; b+=stride) {
                    for (int c = index; c < w.length; c+=stride) {
                        t[(a*100)+(b*10)+(c)] = x[a] + y[b] + w[c];
                    }
                }
            }
        break;
    }
}

Answer 1

在 cuda 中，每个线程都运行您的内核。如果您希望线程做不同的事情，您必须（以某种方式）依赖于threadIdx 和/或blockIdx。

您通过计算index 来做到这一点。内核中的每个线程都有不同的index。现在你必须将你的索引映射到内核应该做的工作。所以你必须将每个index 映射到一个或多个(a,b,c) 的三元组。

您当前的映射类似于：

index -> (index+i*stride,index+j*stride,index+k*stride)

我不相信这是你的意图。