如果正确阻止，CUDA 代码不会处理

Question

卡在 //step 5 正下方的 if 块，问题是代码不会进入给定的 if 块或之后。在开始生成并行代码的任务之前，我需要弄清楚如何解决这个特定问题。如果您运行代码，您将看到一个打印语句指示“one”的值，另外两个指示“i”和“j”的值。在 if 块开始后，没有其他打印语句被命中。结果我被卡住了，我知道这是一个特定的问题，但是，我似乎无法确定它的原因。

任何帮助表示赞赏！ 提前致谢！

输入文件样本。

>386.fasta.screen.Contig1
GAGTTTGATCCTGGCTCAGAATCAACGCTGGCGGCGCGCTTAACACATGC
AAGTCGAACGAGAAAGTGGAGCAATCCATGAGTACAGTGGCGTACGGGTG
AGTAACACGTGGGTAATCTACCTCTTAGTGGGGAATAACTTTGGGAAACC
GAAGCTAATACCGCATAAGCTCGAGAGAGGAAAGCAGCAATGCGCTGAGA
GAGGAGCCCGCGGCCGATTAGCTAGTTGGCAGGGTAAAAGCCTACCAAGG
CAGAGATCGGTAGCCGGCCTGAGAGGGCACACGGCCACACTGGCACTGAA
ACACGGGCCAGACTCCTACGGGAGGCAGCAGTGGGGAATCTTGCACAATG
GGGGCAACCCTGATGCAGCGACGCCGCGTGAGCGATGAAGCCCTTCGGGG
TGTAAAGCTCTTTCGTCAGGGAAGATAGTGACGGTACCTGGAGAAGCAGC
TGCGGCTAACTACGTGCCAGCAGCCGCGGTAATACGTAGGCAGCGAGCGT
TGTTCGGAGTTACTGGGCGTAAAGGGTGTGTAGGCGGTTGTTTAAGTTTG
GTGTGAAATCTCCCGGCTCAACTGGGAGGGTGCGCCGAATACTGAGCGAC
TAGAGTGCGGGAGAGGAAAGTGGAATTCCTGGTGTAGCGGTGAAATGCGT
AGATATCAGGAGGAACACCGGTGGTGTAGACGGCTTTCTGGACCGTAACT
GACGCTGAGACACGAAAGCGTGGGTAGCAAACAGGATTAGATACCCTGGT
AGTCCACGCCCTAAACGATGCATATTTGGTGTGGGCAGTTCATTCTGTCC
GTGCCGGAGCTAACGCGTTAAATATGCCGCCTGGGGAGTACAGTCGCAAG
GCTGAAACTCAAAGGAATTGACGGGGGCCCGCACAAGCGGTGGAGCATGT
GGTTTAATTCGACGCAACGCGAAGAACCTTACCTGGGCTCGAACGGCTTC
CCAACGCCGGTAGAAATATCGGTACCCCGCAAGGGGGTGGAATCGAGGTG
CTGCATGGCTGTCGTCAGCTCGTGTCGTGAGATGTTGGGTTAAGTCCCGC
AACGAGCGCAACCCTTGTCCTGTGTTGCCATGCCGCAAGGCGGCACTCGC
AGGAGACCGCCAGCGATAAGCTGGAGGAAGGTGGGGATGACGTCAAGTCC
TCATGGCCTTTATGTCCAGGGCTACACACGTGCTACAATGGCCGGTACAA
AGCGTCGCTAACCTGCGAAGGGGAGCCAATCGCAAAAAACCGGTCTCAGT
TCGGATTGCAGGCTGCAACCCGCCTGCATGAAGCTGGAATCGCTAGTAAT
GGCAGATCAGCACGCTGCCGTGAATACGTTCCCGGGCCTTGTACACACAT

/********************************
Based on code by:
Lorenzo Seidenari (sixmoney@virgilio.it)
*********************************/

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

#define MAX_SEQUENCE_LENGTH 100000

int  n; 
int  m;
int levenshtein_distance(char *s,char*t);
int minimum(int a,int b,int c);

//-----------------------------------------------------------------------------
void cleanString(char string[]) {
  //Removes all spaces from string pointed to by "string", converts characters
  //to uppercase, and deletes a terminating newline character.
    int i, current;
    int length = strlen(string);

    current = 0;
    for(i=0;i<length;i++) {
        if(string[i]=='\n') {
            string[current++] = '\0';
            break;
        }
        else if(string[i]!=' ') {
            string[current++] = toupper(string[i]);
        }
    }
}
//-----------------------------------------------------------------------------
int importFASTA(char *filename, char *sequence) {
  //Reads a file, located at path specified by "filename", containing a FASTA
  //sequence. It finds the first full, complete sequence in the file, stores
  //it in "sequence", and returns the length of the sequence, or -1 on failure.
    FILE *fastaFile;
    char input[256];
  int readFlag; //set to 1 once a sequence has been read in
  int length;

  //open the file
  if((fastaFile = fopen(filename, "r")) == NULL) {
    return -1;
  }

  sequence[0] = '\0';

  //read the full first sequence, discarding unnecessary headers
  readFlag=0;
  length = 0;
  while(fgets(input,256,fastaFile)!=NULL) {
    //is it a header or a comment?
    if(input[0]=='>' || input[0]==';') {
        if(readFlag) break;
        else continue;
    }
    else readFlag = 1;

    cleanString(input);
    length += strlen(input);

    strncat(sequence,input,MAX_SEQUENCE_LENGTH-length - 1);
  }
  //Add a terminatng null character, just in case
  sequence[length] = '\0';

  fclose(fastaFile);
  return length;
}


/****************************************/
/*Implementation of Levenshtein distance*/
/****************************************/

__global__ void levenshtein_distance(char *s,char*t, int one, int two)
/*Compute levenshtein distance between s and t*/
{
    //Step 1
    int k,i,j,cost,*d;
    int distance = 0;
    if(one!=0&&two!=0)
    {
        d=(int *)malloc((sizeof(int))*(two+1)*(one+1));
        two++;
        one++;
        //Step 2    
        for(k=0;k<one;k++){
            d[k]=k;
        }
        for(k=0;k<two;k++){
            d[k*one]=k;
        }
        //Step 3 and 4  
        for(i=1;i<one;i++){
            for(j=1;j<two;j++)
            {
                //Step 5
                printf("%d  %d  %d\n", one, i, j);
                if(s[i-1]==t[j-1]){
                    cost=0;
                    printf("%d  %d  %d\n", one, i, j);
                }
                else{
                    cost=1;
                    printf("%d  %d  %d\n", one, i, j);
                }
                printf("%d  %d  %d\n", one, i, j);
                //Step 6
                int min = d[(j-1)*one+i]+1;
                if (d[j*one+i-1]+1 < min)
                    min = d[j*one+i-1]+1;
                if (d[(j-1)*one+i-1]+cost < min)
                    min = d[(j-1)*one+i-1]+cost;
                d[j*one+i] = min;        
            }
            distance=d[one*two-1];
            free(d);
            printf("%d\n", distance);
        }
    }
        else
            printf ("-1");
}

int main(int argc, char *argv[]) {
    char A[MAX_SEQUENCE_LENGTH+1];
    char B[MAX_SEQUENCE_LENGTH+1];

    if(argc < 3) {
        printf("Usage: new_edit_distance <sequence1> <sequence2>\n");
        printf("<sequence1>: file containing the first sequence, FASTA format\n");
        printf("<sequence2>: file containing the second sequence, FASTA format\n");
        return EXIT_FAILURE;
    }

    n = importFASTA(argv[1],A);
    m = importFASTA(argv[2],B);

    levenshtein_distance<<<1, 1>>>(A,B, n, m);
    cudaDeviceSynchronize();
    printf ("%s\n", cudaGetErrorString(cudaGetLastError()));

    return EXIT_SUCCESS;
}

Answer 1

我现在明白了。您将直接的串行 C/C++ 代码放入内核中，打算将该内核作为单线程运行，然后想从那里继续。

这个想法是有道理的，但是您错过了关于 CUDA 和 GPU 的一个关键事实：它们不能直接访问主机内存。

所以当你这样设置 A 和 B 时：

char A[MAX_SEQUENCE_LENGTH+1];
char B[MAX_SEQUENCE_LENGTH+1];
....
n = importFASTA(argv[1],A);
m = importFASTA(argv[2],B);

这些是存在于主机内存中的普通变量。 GPU（普通 CUDA）代码不能直接访问主机内存。因此，当您将这些指针传递给这样的内核时：

levenshtein_distance<<<1, 1>>>(A,B, n, m);

GPU 代码将尝试取消引用那些 A 和 B 指针并且会出错（未指定的启动失败）。

每个 CUDA 程序都有以下基本顺序：

1. 将数据复制到 GPU
2. 在 GPU 上执行计算
3. 将结果复制回来

您尝试在没有步骤 1 的情况下执行步骤 2。它不起作用。

由于没有有效的输入文件，我无法运行您的程序，因此我将提出以下建议。我假设您对 CUDA 知之甚少或一无所知。尝试添加这样的行：

n = importFASTA(argv[1],A);              // no change
m = importFASTA(argv[2],B);              // no change

char *d_A, *d_B;                          // add this line
cudaMalloc(&d_A, MAX_SEQUENCE_LENGTH+1);  // add this line
cudaMalloc(&d_B, MAX_SEQUENCE_LENGTH+1);  // add this line

cudaMemcpy(d_A, A, MAX_SEQUENCE_LENGTH+1, cudaMemcpyHostToDevice); // add 
cudaMemcpy(d_B, B, MAX_SEQUENCE_LENGTH+1, cudaMemcpyHostToDevice); // add

levenshtein_distance<<<1, 1>>>(d_A,d_B, n, m); //modify parameters

n 和 m 不需要进行任何不同的处理，因为您通过值传递它们。

并将proper cuda error checking 添加到您的代码中。

编辑：经过进一步分析，很明显这个顺序是不正确的：

        distance=d[one*two-1];
        free(d);
        printf("%d\n", distance);
    }
}

在i 循环的每次迭代中，您都在释放d。这不可能是正确的。我建议你回到第一方，先让你的串行代码在普通的串行 C 代码中工作，然后再以这种方式将它放入 cuda 内核。如果你将free 语句移到i 循环之外，那么你的内核会运行很长时间。请注意，内核内printf 可轻松生成的输出量有限。

我不会再为您调试代码。首先让您的串行代码工作，然后想办法创建一个无需大量打印输出的内核。

最后的评论：我在上面说过你的方法是“合理的”。这意味着它可以工作，即产生与主机上执行的相同代码相同的行为。它确实不意味着它会运行得很快。这不是您从 GPU 中获得加速的方式（运行单个线程的单个块）。我假设您已经根据您的评论“如何在开始生成并行代码的任务之前解决这个特定问题”知道这一点。但我认为免责声明还是合适的。