Passing **char to kernel

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi, i have some problem with passing **char to cuda kernel.

i want to print string and always empty string.

int main () {
  cudaError_t err = cudaSuccess;
  string line,y, line2, x;
  int i = 0, n,j;

  int *len;
  int *score;
  int *dscore;
  int *dlen;

  vector<string> sequence_id;
  vector<string>::iterator it;

  ifstream myfile ("prot_list.csv");
  if (myfile.is_open())
  {
    while ( getline (myfile,line) )
    { 
        sequence_id.push_back(line);
        //cout << line << "\n";
    }
    
    myfile.close();
  }
  //else cout << "Unable to open file";
  cin >> n;
  //cout << n << "\n";
  if(n == 0)
    n = sequence_id.size();
  int size = sequence_id.size();
    cout << size << "\n";
    const char **sequence = new const char*;
    const char **dsequence = new const char*;

     //sequence = (const char**)malloc(size * sizeof(const char));
     score    = (int*)malloc(((n) *(n)) * sizeof(int));
     len    = (int*)malloc(size * sizeof(int));
     
     //score = 10;
     cudaMalloc((void**)&dsequence, size * sizeof(const char*));
     cudaMalloc((void**)&dscore, ((n) *(n)) * sizeof(int));
     cudaMalloc((void**)&dlen,size * sizeof(int));

    for(it = sequence_id.begin(); it < sequence_id.end(); it++){
        ifstream myfile2 ("Fasta/" + *it + ".fasta");
        if (myfile2.is_open())
        {
            const char *tem;
            x = "";
            while ( getline (myfile2,line2) )
            {
            if(line2[0] != '>')  
                    x += line2;
            }
            tem = x.c_str();
            sequence[i] = tem;
            len[i] = x.length();
            //cout << *it << " : " <<  sequence[i] << "\n"; 
            myfile2.close();
        }
        //else cout << *it << " Unable to open file \n"; 
        i++;
    }

    err = cudaMemcpy(dsequence, sequence, size * sizeof(const char*), cudaMemcpyHostToDevice);
    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to copy vector B from host to device (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }
    err = cudaMemcpy(dscore, score, ((n) *  (n)) * sizeof(int), cudaMemcpyHostToDevice);
    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to copy vector B from host to device (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }
    err = cudaMemcpy(dlen, len,size * sizeof(int), cudaMemcpyHostToDevice);
    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to copy vector B from host to device (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }
    
    int threadsPerBlock = n;
    int blocksPerGrid = 1;
    for(j= 0; j < threadsPerBlock; j++){
   	    FillMatrix <<<blocksPerGrid, threadsPerBlock>>>(dsequence,dlen,n,j,dscore);
    }
    err = cudaMemcpy(score, dscore, ((n) *(n)) * sizeof(int), cudaMemcpyDeviceToHost);
    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to copy score from device to host (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }
   

    for(int i = 0; i< n ; i++){
      for(int j = 0; j < n; j++){
          cout << sequence_id.at(i) << "," << sequence_id.at(j) << "," << score[(i*(n)) + j]  << "\n";
      }
    }
    cudaFree(dsequence); cudaFree(dscore); cudaFree(dlen); 
    return 0;
}

this for kernel code

__global__ void FillMatrix(const char **sequence,int *s_length, int n, int a, int *score)
{
  int b = threadIdx.x + blockIdx.x*blockDim.x;
  const int rows = s_length[a],cols = s_length[b];
  score[(a*(n)) + b] = rows;
    printf("string %d = %s %d\n",s_length[b], sequence[a],b);
}
2 
const char **dsequence = new const char*;

you should initialize it to NULL. it has nothing common with your problem, you just leak this memory

the source of problem is that you don’t have a proper picture of pointers and data in your scenario. your text strings are contained in CPU RAM. entries of array “sequence”, also contained in CPU RAM, points to these strings. you copy these pointers into array dsequence, which is contained in GPU RAM. but the copied entries still contains addresses in CPU RAM. you should realize that these pointers are just invalid on GPU. for simplicity, imagine that all pointers on CPU are just numbers starting with 1000, and all pointers on GPU are numbers starting with 2000. When you try to dereference pointers in 1000…1999 range, GPU just goes crazy (it can raise exception or just do Weird Things)

So, first you need to copy data pointed by each sequence[i], not only pointers themself

The second problem in your code is that you create temporary string “x” to store contents read from file, and then you save pointer to data inside “x”, retrieved with x.c_str(). But on the next cycle, you overwrite “x” variable, and the memory area allocated for “x” contents may be freed or reused. You need to better understand how STL containers (in this case std::string) manage their memory

For your goals, i suggest to completely avoid using STL, and use plain C malloc/free machinery

if you want to EFFICICIENTLY copy large string data to GPU, alloc large array of chars and then manually alloc space there for individual strings, and then copy entire buffer to GPU. such code is not short, but that’s the nature of low-level C programming

3

thanks for your explaination. i am so sorry i am newbie here.

i have tried with this code for allocation memory. but i am getting this error

"error: argument of type “const char *” is incompatible with parameter of type “void *” "

for(int i=0;i< size; i++){
      cudaMalloc((void**)&dsequence[i], (len[i]+1) * sizeof(const char));
    }
    cudaMalloc((void**)&dscore, ((n) *(n)) * sizeof(int));
    cudaMalloc((void**)&dlen,size * sizeof(int));

    cudaMemcpy(dsequence, sequence, size * sizeof(const char), cudaMemcpyHostToDevice);
    for(int i=0;i< size; i++){
      err = cudaMemcpy(dsequence[i], sequence[i], (len[i]+1) * sizeof(const char), cudaMemcpyHostToDevice);
      if (err != cudaSuccess)
      {
          fprintf(stderr, "Failed to copy vector B from host to device (error code %s)!\n", cudaGetErrorString(err));
          exit(EXIT_FAILURE);
      }
    }
4

sorry, but i’m not your teaching assistant :) i can help you to understand C/CUDA but now you seems just to make random edits to the code

i will repeat my advice from other topic:

in general, C/C++ books should have a chapter about pointers/arrays. i suggest to find [a book with] exercises on this topic and do them all. after all, it’s the core hard topic of C, so learning this topic as deep as possible will make your future life with C simpler. sorry for unasked advices

5

I am so sorry. but your suggestion is copy data pointed by each sequence[i], not only pointers themself. And i think i do for this.

Sorry if i wrong.

6

you haven’t said where you get this error. do you expect that i will check each line in your code? do you expect that i will do it at all?

if you want to copy each sequence[i] individually, you should do it inside the loop reading input data. because on next loop step, these data are already gone

7

oh i am so really sorry, I never thought like that.
i am getting this error at line 9

cudaMemcpy(dsequence[i], sequence[i], (len[i]+1) * sizeof(const char), cudaMemcpyHostToDevice);
8

it seems that you are applying ‘const’ modifier in a random way. “const char **dsequence” declaration means that you cannot overwrite dsequence[i] contents, and compiler successfully cathed the error

either learn how to use “const” properly, or use “(void*)(dsequence[i])” to suppress this error, or avoid using “const” modifiers in this program at all - do you really need them or just copied from somewhere?