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

Question

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

感谢任何帮助!提前致谢!

输入文件示例。

>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 中获得加速的方式(运行单个线程的单个 block )。我假设您已经根据您的评论“在开始生成并行代码的任务之前解决这个特定问题”而知道这一点。但我认为免责声明无论如何都是合适的。