Aktualisierung Richtlinien OpenACC

Question

Wenn ich einen Kernel auf der oberen Schleife haben, warum ich nicht diese 2-Richtlinien verwenden können:

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

ich diese Variablen aktualisieren hbias, vbias, W in folgenden Code, aber es wird nicht funktionieren:

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; 
    } 

Aber wenn ich auf jedem verschachtelten Schleifen viele getrennte Kerne haben arbeiten, kann ich die Variablen aktualisieren:

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" -Direktiven können nur im Host-Code verwendet werden, da die Datenbewegung vom Host initiiert werden muss. Sie können sie nicht innerhalb einer Rechenregion haben.

Es gibt viele Probleme mit diesem Code. Erstens ist es wahrscheinlich eine schlechte Übung, die gleiche Indexvariable "i" in diesem Fall für verschachtelte Schleifen zu verwenden. Obwohl Scoping-Regeln dies erlauben, ist es schwierig zu sagen, welches "i" der Code verwenden soll.

Die äußere "i" -Schleife ist wahrscheinlich nicht sicher zu parallelisieren, so dass Sie nicht die "Kernel" -Richtlinie außerhalb dieser Schleife setzen sollten. Vielleicht, wenn Sie das "Eingabe" -Array privatisiert haben und dann atomics beim Aktualisieren der Vbias, Hbias, W-Arrays verwendet haben, könnte es funktionieren, aber Ihre Leistung wäre schlecht. (Sie müssen auch bestimmen, ob die anderen Arrays privatisiert werden müssen oder global sind, also atomare Operationen benötigen).

Was ich vorschlagen würde ist mit "#pragma acc parallel loop" um die inneren Schleifen nacheinander zu setzen. Stellen Sie sicher, dass jeder funktioniert, bevor Sie mit dem nächsten fortfahren. Außerdem bezweifle ich stark, dass die "Schritt" -Schleife parallelisierbar ist, so dass Sie wahrscheinlich die Schleifen innerhalb der "gibbs_hvh" -Unterroutine stattdessen parallelisieren müssen.

Da Sie CUDA Unified Memory verwenden (-ta = tesla: verwaltet), ist das Hinzufügen von Datenbereichen wahrscheinlich nicht notwendig. Wenn Sie jedoch beabsichtigen, in Zukunft keinen verwalteten Speicher zu verwenden, besteht der nächste Schritt darin, Datenanweisungen um die äußere "i" -Schleife (oder an einem höheren Punkt im Programm) hinzuzufügen und dann die update-Direktive zu verwenden, um Daten nach dem äußeren zu synchronisieren. ich "Schleife".