对输入数组的不同部分执行不同的内核

Question

假设我有一个简单的内核，它对输入数组的每个元素应用两种不同的数学运算。要应用的操作是根据数组中元素位置的模来选择的：

__global__ void kernel(float* A, int N)
{
    int i = threadIdx.x;    

    if (i < N) {
        float x = A[i];

        if (i % 2 == 0)       x = cosf(sqrtf(x));
        else if (i % 2 == 1)  x = sqrtf(logf(x));

        A[i] = x;
    }    
}

int main()
{
    // some code...

    // Kernel invocation with N threads
    kernel<<<1, N>>>(A, N);

    // ...some more code
}

为了缩短执行时间，是否有一种方法可以通过启动两个内核来绕过逐元素模计算，一个使用 x = cosf(sqrtf(x)) 处理偶数索引，另一个使用 x = sqrtf(logf(x))?

Answer 1

is there perhaps a way to by-pass the element-wise modulo computation by instead launching two kernels, one that processes even indices with x = cosf(sqrtf(x)) and one that processes odd indices with x = sqrtf(logf(x))?

是的。

To improve execution time,

我怀疑它会。

如果你想要一个交替的（even/odd索引）选择，一个非常低成本的操作就是：

if (index & 1)  
   //perform odd function
else
   //perform even function

我认为如果你使用那种低成本的选择，将操作分成两个内核是不太可能的。

但是，如果您想这样做，那将相当简单：

__global__ void even_kernel(float* A, int N)
{
    int i = threadIdx.x*2;    

    if (i < N) {
        float x = A[i];
        x = cosf(sqrtf(x));
        A[i] = x;
    }    
}

__global__ void odd_kernel(float* A, int N)
{
    int i = threadIdx.x*2+1;    

    if (i < N) {
        float x = A[i];
        x = sqrtf(logf(x));
        A[i] = x;
    }    
}

根据您提供的代码，您可以像这样调用它们：

int main()
{
    // some code...
    cudaStream_t s1, s2;
    cudaStreamCreate(&s1);
    cudaStreamCreate(&s2);
    // Kernel invocation with N threads
    odd_kernel<<<1, N/2, 0, s1>>>(A, N);
    even_kernel<<<1, N/2, 0, s2>>>(A, N);

    // ...some more code
}

流处理并不是真正必要的，但对于如此小的网格（每个一个块，N 最多为 1024），运行它们“并发”可能会在性能上有一些小的改进。

正如评论中所建议的，我们可能可以通过让每个线程处理两个元素来解决可能针对您的原始代码的两个性能批评。如果我们为此使用矢量化负载，那么我们可以：

提供合并 loads/stores
避免线程间的条件行为

那个内核可能看起来像这样：

__global__ void kernel(float* A, int N)
{
    int i = threadIdx.x;
    float2* A2 = reinterpret_cast<float2 *>(A);    

    if (i < N/2) {
        float2 x = A2[i];

        x.x = cosf(sqrtf(x.x));
        x.y = sqrtf(logf(x.y));

        A2[i] = x;
    }    
}

int main()
{
    // some code...

    // Kernel invocation with N threads
    kernel<<<1, N/2>>>(A, N);

    // ...some more code
}

这有一些隐含的假设：

A 指针已正确对齐以供 float2 访问。例如，如果 cudaMalloc.

A

要处理的元素总数 (N) 是偶数（在这种简单的情况下为 2048 或更少。）

对输入数组的不同部分执行不同的内核

Execute different kernel for different parts of input array

cuda