c - 使用CUDA内核进行数组操作

Question

我正在编写 CUDA 内核，线程正在执行以下任务:

for example i have array of [1, 2, 3, 4] then I want answer [12, 13, 14, 23, 24, 34]

假设我有一个包含 n 个整数的数组，并且有两个索引 i 和 j。

C 语言的简单解决方案是:

k=0;
for (i = 0; i < n - 1; i++)
    for(j = i+1; j < n-1 ; j++)
       { new_array[k] = array[i]*10 + array[j];
          k++;
        }

在 CUDA 中我尝试过运气:

for(i = threadIdx.x + 1; i < n-1; i++ )
    new_array[i] = array[threadIdx.x] * 10 + array[i];

但我认为这不是完全正确或最佳的方法。谁能提出更好的建议吗？

Answer 1

我假设您想要移植到 CUDA 的代码如下:

#include <stdio.h>
#define N 7

int main(){

    int array[N] = { 1, 2, 3, 4, 5, 6, 7};
    int new_array[(N-1)*N/2] = { 0 };

    int k=0;
    for (int i = 0; i < N; i++)
        for(int j = i+1; j < N; j++)
        { 
            new_array[k] = array[i]*10 + array[j];
            k++;
        }

    for (int i = 0; i < (N-1)*N/2; i++) printf("new_array[%d] = %d\n", i, new_array[i]);

return 0;

}

您可能希望注意，您可以将内部循环重铸为

for (int i = 0; i < N; i++)
    for(int j = i+1; j < N; j++)
        new_array[i*N+(j-(i+1))-(i)*(i+1)/2] = array[i]*10 + array[j];

这将避免显式定义索引变量 k直接使用索引i*N+(j-(i+1))-(i)*(i+1)/2 。这样的观察是有用的，因为如果你解释索引 i和j作为移植代码中的线程索引，那么您将在 2d 之间建立映射。线程索引和访问 __global__ 中的目标数组所需的索引您必须定义的函数。

因此，移植的代码为

#include <stdio.h>
#define N 7

__global__ void kernel(int* new_array_d, int* array_d) {

    int i = threadIdx.x;
    int j = threadIdx.y;

    if (j > i) new_array_d[i*N+(j-(i+1))-(i)*(i+1)/2] = array_d[i]*10 + array_d[j];
}

int main(){

    int array[N] = { 1, 2, 3, 4, 5, 6, 7};
    int new_array[(N-1)*N/2] = { 0 };

    int* array_d;       cudaMalloc((void**)&array_d,N*sizeof(int));
    int* new_array_d;   cudaMalloc((void**)&new_array_d,(N-1)*N/2*sizeof(int));

    cudaMemcpy(array_d,array,N*sizeof(int),cudaMemcpyHostToDevice);

    dim3 grid(1,1);
    dim3 block(N,N);
    kernel<<<grid,block>>>(new_array_d,array_d);

    cudaMemcpy(new_array,new_array_d,(N-1)*N/2*sizeof(int),cudaMemcpyDeviceToHost);

    for (int i = 0; i < (N-1)*N/2; i++) printf("new_array[%d] = %d\n", i, new_array[i]);

    return 0;
}

请添加您自己的 CUDA 错误检查 What is the canonical way to check for errors using the CUDA runtime API? 。此外，您可能希望将上述 CUDA 代码扩展到非统一大小的 block 网格的情况。