CUDA：嵌套循环内核

Question

CUDA：嵌套循环内核

5

我有一些代码想要转换成CUDA核函数。请看：

    for (r = Y; r < Y + H; r+=2)
    {
        ch1RowSum = ch2RowSum = ch3RowSum = 0;
        for (c = X; c < X + W; c+=2)
        {
            chan1Value = //some calc'd value
                            chan3Value = //some calc'd value
            chan2Value = //some calc'd value
            ch2RowSum  += chan2Value;
            ch3RowSum  += chan3Value;
            ch1RowSum  += chan1Value;
        }
        ch1Mean += ch1RowSum / W;
        ch2Mean += ch2RowSum / W;
        ch3Mean += ch3RowSum / W;
    }

这个问题应该分成两个内核，一个用于计算行求和，另一个用于计算均值。但是，由于循环索引不从零开始且不以N结束，我该如何处理呢？

- Derek

尝试只选择一个问题会使选择正确答案变得困难。但是，对于你的第二个问题...具体回答很难，但我认为一旦你在开发内核方面进一步深入，你就会明白。 - jmilloy

你应该使用类似 H 个块和每个块 W 个线程的配置来启动你的内核。然后，您将在内核中通过blockIdx和threadIdx的值计算r和c。可以按任何方式计算r和c...我尝试在下面的答案中解释这个问题... - jmilloy

它看起来像是两个问题，但我不确定如果我把它写成两个问题，上下文是否会存在。 - Derek

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- jmilloy · Accepted Answer

假设您有一个计算三个值的内核。您配置中的每个线程将为每个(r,c)对计算三个值。

__global__ value_kernel(Y, H, X, W)
{
    r = blockIdx.x + Y;
    c = threadIdx.x + W;

    chan1value = ...
    chan2value = ...
    chan3value = ...
}

我不相信你能在上述内核中完全并行地计算总和，至少不能像你上面那样使用 +=。如果每个块（行）中只有一个线程执行总和和平均值的计算，可以将所有内容放在一个内核中，如下所示...

__global__ both_kernel(Y, H, X, W)
{
    r = blockIdx.x + Y;
    c = threadIdx.x + W;

    chan1value = ...
    chan2value = ...
    chan3value = ...

    if(threadIdx.x == 0)
    {
        ch1RowSum = 0;
        ch2RowSum = 0;
        ch3RowSum = 0;

        for(i=0; i<blockDim.x; i++)
        {
            ch1RowSum += chan1value;
            ch2RowSum += chan2value;
            ch3RowSum += chan3value;
        }

        ch1Mean = ch1RowSum / blockDim.x;
        ch2Mean = ch2RowSum / blockDim.x;
        ch3Mean = ch3RowSum / blockDim.x;
    }
}

但最好使用第一个值内核，然后为两个总和和平均数使用第二个内核... 可以进一步并行化下面的内核，如果它是单独的，则可以在准备好时专注于它。

__global__ sum_kernel(Y,W)
{
    r = blockIdx.x + Y;

    ch1RowSum = 0;
    ch2RowSum = 0;
    ch3RowSum = 0;

    for(i=0; i<W; i++)
    {
        ch1RowSum += chan1value;
        ch2RowSum += chan2value;
        ch3RowSum += chan3value;
    }

    ch1Mean = ch1RowSum / W;
    ch2Mean = ch2RowSum / W;
    ch3Mean = ch3RowSum / W;
}