OpenACC 및 GNU 과학 라이브러리 - gsl_matrix의 데이터 이동

Question

기록 된 openacc 개요 코스 비디오를 강의 3에서 살펴본 결과 데이터 이동 표현에 대해 설명했습니다. copy_in()을 사용하면 어떻게 gsl_matrix*을 CPU에서 GPPu로 옮길 것입니까? 예를 들어 CPU의 경우,

gsl_matrix *Z = gsl_matrix_calloc(100, 100), 

과 같이 할 수 있습니다. 그러면 100x100 매트릭스의 0이됩니다. 이제 Z는

typedef struct{ 
    size_t size1; 
    size_t size2; 
    size_t tda; 
    double * data; 
    gsl_block * block; 
    int owner; 
    } gsl_matrix; 

가 어떻게 copyin()를 사용하여 GPU에 대한 CPU에서입니다 (포인터) Z의 데이터 이동을 표현하는 것처럼 보이는 gsl_matrix 구조에 대한 포인터입니다?

Answer 1

OpenACC 데이터와 계산 영역 내에서 GSL을 사용하는 것에 관해 직접 말할 수는 없지만 동적 데이터 멤버를 사용하는 집계 유형에 대한 일반적인 대답을 제공 할 수 있습니다.

PGI 컴파일러와 새로운 NVIDIA 장치를 사용한다고 가정 할 때 가장 먼저 할 일은 CUDA 통합 메모리 (UVM)입니다. 플래그 "-ta = tesla : managed"로 컴파일하면 동적으로 할당 된 모든 데이터가 CUDA 런타임에 의해 관리되므로 데이터 이동을 직접 관리 할 필요가 없습니다. 오버 헤드가 포함되고 경고가 있지만 일이 더 쉬워집니다. PGI 16.9 이상과 함께 제공되는 CUDA 8.0은 UVM 성능을 향상시킵니다.

UVM이 없으면 데이터의 수동 전체 복사본을 수행해야합니다. 다음은 처음으로 장치에서 상위 구조를 만들고 얕은 복사본을 수행하는 기본 개념입니다. 다음으로 디바이스의 "data"라는 동적 배열을 생성하고 배열에 초기 값을 복사 한 다음 디바이스 포인터를 디바이스 구조체의 데이터 포인터에 연결합니다. "블록"자체는 동적 데이터 멤버가있는 구조체의 배열이므로 배열을 반복하여 장치에 데이터 배열을 만들어야합니다.

matrix * mat = (matrix*) malloc(sizeof(matrix)); 
    #pragma acc enter data copyin(mat) 
// Change this to the correct size of "data" and blocks 
    #pragma acc enter data copyin(mat.data[0:dataSize]); 
    #pragma acc enter data copyin(mat.block[0:blockSize]); 
    for (i=0; i < blockSize; ++i) { 
     #pragma acc enter data copyin(mat.block[i].data[0:mat.block[i].size]) 
    } 

은 삭제 상향식 (bottom-up)

for (i=0; i < blockSize; ++i) { 
    #pragma acc exit data delete(mat.block[i].data) 
} 
#pragma acc exit data delete(mat.block); 
#pragma acc exit data delete(mat.data); 
#pragma acc exit data delete(mat); 

업데이트

에서 삭제, 다시 구조를 걸어, 스칼라 또는 기본 데이터 타입의 배열을 업데이트해야합니다. 즉 "데이터"는 업데이트하지만 "차단"하지 않습니다. 업데이트는 얕은 복사본을 수행하므로 "블록"을 업데이트하면 잘못된 주소로 이어지는 호스트 또는 장치 포인터가 업데이트됩니다.

마지막으로 계산 영역에서 사용할 때 행렬 변수를 "present"절에 넣어야합니다.

#pragma acc parallel loop present(mat)