c - 如何修改此组合算法以在启用 cuda 的 gpu 上并行运行？

Question

我目前有一个 c 程序可以生成字符串的所有可能组合。请注意组合，而不是排列。这是程序:

#include <stdio.h>
#include <string.h>
#include <stdlib.h>

//Constants
static const char set[] = "abcd";
static const int setSize = sizeof(set) - 1;

 void brute(char* temp, int index, int max){
     //Declarations
     int i;
     for(i = 0; i < setSize; i++){
         temp[index] = set[i];
         if(index == max - 1){
                printf("%s\n", temp);
         }
         else{
              brute(temp, index + 1, max);
         }
   }
}

void length(int max_len){
   //Declarations
   char* temp = (char *) malloc(max_len + 1);
   int i;

   //Execute
   for(i = 1; i <= max_len; i++){
         memset(temp, 0, max_len + 1);
         brute(temp, 0, i);
   }
   free(temp);
}


int main(void){
   //Execute
   length(2);
   getchar();
   return 0;
}

可以修改组合的最大长度；出于演示目的，它目前设置为 2。考虑到它当前的配置方式，程序输出

a, b, c, d, aa, ab, ac, ad, ba, bb, bc, bd, ca, cb, cc...

我已经设法将这个程序翻译成 cuda:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <string.h>
#include <stdlib.h>

//On-Device Memory
__constant__ char set_d[] = "adcd";
__constant__ int setSize_d = 4;

__device__ void brute(char* temp, int index, int max){
  //Declarations 
  int i;
  for(i = 0; i < setSize_d; i++){
      temp[index] = set_d[i];
      if(index == max - 1){
          printf("%s\n", temp);
      }
      else{
          brute(temp, index + 1, max);
      }
  }
}


__global__ void length_d(int max_len){

      //Declarations
  char* temp = (char *) malloc(max_len + 1);
  int i;
  //Execute
  for(i = 1; i <= max_len; i++){
      memset(temp, 0, max_len+1);
      brute(temp, 0, i);
  }
  free(temp);
}


int main()
{
    //Execute
    cudaSetDevice(0);

    //Launch Kernel
    length_d<<<1, 1>>>(2);
    cudaDeviceSynchronize();

    getchar(); //Keep this console open...
    return 0;
}

原始程序的 cuda 版本基本上是 c 程序的精确副本(注意它是用 -arch=sm_20 编译的。因此，printf 和其他主机功能在 cuda 环境中工作)。

我的目标是计算最大长度为 10 的 a-z、A-Z、0-9 和其他字符 的组合。既然如此，我希望这个程序在我的 gpu 上运行。就像现在一样，它没有利用并行处理——这显然违背了用 cuda 编写程序的全部目的。但是，除了将线程委托(delegate)给特定索引或起点之外，我不确定如何消除程序的递归性质。

欢迎任何建设性的意见。

此外，我在连续编译时偶尔会收到一条警告消息(这意味着它偶尔会出现):警告:无法静态确定入口函数“_Z8length_di”的堆栈大小。

我还没有查明问题所在，但我想我会发布它，以防有人在我之前确定原因。它正在 visual studio 2012 中编译。

注意:我发现这很有趣。就像现在的 cuda 程序一样，它向控制台的输出是周期性的——这意味着它会打印几十个组合、暂停、打印几十个组合、暂停等等。我还在其报告的 gpu 使用情况中观察到这种行为 - 它会周期性地从 5% 摆动到 100%。

Answer 1

我认为您不需要为此使用递归。 (我不会)。

从内核使用 printf 是 problematic for large amounts of output ；它并不是真正为此目的而设计的。内核中的 printf 消除了 GPU 可能具有的任何速度优势。我假设如果你正在测试这样一个大的 vector 空间，你的目标不是打印出每一个组合。即使那是您的目标，内核中的 printf 也不是正确的选择。

您将遇到的另一个问题是您所想到的整个 vector 空间的存储。如果您打算对每个 vector 进行一些处理，然后可以丢弃它，那么存储就不是问题。但是，对于长度为 n=10 且具有 k=62 或更多可能值(a..z、A..Z、0..9 等)的“数字”(元素)的 vector 空间，存储空间将非常庞大。它由 k^n 给出，因此在本例中将是 62^10 个不同的 vector 。如果每个数字都需要一个字节来存储它，那将超过 7 万亿千兆字节。所以这几乎表明整个 vector 空间的存储是不可能的。无论您要做什么工作，都必须即时完成。

鉴于上述讨论，this answer 应该拥有您需要的几乎所有内容。 vector 数字作为 unsigned int 处理，您可以在 unsigned int 和您的“数字”(即 a..z、A..Z、0. .9 等)在该示例中，对每个 vector 执行的函数正在测试数字总和是否匹配特定值，因此您可以替换内核中的这一行:

  if (vec_sum(vec, n) == sum) atomicAdd(count, 1UL);

无论您想对生成的每个 vector 应用什么函数和处理。您甚至可以在此处放置 printf，但对于较大的空间，输出将是零散且不完整的。