Dijkstra 算法的 CUDA 并行程序流程

Question

此处的代码是 cuda 代码，旨在使用 Dijkstra 算法找到最短路径对。

我的代码逻辑在 c 程序中完美运行，而不是在 Cuda 中。我正在使用带有 N 个线程的 1 个块，N 是用户输入的。

第一个疑问，除了共享变量temp之外，每个线程都有自己的变量副本。正确吗？

当我打印结果时，我将所有值存储在数组 d 中并打印其值为零的值。仅当控制流在 s = threadIdx.x.

后不进入循环时才有可能

请帮忙，自上次 24 小时以来一直在调试此问题。

给定的输入是：

Number of vertices: 4

enter the source,destination and cost of the edge\n Enter -1 to end Input\n Edges start from Zero : 0 1 1

enter the source,destination and cost of the edge\n Enter -1 to end Input\n Edges start from Zero : 0 2 5

enter the source,destination and cost of the edge\n Enter -1 to end Input\n Edges start from Zero : 0 3 2

enter the source,destination and cost of the edge\n Enter -1 to end Input\n Edges start from Zero : 1 3 4

enter the source,destination and cost of the edge\n Enter -1 to end Input\n Edges start from Zero : 2 3 7

enter the source,destination and cost of the edge\n Enter -1 to end Input\n Edges start from Zero : -1 -1 -1

#include<stdio.h>
#include<stdlib.h>
#include<time.h>
#include<sys/time.h>
#define nano 1000000L

__global__ void dijkstras(int *a, int *b, int *n)
{
    int i;
    int d[10],p[10],v[10];
    // d stores distnce/cost of each path
    // p stores path taken
    // v stores the nodes already travelled to
    int k,u,s;
    int check =0;

    // shared memory on cuda device
    __shared__ int temp[20];
    for(i=0; i < (*n)*(*n); i++)
    {
        temp[i] = a[i];
    }
    check = check + 1;
    __syncthreads();

    // were passing int s -- node from which distances are calculated
    s = threadIdx.x;
    for(i=0; i<(*n); i++)
    {
        d[i]=temp[s*(*n)+i];
        if(d[i]!=9999)
            p[i]=1;
        else
            p[i]=0;
        v[i]=0;
    }
    p[s]=0;
    v[s]=1;
    for(i=0; i<((*n)-1); i++)
    {
    // findmin starts here
        int i1,j1,min=0;
        for(i1=0;i1<(*n);i1++)
        {
            if(v[i1]==0)
            {
                min=i1;
                break;
            }
        }
        for(j1=min+1;j1<(*n);j1++)
        {
            if((v[j1]==0) && (d[j1]<d[min]))
                min=j1;
        }
        k = min;
    // findmin ends here
        v[k]=1;
        for(u=0; u<(*n); u++)
        {
            if((v[u]==0) && (temp[k*(*n)+u]!=9999))
            { 
                if(d[u]>d[k]+temp[k*(*n)+u])
                {
                    d[u]=d[k]+temp[k*(*n)+u];
                    p[u]=k;
                }
            }
        }
    //storing output
        int count = 0;
        for(i = (s*(*n)); i< (s+1) * (*n); i++)
        {
            b[i] = d[count];
            count++;
        }
    }
    *n = check; 
}



main()
{
    int *a, *b, *n;
    int *d_a, *d_b, *d_n;
    int i,j,c;
    int check = 0;
    printf("enter the number of vertices.... : ");
    n = (int*)malloc(sizeof(int));
    scanf("%d",n);
    int size = (*n) * (*n) * sizeof(int);

    //allocating device memory
    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_n, sizeof(int));

    a = (int*)malloc(size);
    b = (int*)malloc(size);

    check = check +1;
    for(i=0; i<(*n); i++)
        for(j=0; j<=i; j++)
            if(i==j)
                a[(i*(*n) + j)]=0;
            else
                a[(i*(*n) + j)]=a[(j*(*n) + i)]=9999;

    printf("\nInitial matrix is\n");
    for(i=0;i<(*n);i++)
    {
        for(j=0;j<(*n);j++)
        {
            printf("%d ",a[i*(*n)+j]);
        }
        printf("\n");
    }

    while(1)
    {
        printf("\n enter the source,destination and cost of the edge\n Enter -1 to end Input\n Edges start from Zero : \n");
        scanf("%d %d %d",&i,&j,&c);
        if(i==-1)
            break;
        a[(i*(*n) + j)]=a[(j*(*n) + i)]=c;
    }

    printf("\nInput matrix is\n");
    for(i=0;i<(*n);i++)
    {
        for(j=0;j<(*n);j++)
        {
            printf("%d ",a[i*(*n)+j]);
        }
        printf("\n");
    }

    check = check +1;
    // copying input matrix to device
    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_n, n, sizeof(int), cudaMemcpyHostToDevice);
    check++;
    struct timeval start,stop;
    double time;
    int N = *n;
    gettimeofday(&start,NULL);
    dijkstras<<<1,N>>>(d_a, d_b, d_n);
    gettimeofday(&stop,NULL);
    time=(double)(stop.tv_sec-start.tv_sec)+(double)(stop.tv_usec-start.tv_usec)/(double)nano;
    printf("\n TIME TAKEN: %lf\n",time);
    check++;

    // copying result from device to host
    cudaMemcpy(b, d_b, size, cudaMemcpyDeviceToHost);
    cudaMemcpy(n, d_n, sizeof(int), cudaMemcpyDeviceToHost);

    check++;
    //  printing result
    printf("the shortest paths are....");
    for(i=0; i<(N); i++)
    {
        for(j=0; j<(N); j++)
        {
            if(i != j)
                printf("\n the cost of the path from %d to %d = %d\n",i,j,b[i*(N) + j]);
        }
        printf("\n\n");
    }

    printf("your debug value of check in main is %d\n",check);  //5

    printf("your debug value of check in device is %d\n",*n);       // 1+ 7+ 10

    free(a); free(b);free(n);
    cudaFree(d_a); cudaFree(d_b);cudaFree(d_n);
}

Answer 1

这个问题的根本原因是提供了一个未初始化的设备变量作为内核参数。在这个内核调用中：

dijkstras<<<1,N>>>(d_a, d_b, d_n);

d_n 已分配内存，但从未分配值，导致内核中出现未定义的行为。

我认为由于内核本身的设计决策不佳，原始发布者很难检测到这一点。在这个原型中：

__global__ void dijkstras(int *a, int *b, int *n)

n 被用作具有两种完全不同含义的输入和输出，这使得检测调用问题变得更加困难。如果原型是：

__global__ void dijkstras(int *a, int *b, int n, *int check)

那么n和check的作用就更明确了，调用内核时出错，调试时漏掉的可能性也会降低。

Dijkstra 算法的 CUDA 并行程序流程

CUDA parallel program flow for Dijkstra Algorithm

c++

cuda