使用 CUDA 生成排列

Question

我正在阅读有关 CUDA 的内容，我尝试实现一个简单的代码来创建 array {a,b,c,d} 的所有可能排列，但我不知道如何实现 CUDA 方式（因为我阅读的所有示例都是形成a[blockIdx.x] = b[blockIdx.x] + c[blockIdx.x])。

任何帮助将不胜感激。

Answer 1

下面是 CUDA 中并行排列生成器的一个有点幼稚的实现。该示例旨在生成ABCD 的所有可能排列。

由于可以通过将第一个符号固定为 X 并附加其余符号的所有可能排列来生成所有可能的排列，然后将第一个符号更改为 Y 并再次执行上述过程，代码背后的简单想法是指定4线程来完成这项工作，每个线程引用不同的初始符号。

第一个符号之后的排列以规范的方式进行评估，即通过递归。

显然，下面的代码可以变得更通用，或许还可以改进，但它应该会给你一个初步的大致思路。

#include <stdio.h>

inline void GPUassert(cudaError_t code, char * file, int line, bool Abort=true)
{
    if (code != 0) {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),file,line);
        if (Abort) exit(code);
    }       
}

#define GPUerrchk(ans) { GPUassert((ans), __FILE__, __LINE__); }

__host__ __device__ void swap(char *x, char *y)
{
    char temp;
    temp = *x;
    *x = *y;
    *y = temp;
}

__device__ void permute_device(char *a, int i, int n, int tid, int* count) 
{
    if (i == n) { 
        char b[4]; char* c = a - 1; 
        b[0] = c[0]; b[1] = c[1]; b[2] = c[2]; b[3] = c[3];
        printf("Permutation nr. %i from thread nr. %i is equal to %s\n", count[0], tid, c); count[0] = count[0] + 1; 
    }
    else
    {
        for (int j = i; j <= n; j++)
        {
            swap((a+i), (a+j));
            permute_device(a, i+1, n, tid, count);
            swap((a+i), (a+j)); //backtrack
        }
    }
} 

__global__ void permute_kernel(char* d_A, int size) {

    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int count[1]; count[0] = 0;

    char local_array[4];

    for (int i=0; i<size; i++) local_array[i] = d_A[i];

    swap(local_array+threadIdx.x,local_array);
    permute_device(local_array+1,0,2,tid,count);

}

int main()
{
    char h_a[] = "ABCD";

    char* d_a; cudaMalloc((void**)&d_a,sizeof(h_a));
    GPUerrchk(cudaMemcpy(d_a, h_a, sizeof(h_a), cudaMemcpyHostToDevice));

    printf("\n\n Permutations on GPU\n");
    permute_kernel<<<1,4>>>(d_a, 4);
    GPUerrchk(cudaPeekAtLastError());
    GPUerrchk(cudaDeviceSynchronize());

    getchar();
    return 0;
}

Answer 2

我想知道你的组合是不是成对的。

1. 想想数组的长度。(n)
2. 想想组合复合体。 （n^2，成对）
3. 选择一种并行计算的方法。 （以block id为add offset，thread id为base one.array(threadid) + array(threadid+offset)为例）

所以你的内核应该这样写：

template<typename T>
__global__ void 
combination(T* in, T* out) {
  int tid = threadId.x;
  int bid = blockId.x+1;
  out[tid+bid*blockDim.x] = in[tid]+in[(tid+bid)%arrlen];     //0+1, 1+2,..., 0+2, 1+3
}

您可以使用gridDim(arrlen) 和blockDim(arrlen) 调用内核。