如何将字符串矩阵从C++传递到Cuda内核

How to pass string matrix to Cuda kernel from C++

本文关键字：Cuda 内核 C++ 字符串      更新时间：2023-10-16

问题：

我在C++中有一个矩阵，里面充满了字符串，我想把它传递给cuda内核函数。我知道CUDA不能处理字符串，所以经过一些研究，我尝试了下面列出的一些解决方案。

尝试次数：

  1. 在C++中定义一个指针数组，每个单元格包含一个指针字符（为了简单起见，tmp[i]用前面引用的矩阵中包含的字符串填充）

    C++部分

     char *tmp[3];
 int text_length, array_length;
 text_length = 4;
 array_length = 3;
 tmp[0] = (char*) malloc(text_length*sizeof(char));
 tmp[1] = (char*) malloc(text_length*sizeof(char));
 tmp[2] = (char*) malloc(text_length*sizeof(char));
 tmp[0] = "some";
 tmp[1] = "rand";
 tmp[2] = "text";
 char *a[3];
 for(int i=0;i<array_length;i++)
 {
   cudaMalloc((void**) &a[i],text_length*sizeof(char));
   cudaMemcpy(&a[i],&tmp[i],text_length*sizeof(char),cudaMemcpyHostToDevice);
 }
 func<<<blocksPerGrid, threadsPerBlock>>>(a);

    CUDA部分

     __global__ void func(char* a[]){
 for(int i=0;i<3;i++)
   printf("value[%d] = %s n",i, a[i]);
 }

    输出

     value[0] = (null)
 value[1] = (null)
 value[2] = (null)

  2. 将充满字符串的矩阵扩展到一个char指针，并将其传递给cuda内核，然后尝试检索字符串（再次在C++中简化代码）

    C++部分

     char *a;
 int index[6];
 a = "somerandtext";
 index[0] = 0; // first word start
 index[1] = 3; // first word end
 index[2] = 4; // same as first word 
 index[3] = 7;
 index[4] = 8;
 index[5] = 1;
 func<<<blocksPerGrid, threadsPerBlock>>>(a,index);

    CUDA部分

     __global__ void func(char* a,int index[]){
 int first_word_start = index[0];
 int first_word_end = index[1];
 // print first word
 for(int i=first_word_start;i<=first_word_end;i++)
   printf("%c",a[i]);
 }

    输出

     no output produced

我已经尝试了很多其他解决方案，但没有一个适合我……这个问题也可以重新提出，问：我如何将"n"个字符串传递到cuda内核，并在那里打印（和比较）所有字符串（请记住，我不能传递"n"变量）。

您所展示的代码中有许多是完整的，您遗漏的内容可能很重要。如果你展示了完整的代码，你会让其他人更容易帮助你。此外，每当你在使用CUDA代码时，最好使用正确的CUDA错误检查，这通常会让你发现哪些代码不起作用（我怀疑这可能有助于你的第二次尝试）。此外，使用cuda-memcheck运行代码通常很有指导意义。

在第一次尝试中，您遇到了CUDA和嵌套指针的经典问题（a是指向指针数组的指示器）。这个问题几乎在其他数据结构中隐藏指针的任何时候都会发生。将这样的数据结构从主机复制到设备需要"深度复制"操作，该操作有多个步骤。要了解更多信息，请搜索"CUDA 2D数组"（我认为规范答案是talonmies在这里给出的答案），或者在这里和这里查看我的答案。

还要注意的是，使用CUDA6，如果能够使用统一内存，"深度复制"在概念上对程序员来说会容易得多。

你的第二次尝试似乎是沿着一条"压平"你的2D或指向char的ponter数组的指针的路径前进。这是深度复制"问题"的典型解决方案，可以降低代码复杂性，并可能提高性能。这里有一个完整的例子，融合了你第一次和第二次尝试的想法，这似乎对我有效：

$ cat t389.cu
#include <stdio.h>
 __global__ void func(char* a, int *indexes, int num_strings){

 for(int i=0;i<num_strings;i++){
   printf("string[%d]: ", i);
   for (int j=indexes[2*i]; j < indexes[2*i+1]; j++)
     printf("%c", a[j]);
   printf("n");
 }
}
int main(){
 int max_text_length, num_str;
 num_str = 3;
 char *tmp[num_str];
 max_text_length = 12;
 tmp[0] = (char*) malloc(max_text_length*sizeof(char));
 tmp[1] = (char*) malloc(max_text_length*sizeof(char));
 tmp[2] = (char*) malloc(max_text_length*sizeof(char));
 tmp[0] = "some text";
 tmp[1] = "rand txt";
 tmp[2] = "text";
 int stridx[2*num_str];
 int *d_stridx;
 stridx[0] = 0;
 stridx[1] = 9;
 stridx[2] = 9;
 stridx[3] = 17;
 stridx[4] = 17;
 stridx[5] = 21;
 char *a, *d_a;
 a = (char *)malloc(num_str*max_text_length*sizeof(char));
 //flatten
 int subidx = 0;
 for(int i=0;i<num_str;i++)
 {
   for (int j=stridx[2*i]; j<stridx[2*i+1]; j++)
     a[j] = tmp[i][subidx++];
   subidx = 0;
 }
 cudaMalloc((void**)&d_a,num_str*max_text_length*sizeof(char));
 cudaMemcpy(d_a, a,num_str*max_text_length*sizeof(char),cudaMemcpyHostToDevice);
 cudaMalloc((void**)&d_stridx,num_str*2*sizeof(int));
 cudaMemcpy(d_stridx, stridx,2*num_str*sizeof(int),cudaMemcpyHostToDevice);

 func<<<1,1>>>(d_a, d_stridx, num_str);
 cudaDeviceSynchronize();
}
$ nvcc -arch=sm_20 -o t389 t389.cu
$ cuda-memcheck ./t389
========= CUDA-MEMCHECK
string[0]: some text
string[1]: rand txt
string[2]: text
========= ERROR SUMMARY: 0 errors
$
相关文章：