Поток параллельной программы CUDA для алгоритма Дейкстры

код здесь представляет собой код cuda и предназначен для поиска кратчайшего пути пары с использованием алгоритма Дейкстры.

Моя логика кода прекрасно работает в программе c, а не в Cuda. Я использую 1 блок с N потоков, N вводится пользователем.

Первое сомнение: каждый поток имеет свою собственную копию переменных, кроме общей переменной temp. Правильный ?

Когда я печатаю результаты, я сохраняю все значения в массиве d и печатаю его значение, равное нулю для всех. Это возможно, только если поток управления не входит в цикл после s = threadIdx.x.

Пожалуйста, помогите, отлаживал это с последних 24 часов.

Данный вход является:

Количество вершин: 4

введите источник, назначение и стоимость края \ n Введите -1, чтобы закончить
Ввод \ n Ребра начинаются с нуля: 0 1 1

введите источник, назначение и стоимость края \ n Введите -1, чтобы закончить
Ввод \ n Ребра начинаются с нуля: 0 2 5

введите источник, назначение и стоимость края \ n Введите -1, чтобы закончить
Ввод \ n Ребра начинаются с нуля: 0 3 2

введите источник, назначение и стоимость края \ n Введите -1, чтобы закончить
Ввод \ n Ребра начинаются с нуля: 1 3 4

введите источник, назначение и стоимость края \ n Введите -1, чтобы закончить
Ввод \ n Ребра начинаются с нуля: 2 3 7

введите источник, назначение и стоимость края \ n Введите -1, чтобы закончить
Ввод \ n края начинаются с нуля: -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);
}

Вывод этого кода

-3

Решение

Основной причиной этой проблемы было предоставление неинициализированной переменной устройства в качестве аргумента ядра. В этом ядре вызов:

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было бы гораздо яснее, и вероятность ошибки при вызове ядра и ее отсутствия при отладке была бы уменьшена.

1

Другие решения


По вопросам рекламы ammmcru@yandex.ru
Adblock
detector