Как не распараллелить внутренние циклы в OpenACC - PullRequest
Новая
       100

Как не распараллелить внутренние циклы в OpenACC

0 голосов
/ 30 апреля 2020

Я новичок в программировании на GPU с OpenA CC. Я пытался сделать прямую свертку. Свертка состоит из 6 вложенных циклов. Я только хочу, чтобы первый l oop был распараллелен. Я дал прагму #pragma cc l oop для первого l oop и #pragma cc l oop seq для остальных. Но вывод, который я получаю, не верен. Является ли мой подход к распараллеливанию l oop правильным? Спецификации для свертки: Входные каналы-3, Входной размер - 224X224X3, Выходные каналы - 64, Выходной размер - 111X111X64, размер фильтра - 3X3X3X64. Ниже приводится ссылка на файлы заголовков dog.h и squeezenet_params.h. https://drive.google.com/drive/folders/1a9XRjBTrEFIorrLTPFHS4atBOPrG886i

# include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include "squeezenet_params.h"
#include "dog.h"

void conv3x3(
  const int input_channels, const int input_size,
  const int pad, const int stride, const int start_channel,
  const int output_size, const float* restrict input_im, const float* restrict filter_weight,
  const float* restrict filter_bias, float* restrict output_im){
    #pragma acc data copyin (input_im[0:150527],filter_weight[0:1727],filter_bias[0:63]) copyout(output_im[0:788543])
    {
    #pragma acc parallel
    {
      #pragma acc loop
      for(int p=0;p<64;++p){
        filter_weight += p * input_channels * 9;
        float bias = filter_bias[p];
        output_im += (start_channel + p) * output_size * output_size;

        //loop over output feature map
        #pragma acc loop seq
        for(int i = 0; i < output_size; i++)
        {
          #pragma acc loop seq
          for(int j = 0; j < output_size; j++)
          {
            //compute one element in the output feature map
            float tmp = bias;

            //compute dot product of 2 input_channels x 3 x 3 matrix
            #pragma acc loop seq
            for(int k = 0; k < input_channels; k++)
            {
              #pragma acc loop seq
              for(int l = 0; l < 3; l++)
              {
                int h = i * stride + l - pad;
                #pragma acc loop seq
                for(int m = 0; m < 3; m++)
                {
                  int w = j * stride + m - pad;
                  if((h >= 0) && (h < input_size) && (w >= 0) && (w < input_size))
                  {
                    tmp += input_im[k * input_size * input_size + (i * stride + l - pad) * input_size + j * stride + m - pad] \
                                     * filter_weight[9 * k + 3 * l + m];
                  }
                }
              }
            }

            //add relu activation after conv
            output_im[i * output_size + j] = (tmp > 0.0) ? tmp : 0.0;
          }
        }
      }
    }
    }
  }

  void main(){
    float * result = (float*)malloc(sizeof(float) * (1 * 64 * 111 * 111));
    conv3x3(3,224,0,2,0,111,sample,conv1_weight,conv1_bias,result);
    for(int i=0;i<64 * 111 * 111;++i){
      //if(result[i]>0)
        printf("%f:%d\n",result[i],i);
    }
  }

1 Ответ

0 голосов
/ 30 апреля 2020

Автор опубликовал тот же вопрос на форумах пользователей PGI, где я ответил. (См .: https://www.pgroup.com/userforum/viewtopic.php?f=4&t=7614). Вопрос topi c является неправильным в том смысле, что внутренние циклы не распараллеливаются и не являются причиной проблемы.

Проблема здесь в том, что в коде есть условие гонки на общем указателе «output_im». Мое предлагаемое решение состоит в том, чтобы вычислять смещение на поток в массив, а не пытаться манипулировать самим указателем. 

  for(int p=0;p<64;++p){
    filter_weight += p * input_channels * 9;
    float bias = filter_bias[p];
    int offset;
    offset = (start_channel + p) * output_size * output_size;

    //loop over output feature map
    #pragma acc loop vector collapse(2)
    for(int i = 0; i < output_size; i++)
    {
      for(int j = 0; j < output_size; j++)
      {
         ... cut ... 
          }
        }

        //add relu activation after conv
        int idx = offset + (i * output_size + j);
        output_im[idx] = (tmp > 0.0) ? tmp : 0.0;
      }
    }
...