如何不并行化 OpenACC 中的内部循环

Question

我是使用 OpenACC 进行 GPU 编程的初学者。我试图做一个直接卷积。卷积由 6 个嵌套循环组成。我只希望第一个循环被并行化。我为第一个循环提供了 pragma #pragma acc 循环，为其余循环提供了 #pragma acc loop seq。但是我得到的输出不正确。我采用的并行化循环的方法是否正确？卷积规格：输入通道 - 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);
    }
  }

Answer 1

投稿人在我回答过的 PGI 用户论坛上发布了相同的问题。 （见：https://www.pgroup.com/userforum/viewtopic.php?f=4&t=7614）。主题问题是不正确的，因为内部循环没有得到并行化，也不是问题的原因。

这里的问题是代码在共享的“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;
      }
    }