Использование openmp для распределения умножения матриц по нескольким графическим процессорам через openacc с использованием C - PullRequest
Новая
       69

Использование openmp для распределения умножения матриц по нескольким графическим процессорам через openacc с использованием C

0 голосов
/ 20 апреля 2019

Я пытаюсь распределить работу по умножению двух матриц NxN на 3 графических процессора nVidia с использованием 3 потоков OpenMP. (Значения матрицы станут большими, поэтому тип данных long long.) Однако у меня возникли проблемы с размещением #pragma acc parallel loop в правильном месте. Я использовал некоторые примеры в PDF-файлах nVidia, но не к счастью. Я знаю, что самый внутренний цикл не может быть распараллелен. Но я бы хотел, чтобы каждый из трех потоков владел графическим процессором и выполнял часть работы. Обратите внимание, что матрицы ввода и вывода определены как глобальные переменные, так как я постоянно исчерпывал память стека.

Я попробовал приведенный ниже код, но все ошибки компиляции указывают на строку 75, которая является #pragma acc parallel loop строкой 

[test@server ~]pgcc -acc -mp -ta=tesla:cc60 -Minfo=all -o testGPU matrixMultiplyopenmp.c

PGC-S-0035-Syntax error: Recovery attempted by replacing keyword for by keyword barrier (matrixMultiplyopenmp.c: 75)

PGC-S-0035-Syntax error: Recovery attempted by replacing acc by keyword enum (matrixMultiplyopenmp.c: 76)

PGC-S-0036-Syntax error: Recovery attempted by inserting ';' before keyword for (matrixMultiplyopenmp.c: 77)

PGC/x86-64 Linux 18.10-1: compilation completed with severe errors

Функция: 

void multiplyMatrix(long long int matrixA[SIZE][SIZE], long long int matrixB[SIZE][SIZE], long long int matrixProduct[SIZE][SIZE])
{
    // Get Nvidia device type
    acc_init(acc_device_nvidia);

    // Get Number of GPUs in system
    int num_gpus = acc_get_num_devices(acc_device_nvidia);

    //Set the number of OpenMP thread to the number of GPUs
    #pragma omp parallel num_threads(num_gpus)
    {
        //Get thread openMP number and set the GPU device to that number
        int threadNum = omp_get_thread_num();
        acc_set_device_num(threadNum, acc_device_nvidia);

        int row;
        int col;
        int key;

        #pragma omp for
        #pragma acc parallel loop
        for (row = 0; row < SIZE; row++)
            for (col = 0; col < SIZE; col++)
                for (key = 0; key < SIZE; key++)
                    matrixProduct[row][col] = matrixProduct[row][col] + (matrixA[row][key] * matrixB[key][col]);
    }
}

1 Ответ

0 голосов
/ 22 апреля 2019

Как указывает fisehara, вы не можете одновременно использовать цикл OpenMP "for" и параллельный цикл OpenACC для цикла for. Вместо этого вам нужно вручную разложить работу по потокам OpenMP. Пример ниже.

Есть ли причина, по которой вы хотите использовать здесь несколько графических процессоров? Скорее всего, умножение матриц будет соответствовать одному графическому процессору, поэтому нет необходимости в дополнительных затратах на внедрение параллелизации на стороне хоста.

Кроме того, я обычно рекомендую использовать MPI + OpenACC для программирования нескольких графических процессоров. Разложение домена, естественно, является частью MPI, но не присуще OpenMP. Кроме того, MPI дает вам взаимно-однозначное соотношение между хост-процессом и ускорителем, позволяет масштабировать за пределы одного узла, и вы можете воспользоваться преимуществами CUDA Aware MPI для прямой передачи данных с GPU на GPU. Для получения дополнительной информации выполните поиск в Интернете по запросу «MPI OpenACC», и вы найдете несколько учебных пособий. Класс # 2 на https://developer.nvidia.com/openacc-advanced-course - хороший ресурс. 

% cat test.c
#include <stdlib.h>
#include <stdio.h>
#include <omp.h>
#ifdef _OPENACC
#include <openacc.h>
#endif

#define SIZE 130

void multiplyMatrix(long long int matrixA[SIZE][SIZE], long long int matrixB[SIZE][SIZE], long long int matrixProduct[SIZE][SIZE])
{

#ifdef _OPENACC
    // Get Nvidia device type
    acc_init(acc_device_nvidia);
    // Get Number of GPUs in system
    int num_gpus = acc_get_num_devices(acc_device_nvidia);
#else
    int num_gpus = omp_get_max_threads();
#endif
    if (SIZE<num_gpus) {
        num_gpus=SIZE;
    }
    printf("Num Threads: %d\n",num_gpus);

    //Set the number of OpenMP thread to the number of GPUs
    #pragma omp parallel num_threads(num_gpus)
    {
        //Get thread openMP number and set the GPU device to that number
        int threadNum = omp_get_thread_num();
#ifdef _OPENACC
        acc_set_device_num(threadNum, acc_device_nvidia);
        printf("THID %d using GPU: %d\n",threadNum,threadNum);
#endif
        int row;
        int col;
        int key;
        int start, end;
        int block_size;
        block_size = SIZE/num_gpus;
        start = threadNum*block_size;
        end = start+block_size;
        if (threadNum==(num_gpus-1)) {
           // add the residual to the last thread
           end = SIZE;
        }
        printf("THID: %d, Start: %d End: %d\n",threadNum,start,end-1);

        #pragma acc parallel loop \
          copy(matrixProduct[start:end-start][:SIZE]), \
          copyin(matrixA[start:end-start][:SIZE],matrixB[:SIZE][:SIZE])
        for (row = start; row < end; row++) {
            #pragma acc loop vector
            for (col = 0; col < SIZE; col++) {
                for (key = 0; key < SIZE; key++) {
                    matrixProduct[row][col] = matrixProduct[row][col] + (matrixA[row][key] * matrixB[key][col]);
        }}}
    }
}

int main() {
   long long int matrixA[SIZE][SIZE];
   long long int matrixB[SIZE][SIZE];
   long long int matrixProduct[SIZE][SIZE];
   int i,j;
   for(i=0;i<SIZE;++i) {
     for(j=0;j<SIZE;++j) {
        matrixA[i][j] = (i*SIZE)+j;
        matrixB[i][j] = (j*SIZE)+i;
        matrixProduct[i][j]=0;
     }
   }
   multiplyMatrix(matrixA,matrixB,matrixProduct);
   printf("Result:\n");
   for(i=0;i<SIZE;++i) {
      printf("%d: %ld %ld\n",i,matrixProduct[i][0],matrixProduct[i][SIZE-1]);
   }

}
% pgcc test.c -mp -ta=tesla -Minfo=accel,mp
multiplyMatrix:
     28, Parallel region activated
     49, Generating copyin(matrixB[:130][:])
         Generating copy(matrixProduct[start:end-start][:131])
         Generating copyin(matrixA[start:end-start][:131])
         Generating Tesla code
         52, #pragma acc loop gang /* blockIdx.x */
         54, #pragma acc loop vector(128) /* threadIdx.x */
         55, #pragma acc loop seq
     54, Loop is parallelizable
     55, Complex loop carried dependence of matrixA->,matrixProduct->,matrixB-> prevents parallelization
         Loop carried dependence of matrixProduct-> prevents parallelization
         Loop carried backward dependence of matrixProduct-> prevents vectorization
     59, Parallel region terminated
% a.out
Num Threads: 4
THID 0 using GPU: 0
THID: 0, Start: 0 End: 31
THID 1 using GPU: 1
THID: 1, Start: 32 End: 63
THID 3 using GPU: 3
THID: 3, Start: 96 End: 129
THID 2 using GPU: 2
THID: 2, Start: 64 End: 95
Result:
0: 723905 141340355
1: 1813955 425843405
2: 2904005 710346455
3: 3994055 994849505
...
126: 138070205 35988724655
127: 139160255 36273227705
128: 140250305 36557730755
129: 141340355 36842233805
...