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

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

也许有一种方法可以通过启动两个内核来绕过逐元素模计算，一个使用 x = cosf(sqrtf(x)) 处理偶数索引，另一个使用 x = sqrtf( logf(x))?

是的。

为了缩短执行时间，

我怀疑它会。

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

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），通过“并发”运行它们可能会稍微提高性能。

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

• 提供合并的负载/存储
• 避免线程间的条件行为

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

__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
}

这有一些隐含的假设：

1. A 指针针对float2 访问正确对齐。例如，如果 A 由 cudaMalloc （直接）返回，则可以满足此要求。
2. 要处理的元素总数 (N) 是偶数（在这种简单的情况下，为 2048 或更少。）