c - 为什么只有少量线程启动？

Question

该代码用于并行化字符串匹配算法(暴力)。为什么只启动 5 个线程而不是 17 个？

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

__global__ void pattern_search(char* d_txt,char* d_pat,int* d_result,int N,int M){
    int id=threadIdx.x+blockIdx.x*blockDim.x;
    if(id<=(N)){
        int j=0;
        for(j=0;j<M;j++){
            if(d_txt[id+j]!=d_pat[j]){
                break;
            }
        }
        //if(j==M){
            d_result[id]=id;
        //}

    }
}

int main(){
    char txt[]="AABAACAADAABAAABAA";
    char pat[]="AABA";

    int N=strlen(txt);
    int M=strlen(pat);
    char* d_pat;
    cudaMalloc((void **)&d_pat,M);
    char* d_txt;
    cudaMalloc((void **)&d_txt,N);

    int result[N];
    for(int i=0;i<N;i++){
        result[i]=0;
    }

    int* d_result;
    cudaMalloc((void**)&d_result,N);
    cudaMemcpy(d_txt,txt,N,cudaMemcpyHostToDevice);
    cudaMemcpy(d_pat,pat,M,cudaMemcpyHostToDevice);
    cudaMemcpy(d_result,result,N,cudaMemcpyHostToDevice);
    pattern_search<<<1,50>>>(d_txt,d_pat,d_result,N,M);

    cudaMemcpy(result,d_result,N,cudaMemcpyDeviceToHost);

    for(int k=0;k<N;k++){

        printf("pattern found at:%d\n",result[k]);
    }
}

Answer 1

1. 每当您在使用 CUDA 代码时遇到问题时，演示和使用 proper cuda error checking 是一个很好的做法。 并使用 cuda-memcheck 运行您的代码。如果您使用 cuda-memcheck 运行此代码，它会报告可能具有指导意义的错误。

2. 不清楚为什么您会询问 17 个线程。在 CUDA 中，启动的线程数可以从内核启动中推断出来，简单来说，它是内核启动配置中前两个数字的乘积( <<<...>>> ):

   pattern_search<<<1,50>>>(d_txt,d_pat,d_result,N,M);
   

   因此在这种情况下，它应该启动 50 个线程。即使您的问题仅限于 N ，号码N因为您发布的代码是 18，而不是 17。

3. cudaMalloc ，如主机端malloc ，以字节为单位分配内存。因此，对于这种情况，这样的用法是不正确的:

   cudaMalloc((void**)&d_result,N);
   

   相反，你应该这样做:

   cudaMalloc((void**)&d_result,N*sizeof(int));
   

   因为在这种情况下您想要存储 N尺寸数量int 。此错误可通过 cuda-memcheck 发现这将报告无效 __global__由于此分配错误而写入。

4. 您的 cudaMemcpy 上也出现类似的问题调用，它也在字节上运行(就像主机 memcpy )。而不是这个:

   cudaMemcpy(d_result,result,N,cudaMemcpyHostToDevice);
   

   我们想要这个:

   cudaMemcpy(d_result,result,N*sizeof(int),cudaMemcpyHostToDevice);
   

   并且需要对内核之后的调用进行类似的更正。

5. 您的内核存在索引越界错误:

   if(id<=(N)){
       int j=0;
       for(j=0;j<M;j++){
           if(d_txt[id+j]!=d_pat[j]){
   

   上面的代码将允许 for 循环索引超出数组的末尾 d_txt长度限制为 N 。为了解决这个问题，我们可以将循环行为限制为仅在有足够的“索引空间”用于 j 循环的完整迭代时才运行:

   if((id+M)<=(N)){
       int j=0;
       for(j=0;j<M;j++){
           if(d_txt[id+j]!=d_pat[j]){
   

   (并且可能还有许多其他方法来解决此问题。)此错误可能已被无效的 __global__ 发现。读取cuda-memcheck报道.

以下代码解决了上述问题，并且运行时没有错误:

$ cat t964.cu
#include<stdio.h>
#include<string.h>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


__global__ void pattern_search(char* d_txt,char* d_pat,int* d_result,int N,int M){
    int id=threadIdx.x+blockIdx.x*blockDim.x;
    if((id+M)<=(N)){
        int j=0;
        for(j=0;j<M;j++){
            if(d_txt[id+j]!=d_pat[j]){
                break;
            }
        }
        //if(j==M){
            d_result[id]=id;
        //}

    }
}

int main(){
    char txt[]="AABAACAADAABAAABAA";
    char pat[]="AABA";

    int N=strlen(txt);
    int M=strlen(pat);
    char* d_pat;
    cudaMalloc((void **)&d_pat,M);
    char* d_txt;
    cudaMalloc((void **)&d_txt,N);

    int result[N];
    for(int i=0;i<N;i++){
        result[i]=0;
    }

    int* d_result;
    cudaMalloc((void**)&d_result,N*sizeof(int));
    cudaMemcpy(d_txt,txt,N,cudaMemcpyHostToDevice);
    cudaMemcpy(d_pat,pat,M,cudaMemcpyHostToDevice);
    cudaMemcpy(d_result,result,N*sizeof(int),cudaMemcpyHostToDevice);
    cudaCheckErrors("1");
    pattern_search<<<1,50>>>(d_txt,d_pat,d_result,N,M);

    cudaMemcpy(result,d_result,N*sizeof(int),cudaMemcpyDeviceToHost);
    cudaCheckErrors("2");
    for(int k=0;k<N;k++){

        printf("pattern found at:%d\n",result[k]);
    }
}

您没有准确指出您期望的输出，但结果对我来说似乎是合理的。

请注意，我对上面第 5 项应用的修复意味着只有第一个 N-M+1线程将报告结果。如果您想要某种不同的行为(不确定您想要哪种模式匹配)，当然还有其他方法可以修改它。

将来，如果您想避免否决票和接近票，我的建议是注意您的问题本质上是在请求调试帮助(“为什么这段代码不起作用？”)，所以它属于 SO expects 的问题类别一个MCVE 。您提供了完整的代码，这很好。可能还有其他一些可以改进的地方:

1. 更清楚地说明问题所在。关于有多少线程正在运行的问题不是很清楚。对于一个好的 MCVE，您应该解释预期结果并显示实际结果。在某些情况下，提及 CUDA 版本、编译命令行以及您正在使用的平台(主机操作系统)也可能很有用。

2. 演示并使用正确的 cuda 错误检查和 cuda-memcheck 的使用。即使您不理解错误输出，也可以在您的问题中描述或包含它 - 这对于其他试图帮助您的人来说会很有用。