Блог

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

Я попробовал приведенный ниже код, но я получаю ошибки компиляции, указывающие на строку 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]);
    }
}
  

Как указывает Фишехара, вы не можете комбинировать цикл OpenMP «for» с параллельным циклом OpenACC в одном и том же цикле for. Вместо этого вам нужно вручную разложить работу по потокам OpenMP. Пример ниже.

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

Кроме того, я обычно рекомендую использовать MPI OpenACC для программирования с несколькими графическими процессорами. Декомпозиция домена, естественно, является частью MPI, но не присуща OpenMP. Кроме того, MPI обеспечивает однозначную связь между хост-процессом и ускорителем, позволяет масштабировать за пределы одного узла, и вы можете воспользоваться поддержкой CUDA 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: %dn",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: %dn",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: %dn",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 %ldn",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
  

Я столкнулся с проблемой компиляции MPI OpenACC в общей системе, к которой я был ограничен, и не смог обновить компилятор. Решение, которое я в итоге использовал, состояло в том, чтобы сначала выполнить работу с OMP, а затем вызвать функцию OpenACC следующим образом:

 //Main code
pragma omp parallel num_threads(num_gpus)
    {
        #pragma omp for private(tid)
        for (tid = 0; tid < num_gpus; tid  )
        {
            //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);

            // check with thread is using which GPU
            int gpu_num = acc_get_device_num(acc_device_nvidia);
            printf("Thread # %d is going to use GPU # %d n", threadNum, gpu_num);

            //distribute the uneven rows
            if (threadNum < extraRows)
            {
                startRow = threadNum * (rowsPerThread   1);
                stopRow = startRow   rowsPerThread;
            }
            else
            {
                startRow = threadNum * rowsPerThread   extraRows;
                stopRow = startRow   (rowsPerThread - 1);
            }
            // Debug to check allocation of data to threads
            //printf("Start row is %d, and Stop rows is %d n", startRow, stopRow);

            GPUmultiplyMatrix(matrixA, matrixB, matrixProduct, startRow, stopRow);
        }
    }
void GPUmultiplyMatrix(long long int matrixA[SIZE][SIZE], long long int 
matrixB[SIZE][SIZE], long long int matrixProduct[SIZE][SIZE], int 
startRow, int stopRow)
    {
        int row;
        int col;
        int key;

        #pragma acc parallel loop collapse (2)
        for (row = startRow; row <= stopRow; 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]);
    }