Cuda передавая массив структур

Question

Я новичок в cuda и пытаюсь распараллелить очень простую программу, показанную ниже, которая была создана по этой ссылке: https://devblogs.nvidia.com/even-easier-introduction-cuda/

typedef struct{
    int temp;
    int newtemp;
    int neighbors[20];
} S;
void add(int n, S * s){
    for(int i = 0; i < n; i++){
        int newTemp = 0;
        for(int j = 0; j < 20; j++){
            newTemp += s[s[i].neighbors[j]].temp;
        }
        newTemp /= 3;
        s[i].newtemp = newTemp;
    }
}
int main(int argc, char *argv[]){
    int n = 1<<21;
    S grid[n];

    for(int i = 0; i < n; i++){
        S tmp1;
        tmp1.temp = rand();
        for(int j = 0; j<20; j++){
            tmp1.neighbors[j] = rand()%n;
        }
        grid[i] = tmp1;
    }
    struct timespec start, end;
    double gettime_diff, time_diff; 
    clock_t t, starttime, endtime; 

    clock_gettime(CLOCK_REALTIME, &start);
    t = clock(); 
    time(&starttime);

    add(n,grid);

    for(int i = 0; i < n; i++){
        grid[i].temp = grid[i].newtemp;
        if(i%83940==1)printf("%d\n",grid[i].temp);
    }
    return 0;
}

Однако я не получаю желаемых результатов, как при обновленииtemp все новые значения равны 0. Я думаю, что проблема в том, что массив структур, которые я передаю моей функции добавления, не может быть доступен в памяти устройства. Мне, однако, трудно понять, как это исправить. Я нашел этот пост на stackoverflow и немного не уверен, что предложенный ответ сделал, чтобы решить проблему: Массив структур массивов CUDA C

Код CUDA, который я имею для справки, находится здесь:

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#define SIZE 1000
#define NS_PER_US 1000

typedef struct{
    int temp;
    int newtemp;
    int neighbors[20];
} S;
__global__ void add(int n, S * s){
    int index = threadIdx.x;
    int stride = blockDim.x;
    //printf("%d\n",(n-index)/stride);
    //printf("%d\n",s[0].temp);
    for(int i = index; i < n; i+=stride){
        printf("%d\n",index);
        int newTemp = 0;
        for(int j = 0; j < 20; j++){
            newTemp += s[s[i].neighbors[j]].temp;
        }
        printf("%d\n",index);
        newTemp /= 3;
        s[i].newtemp = newTemp;
    }
}

int main(int argc, char *argv[]){
    int  *h_a;
    int  *d_a;
    int  num_blocks= 2;
    int  num_th_per_blk= 5;

    int n = 1<<21;
    S grid[n];

    for(int i = 0; i < n; i++){
        S tmp1;
        tmp1.temp = rand();
        for(int j = 0; j<20; j++){
            tmp1.neighbors[j] = rand()%n;
        }
        grid[i] = tmp1;
    }
    struct timespec start, end;
    double gettime_diff, time_diff; 
    clock_t t, starttime, endtime; 

    clock_gettime(CLOCK_REALTIME, &start);
    t = clock(); 
    time(&starttime);

    size_t  memSize;
    memSize = num_blocks* num_th_per_blk* sizeof(int);
    h_a= (int*) malloc(memSize);

    cudaMallocManaged((void **)&grid, n * sizeof(S));
    cudaMalloc( (void**) &d_a, memSize);
    dim3  dimGrid(num_blocks);
    dim3  dimBlock(num_th_per_blk);    

    add<<< dimGrid, dimBlock >>>(n,grid);

    cudaMemcpy( h_a, d_a, memSize,cudaMemcpyDeviceToHost);

    for(int i = 0; i < n; i++){
        grid[i].temp = grid[i].newtemp;
        if(i%83940==1)printf("%d\n",grid[i].newtemp);
    }
    clock_gettime(CLOCK_REALTIME, &end); 
    t = clock() - t; 
    time(&endtime);
    gettime_diff = (double) ((end.tv_sec - start.tv_sec)*CLOCKS_PER_SEC) + (double)((end.tv_nsec - start.tv_nsec)/NS_PER_US);
    time_diff = difftime(endtime, starttime);

    printf("\ttime (clock_gettime) %f\n", gettime_diff);
    printf("\ttime (clock) %f\n", ((float)t)/CLOCKS_PER_SEC);
    printf("\ttime (time) %f\n", time_diff); 

    return 0;
}

Я чувствую, что здесь есть простое исправление, которого я не вижу, или, может быть, мне не хватает ключевой концепции. В любом случае любая помощь будет принята с благодарностью.

Answer 1

В вашем коде действительно много неправильного, настолько, что легче опубликовать рабочую версию, чем указывать на все отдельные ошибки:

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#define NS_PER_US 1000

typedef struct{
    int temp;
    int newtemp;
    int neighbors[20];
} S;

__global__ 
void add(int n, S * s)
{
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    for(int i = index; i < n; i+=stride){
        int newTemp = 0;
        for(int j = 0; j < 20; j++){
            newTemp += s[s[i].neighbors[j]].temp;
        }
        newTemp /= 3;
        s[i].newtemp = newTemp;
    }
}

int main(int argc, char *argv[]){
    int n = 1<<10;
    S* grid;
    cudaMallocManaged((void **)&grid, n * sizeof(S));

    for(int i = 0; i < n; i++){
        S tmp1;
        tmp1.temp = rand()%n;
        for(int j = 0; j<20; j++){
            tmp1.neighbors[j] = rand()%n;
        }
        grid[i] = tmp1;
    }
    struct timespec start, end;
    double gettime_diff, time_diff; 
    clock_t t, starttime, endtime; 

    clock_gettime(CLOCK_REALTIME, &start);
    t = clock(); 
    time(&starttime);

    int  num_th_per_blk= 32;
    int  num_blocks= (n / num_th_per_blk) + (n % num_th_per_blk > 0) ? 1 : 0;

    dim3  dimGrid(num_blocks);
    dim3  dimBlock(num_th_per_blk);    

    add<<< dimGrid, dimBlock >>>(n,grid);
    cudaDeviceSynchronize();

    for(int i = 0; i < n; i++){
        grid[i].temp = grid[i].newtemp;
        if(i%10==1)printf("%d %d\n",i,grid[i].temp);
    }
    clock_gettime(CLOCK_REALTIME, &end); 
    t = clock() - t; 
    time(&endtime);
    gettime_diff = (double) ((end.tv_sec - start.tv_sec)*CLOCKS_PER_SEC) + (double)((end.tv_nsec - start.tv_nsec)/NS_PER_US);
    time_diff = difftime(endtime, starttime);

    printf("\ttime (clock_gettime) %f\n", gettime_diff);
    printf("\ttime (clock) %f\n", ((float)t)/CLOCKS_PER_SEC);
    printf("\ttime (time) %f\n", time_diff); 

    return 0;
}

Самая вопиющая ошибка заключается в том, как вы обрабатываетеgrid в коде хоста. Выполнение этого:

S grid[n];

// code initializing grid

cudaMallocManaged((void **)&grid, n * sizeof(S));

является недопустимым (вы не должны пытаться установить для сетки другое значение указателя, это не указатель) и бессмысленным. cudaMallocManaged выделяет новую память, поэтому все, что вы делаете, это инициализируете grid, затем отбрасываете всю тщательно инициализированную память и заменяете ее неинициализированной памятью, которую вы передаете ядру. Затем ядро ​​оперирует случайными данными. Также обратите внимание, что цикл шага сетки в ядре также неверен, и исходный код и версия CUDA потенциально страдают от целочисленного переполнения из-за того, как вы инициализируете временные члены структуры в обеих версиях, используя rand().