Execution Model
- parallel regions : 일반적으로 work-sharing loops를 포함한다.
- kernels regions : 일반적으로 하나 또는 그 이상의 loops을 포함한다. 이는 kernels로써 실행된다.
- serial regions : 순차 코드 블록이다.
- OpenACC의 3 levels of parallelism
- gang : coarse-grain, 많은 gang들이 accelerator에서 실행된다.
- worker : fine-grain, 각 gang은 하나 혹은 그 이상의 worker를 가진다.
- vector : for SIMD 또는 worker 내에서의 vector operations.
- Device에서 compute region을 실행할 때, 하나 혹은 그 이상의 gangs가 실행된다. 이 때 이 gangs는 gang-redundant mode(GR mode)로 실행되며, 이는 각 gang의 하나의 worker의 하나의 vector lane이 동일한 코드를 실행하는 것을 의미한다.
-----------------------------------------------------------------------------------------------------------
Compute Directive
Compute directive에는 parallel, kernels, serial 이 있다.
Parallel
- parallel directive는 여러 gangs를 병렬로 실행하며, 각 gang은 여러개의 workers를 가지며, 각 worker에는 SIMD나 vector operation을 위한 vector가 있다.
- Syntax
// C/C++
#pragma acc parallel [clause [[,] clause]…] new-line
{ structured block }
// FORTRAN
!$acc parallel [clause [[,] clause]…]
structured block
!$acc end parallel- clause는 아래와 같다.
- if( condition )
- self [( condition )]
- default( none )
- default( present )
- device_type or dtype( [ * | device-type-list ] )
- async [( expression )]
- wait [( expression-list )]
- num_gangs( expression )
- num_workers( expression )
- vector_length( expression )
- reduction( operator: list )
- private( list )
- firstprivate( list )
- copy( list )
- copyin( [readonly:] list )
- copyout( list )
- create( list )
- no_create( list )
- present( list )
- deviceptr( list )
- attach( list )
Kernels Directive
- kernels directive는 device에서 실행되는 일반적으로 kernel operations의 순서인 loops를 구성한다.
- Syntax
// C/C++
#pragma acc kernels [clause [[,] clause]…] new-line
{ structured block }
// FORTRAN
!$acc kernels [clause [[,] clause]…]
structured block
!$acc end kernels- clause는 아래와 같다.
- if( condition )
- self [( condition )]
- default( none )
- default( present )
- device_type or dtype( [ * | device-type-list ] )
- async [( expression )]
- wait [( expression-list )]
- num_gangs( expression )
- num_workers( expression )
- vector_length( expression )
- copy( list )
- copyin( [readonly:] list )
- copyout( list )
- create( list )
- no_create( list )
- present( list )
- deviceptr( list )
- attach( list )
Serial Directive
- serial directive는 device에서 serial하게 실행할 loops 또는 코드를 구성한다.
- Syntax
// C/C++
#pragma acc serial [clause [[,] clause]…] new-line
{ structured block }
// FORTRAN
!$acc serial [clause [[,] clause]…]
structured block
!$acc end serial- clause는 아래와 같다.
- if( condition )
- self [( condition )]
- default( none )
- default( present )
- device_type or dtype( [ * | device-type-list ] )
- async [( expression )]
- wait [( expression-list )]
- reduction( operator: list )
- private( list )
- firstprivate( list )
- copy( list )
- copyin( [readonly:] list )
- copyout( list )
- create( list )
- no_create( list )
- present( list )
- deviceptr( list )
- attach( list )
When the program encounters an accelerator serial construct, one gang of one worker with a vector length of one is created to execute the accelerator serial region sequentially.
Data directive
- data directive는 device에서 데이터에 접근할 수 있는 영역을 정의한다.
- Syntax
// C/C++
#pragma acc data [clause[[,] clause]…] new-line
{ structured block }
// FORTRAN
!$acc data [clause[[,] clause]…]
structured block
!$acc end data- clause는 아래와 같다.
- if( condition ) : When the condition is zero or .FALSE. no data will be allocated or moved to or from the device.
- copy( list )
- copyin( [readonly:] list )
- copyout( list )
- create( list )
- no_create( list )
- present( list )
- deviceptr( list )
- attach( list )
- default( none )
- default( present )
Enter Data Directive / Exit Data Directive
- enter data : exit data directive를 통해 data를 deallocation을 하기 전까지 device memory에 데이터를 할당 및 이동하는데 사용된다.
- exit data : enter data directive를 통해 생성된 데이터를 device memory에서 이동하고 deallocate.
- Syntax
// C/C++
#pragma acc enter data [clause[[,] clause]…] new-line
// FORTRAN
!$acc enter data [clause[[,] clause]…]- enter data의 clause는 아래와 같다.
- if( condition )
- async [( expression )]
- wait [( expression-list )]
- copyin( list )
- create( list )
- attach( list )
- exit data의 clause
- if( condition )
- async [( expression )]
- wait [( expression-list )]
- finalize : Sets the dynamic reference count to zero.
- copyout( list )
- delete( list )
- detach( list )
data 관련 clause 설명
- no_create : data가 이미 디바이스 메모리에 있으면 ref.count 증가하고 copy, 아니면 아무일도 하지않음. 이후에 data를 사용하려면 local memory address를 사용함.
- deviceptr : C/C++에서 list에 있는 엔트리들은 반드시 디바이스 주소를 가지는 포인터 변수여야함.
- attach : list에 있는 포인터가 이미 their targets??에 attach되어있다면 attachment count 증가, 아니라면 device pointer를 their device target에 attach하고 attachment count를 1로 만듬. region을 나가면 카운트수를 줄이고, 그 카운트가 0이 되면 detach.
- detach : finalize clause와 같이 사용되지 않을 때에는 attachment count를 감소, 같이 사용되었을 경우 0으로 만듬. 그 카운트가 0이 되면 detach.
Loop directive
- applies to the immediately following loop or tightly nested loops
- describes the type of device parallelism to use to execute the iterations of the loop
- Syntax
// C/C++
#pragma acc loop [clause [[,] clause]…] new-line
// FORTRAN
!$acc loop [clause [[,] clause]…]- clause
- collapse( n ) : Applies the associated directive to the following n tightly nested loops.
- seq : Executes the loop or loops sequentially.
- auto : Instructs the compiler to analyze the loop or loops to determine whether it can be safely executed in parallel, and if so, to apply gang, worker or vector parallelism.
- independent
- tile( expression-list )
- device_type or dtype( [ * | device-type-list ] )
- private( list )
- reduction( operator: list )
LOOP CLAUSES WITHIN A PARALLEL CONSTRUCT OR ORPHANED LOOP DIRECTIVE
gang : Shares the iterations of the loop or loops across the gangs of the parallel region.
worker : Shares the iterations of the loop or loops across the workers of the gang.
vector : Executes the iterations of the loop or loops in SIMD or vector mode.
LOOP CLAUSES WITHIN KERNELS CONSTRUCT
gang [( num_gangs )] : Executes the iterations of the loop or loops in parallel across at most num_gangs gangs.
worker [( num_workers )] : Executes the iterations of the loop or loops in parallel across at most num_workers workers of a single gang.
vector [( vector_length )] : Executes the iterations of the loop or loops in SIMD or vector mode, with a maximum vector_length.
Device-Specific Clauses
device_type :
- Clauses associated with a device_type apply only when compiling for the device type named.
- Clauses associated with a device_type that has an asterisk argument apply to any device type that was not named in any device_type on that directive.
num_gangs clause가 없을 경우
If the clause is not specified, an implementation-defined default will be used; the default may depend on the code within the construct. The implementation may use a lower value than specified based on limitations imposed by the target architecture.
num workers clause가 없을 경우
If the clause is not specified, an implementation-defined default will be used; the default value may be 1, and may be different for each parallel construct or for each kernel created for a kernels construct. The implementation may use a different value than specified based on limitations imposed by the target architecture.
vector length clause가 없을 경우
If the clause is not specified, an implementation-defined default will be used. This vector length will be used for loop constructs annotated with the vector clause, as well as loops automatically vectorized by the compiler. The implementation may use a different value than specified based on limitations imposed by the target architecture.
Attachment Count
Since multiple pointers can target the same address, each pointer in device memory is associated with an attachment counter per device.
loop
A loop construct with the gang clause transitions a compute region from gang-redundant mode to gang-partitioned mode.
--------------------------------------------------------------------------------------------------------------
커널지시자
SAXPY C 코드
restrict : 아래에 x, y가 같은 주소를 가리키면(겹치면) 루프를 조각으로 나누는데 어렵게 만듬. 없으면 루프는 느리게 동작. 표준 c에서 이 키워드는 다른 어떤것도 y를 가리키지 않는다. y를 독립적으로 취급해라. y가 다른 변수와 섞이는 걸 신경쓰지마라. 는 뜻. x와 y는 관련이 없고 메모리에서 겹치지 않는다고 컴파일러에게 알려줌.
pgcc에서 빌드할 때 Minfo 옵션은 피드백용 (Accelerator kernel generated)
그러나 루프 반복이 비독립적일 경우, 루프 반복을 쪼개서 각각 따로 작업할 수 없는 경우, 데이터 의존성을 가짐
이럴 땐 kernel을 추가해도 속도를 못 냄
라플라스 솔버, 금속 판 위 온도에관한 간단한 물리학 문제
라플라스 방정식, 기우기, 어떤 필드의 델의 제곱은 0과 같다.. 원천(source)이 없는 문제를 푼다는 것이다. 여기 있는 금속 판 가운데는 열 원천이 없다.
정상상태 열 방정식 = 금속판에 온도를 적용한다. 그러면 그것이 평형상태에 도달한 후의 최대 정상 상태는 어떻게 나타나는가?를 표현한다.
자코비 반복법, 4개의 이웃을 더하고 4로 나눈다.. 값이 변화하기를 멈출 때까지 금속판위 점들을 계속 계산
------------------------------------
