更新ディレクティブOpenACC

Question

私はこれらの2つのディレクティブを使用することはできませんなぜトップループ、上のカーネル持っている：私はこれらの変数を更新する必要があり

#pragma acc update device(hbias[0:n_hidden],W[0:n_hidden][0:n_visible]) 
#pragma acc update device(vbias[0:n_visible) 

をhbias、以下のコードでvbias、W、それ動作しません。

void RBM::contrastive_divergence(int train_X[6][6], double learning_rate, int k) { 
    double r= rand()/(RAND_MAX + 1.0); 

     int * input = new int[n_visible]; 
     double *ph_mean = new double[n_hidden]; 
     int *ph_sample = new int[n_hidden]; 
     double *nv_means = new double[n_visible]; 
     int *nv_samples = new int[n_visible]; 
     double *nh_means = new double[n_hidden]; 
     int *nh_samples = new int[n_hidden]; 

     #pragma acc kernels 
     for (int i = 0; i<train_N; i++) { 


      for (int j = 0; j< n_visible; j++){ 
       input[j] = train_X[i][j]; 
      } 


      sample_h_given_v(input, ph_mean, ph_sample,r); 

      for (int step = 0; step<k; step++) { 
       if (step == 0) { 
        gibbs_hvh(ph_sample, nv_means, nv_samples, nh_means, nh_samples,r); 
       } 
       else { 
        gibbs_hvh(nh_samples, nv_means, nv_samples, nh_means, nh_samples,r); 
       } 
      } 


      for (int i = 0; i<n_hidden; i++) { 
       for (int j = 0; j<n_visible; j++) { 

       W[i][j] += learning_rate * (ph_mean[i] * input[j] - nh_means[i] * nv_samples[j])/N; 

       } 
       hbias[i] += learning_rate * (ph_sample[i] - nh_means[i])/N; 

      } 
    //this directive 
     #pragma acc update device(hbias[0:n_hidden],W[0:n_hidden][0:n_visible]) 


      for (int i = 0; i<n_visible; i++) { 
       vbias[i] += learning_rate * (input[i] - nv_samples[i])/N; 
      } 
    //and this directive 
     #pragma acc update device(vbias[0:n_visible) 
    } 

     delete[] input; 
     delete[] ph_mean; 
     delete[] ph_sample; 
     delete[] nv_means; 
     delete[] nv_samples; 
     delete[] nh_means; 
     delete[] nh_samples; 
    } 

しかし、私は、各ネストされたループに取り組んで多くの分離カーネルを持っているとき、私は、変数を更新することができます。

void RBM::contrastive_divergence(int train_X[6][6], double learning_rate, int k) { 
    double r= rand()/(RAND_MAX + 1.0); 

     int * input = new int[n_visible]; 
     double *ph_mean = new double[n_hidden]; 
     int *ph_sample = new int[n_hidden]; 
     double *nv_means = new double[n_visible]; 
     int *nv_samples = new int[n_visible]; 
     double *nh_means = new double[n_hidden]; 
     int *nh_samples = new int[n_hidden]; 


    for (int i = 0; i<train_N; i++) { 

      #pragma acc kernels 
       for (int j = 0; j< n_visible; j++){ 
        input[j] = train_X[i][j]; 
       } 


       sample_h_given_v(input, ph_mean, ph_sample,r); 
      #pragma acc kernels 
       for (int step = 0; step<k; step++) { 
        if (step == 0) { 
         gibbs_hvh(ph_sample, nv_means, nv_samples, nh_means, nh_samples,r); 
        } 
        else { 
         gibbs_hvh(nh_samples, nv_means, nv_samples, nh_means, nh_samples,r); 
        } 
       } 

      #pragma acc kernels 
      { 
       for (int i = 0; i<unhidden; i++) { 
        for (int j = 0; j<n_visible; j++) { 

         W[i][j] += learning_rate * (ph_mean[i] * input[j] - nh_means[i] * nv_samples[j])/N; 

        } 
       hbias[i] += learning_rate * (ph_sample[i] - nh_means[i])/N; 

       } 
     //this directive 
      #pragma acc update device(hbias[0:n_hidden],W[0:n_hidden][0:n_visible]) 
      } 


      #pragma acc kernels 
      { 
       for (int i = 0; i<n_visible; i++) { 
        vbias[i] += learning_rate * (input[i] - nv_samples[i])/N; 
       } 

      //and this directive 
       #pragma acc update device(vbias[0:n_visible) 
      } 
    } 

     delete[] input; 
     delete[] ph_mean; 
     delete[] ph_sample; 
     delete[] nv_means; 
     delete[] nv_samples; 
     delete[] nh_means; 
     delete[] nh_samples; 
    } 

Answer 1

"Update"ディレクティブは、データの移動をホストから開始する必要があるため、ホストコードでのみ使用できます。計算領域内にそれらを置くことはできません。

このコードには多くの問題があります。まず、ネストされたループに対して同じインデックス変数 "i"を使用することはおそらく貧弱な方法です。スコープルールでは許可されていますが、どのコードを使用しているのかを判断することは困難です。

外側の "i"ループはおそらく並列化が安全でないため、このループの外側に "kernels"ディレクティブを置いてはいけません。たぶんあなたが "入力"配列を民営化し、vbias、hbias、W配列を更新するときにアトミックを使用しても動作するかもしれませんが、パフォーマンスは悪くなります。 （他の配列を民営化する必要があるのか​​、それともグローバルなのでアトミック操作が必要なのかを判断する必要もあります）。

私が示唆したいのは、「#pragma acc parallel loop」を内部ループの周りに1つずつ配置することから始めることです。次のものに移動する前に、それぞれが動作することを確認してください。また、私は "step"ループが並列化可能であることを非常に疑うので、代わりに "gibbs_hvh"サブルーチンの中でループを並列化する必要があります。

CUDAユニファイドメモリ（-ta = tesla：managed）を使用しているので、データ領域の追加はおそらく必要ありません。しかし、将来、マネージメモリを使用しないことを計画している場合は、外側の "i"ループ（またはプログラムの上位ポイント）にデータディレクティブを追加してから、updateディレクティブを使用して、 i "ループ）。