업데이트 지침 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"가 사용한다고 가정하는 것을 어렵게 만듭니다.

외부 "i"루프는 병렬 처리가 안전하지 않으므로 "kernels"지시문을이 루프 외부에 두어서는 안됩니다. 어쩌면 당신이 "입력"배열을 민영화 한 다음 vbias, hbias, W 배열을 업데이트 할 때 atomics를 사용하면 작동하지만 성능이 떨어질 수 있습니다. (당신은 또한 다른 배열이 사유화되어야하는지 아니면 전역 적이기 때문에 원자 연산이 필요한지를 결정할 필요가있다).

내가 제안하는 것은 내부 루프 주위에 "#pragma acc parallel loop"를 한 번에 하나씩 넣는 것부터 시작하는 것이 좋습니다. 다음 작업을 수행하기 전에 각각 작동하는지 확인하십시오. 또한 나는 "단계"루프가 병렬화 될 수 있는지 의심 스럽기 때문에 대신 "gibbs_hvh"서브 루틴 내에서 루프를 병렬 처리해야합니다.

CUDA 통합 메모리 (-ta = tesla : managed)를 사용하고 있으므로 데이터 영역을 추가하는 것이 아마도 필요하지 않습니다. 그러나 나중에 관리 메모리를 사용하지 않을 계획이라면 다음 단계는 외부 "i"루프 (또는 프로그램의 상위 지점)에 데이터 지시문을 추가 한 다음 update 지시문을 사용하여 외부 " i "루프).