Fractional GPUs : Software-based Compute and Memory Bandwidth Reservation for GPUs

Abstract

 GPU는 큰 성능 이익때문에 자율주행같은 real-time system(이하 RT, RTS)에 점차 사용되고있다. 하지만 RTS는 예측가능한 성능을 요구한다. Nvidia는 최근 Multi-Process Service(MPS)라 불리는 closed-source kernel module 형태의 feature를 제공했다. 하지만 MPS는 GPU의 compute자원을 분할하는 능력만을 제공할뿐 shared memory hierarchy내 어플리케이션간 충돌을 피하는 메커니즘은 제공하지 않았다. 실험상 최악의 경우에, 동식 동작하는 GPU task로부터의 간섭 때문에 read/write transaction이 10x 이상 느려짐을 확인할수 있었다.

 

 이 논문에서, performance isolation으로 GPU 작업의 병렬 실행을 할수있게 GPU의 compute and memory resources를 분할하는 software-only mechanism인 Fractional GPUs(FGPUs)를 썼다. GPU memory hierarchy는 page coloring을 위해 잘 짜여진 CPU와는 다르다는것을 알아냈다. 발견을 토대로 여러 Nvidia GPU에서 L2 cache와 DRAM 둘을 분리할수 있었다. 나아가, MPS를 사용한것보다 compute resources를 분리하는데 더 나은 전략이 있다는것을 보인다. FGPU는 더 나은 isolation을 제공하기 위해 이 전략과 memory coloring을 결합한다. MPS와 비교했을때, FGPU는 multi-tenancy 환경의 어플리케이션 런타임부분에서 평균 variation을 135%에서 9%까지 줄였다. FGPU를 사용할 multiple application들을 완벽하게 하기 위해서Caffe를 포팅했다.

 

Introduction

 GPU는 새로운 세대와 아키텍쳐에 더 영향을 미치고 있다. 또한, real-time systems는 점점 GPU를 사용하는 어플리케이션들을 확장중이다. 이는 특히 GPU로부터 많은 물량의 병렬연산을 착취하는 자율주행같은 도메인에서 머신러닝 어플리케이션의 인기를 증가시키고 있다. 하나의 어플리케이션은 전체의 GPU를 사용할순 없을것이다. 하지만 여러 어플리케이션들은 GPU를 사용하는데 장점이 있다. 이런 두 트렌드가 다수의 어플리케이션들을 GPU에서 동작하게 하는것을 중요하게 만든다. RT어플리케이션들이 엄격한 데드라인을 가지고 있기 때문에 특히 안전에 중요한 어플리케이션들을 위해 워스트 시나리오의 케이스라도 GPU들은 동시에 예측가능한 어플리케이션 퍼포먼스를 제공받기를 원한다. 

 

 이런 요구들을 충족시키기 위해서, Nvidia는 GPU에서 여러 어플리케이션들이 co-run 할수 있게 MPS를 지원한다. 그들은 최근에 MPS에 개발자가 어플리케이션 한개당 사용가능한 compute bandwidth 제한을 각 어플리케이션에 사용가능한 GPU thread 수에 대해 상한선을 지정할 수 있게 하는 새 QoS feature를 추가했다. 이것은 사용가능한 스레드의 비율을 제한하는것이 어플리케이션간 파괴적인 간섭을 줄일것이다는 아이디어이다. 하지만, MPS QoS feature에 관한 세가지 이슈가 있다. : 1) 오직 최신 Nvidia GPU에서 사용 가능하다, 2) 이 소스 코드를 사용할 수 없으므로 블랙박스로 사용할 수 없으며 잠재적으로 신뢰할 수 없다, 3)MPS는 GPU의 컴퓨팅 리소스를 분할하는 것만을 허용하며, 메모리 계층 내에서 애플리케이션 간 충돌을 피하기 위한 메커니즘을 제공하지 않는다.

 이전 작업 [13] [19] [20] [26] [30] [31] [35] 에서는 CPU의 서로 다른 코어에서 실행되는 두 애플리케이션이 다음과 같은 이유로 인해 메모리 계층(주로 공유 캐시 및 DRAM)의 충돌로 인해 여전히 서로의 런타임에 영향을 미칠 수 있다는 것을 보여주었다.

 

1) Cache set conflicts in the shared cache,
2) Miss Status Holding Registers (MSHR) contention in
the shared cache,
3) Reordering of requests in the memory controller,
4) DRAM bus contention, and
5) Row buffer conflicts in DRAM.

 

[30]은 이러한 CPU 충돌 원인을 조사하는 포괄적인 연구 세트를 제시한다. [30]을 포함한 각 논문들은 CPU에서 이러한 충돌의 원천들 중 하나 이상을 제거하기 위한 메커니즘을 설계한다. 그러나 어떤 작업도 이러한 충돌의 원천들을 모두 제거할 수 없었고, 게다가 GPU에서는 그러한 작업이 시도되지 않았다. 본 논문에서는 먼저 GPU 메모리 계층을 조사한 다음, 이러한 모든 충돌을 해결할 페이지 컬러링 메커니즘을 제시하며  GPU에서 격리된 상태로 애플리케이션을 실행한다.

 

 본 논문의 주요 공헌은 다음과 같다. :

 

1) 우리는 예측가능한 성능을 갖는 다수의 어플리케이션을 병렬로 실행하게 하는 소프트웨어 기반 GPU 파티셔닝 메커니즘을 구현했다. 구체적으로, 우리는 단일 GPU를 더 작은 부분적인 GPU로 분할하기 위해 컴퓨팅과 메모리 대역폭 분리 메커니즘을 결합한다. 우리는 각각 높은 수준의 분리도를 가진 부분적 GPU 내에서 여러 애플리케이션을 병렬로 실행할 수 있다는 것을 보여준다.

2) 우리는 Nvidia GPU에 있는 L2 캐시와 D램의 공개되지 않은 아키텍처 세부사항을 역설계하기 위한 일반 알고리즘을 구축한다. 이전 연구[17] [21] [22]는 다양한 Nvidia GPU의 메모리 계층 구조를 분석하려고 시도했지만, 어떤 작업도 L2 캐시/DRAM의 구조를 설명할 수 없었다(일찍은 GPU [32]). [33]). 마이크로벤치마크를 통해 복수의 Nvidia GPU에서 L2 캐시와 D램의 아키텍처를 역설계한다. 우리는 Nvidia GPU L2 캐시가 기존의 CPU L2 캐시와 같은 방식으로 구조화되지 않았으며, 보다 정교하다는 것을 보여준다.

3)우리가 아는 한, 우리는 GPU에서 페이지 컬러링을 가장 먼저 구현했다. 우리는 Nvidia GPU의 L2 캐시/DRAM 구조가 페이지 컬러링에 잘 적합하다는 것을 보여준다. Nvidia 장치 드라이버가 부분적으로 닫혔음에도 불구하고, 우리는 페이지 컬러링을 실행할 수 있다.

4)우리는 FGPU abstraction를 이용하기 위해 Caffe를 포팅시킨다. Caffe는 DNN(Deep Neural Networks)에 널리 사용되는 GPU 프레임워크다.FGPUs에 Caffe를 포팅하는 것은 예측 가능한 런타임과 병행하여 여러 DNN을 단일 GPU에서 실행할 수 있도록 한다.

 

 본 논문은 고성능 컴퓨팅을 위한 선도적인 플랫폼이기 때문에 Nvidia GPU (그리고 대응하는 SDK, 즉 CUDA [23])에 초점을 맞추고 있다. 구체적으로 GTX 1070, GTX 1080, 테슬라 V100을 선택했다. 이러한 GPU에는 최신 아키텍처가 있고 드라이브 PX2 및 드라이브 PX 페가수스[1]와 같은 새로운 자동차 GPU 보드가 있다. 우리는 우리의 역설계 알고리즘이 다른 GPU에 적용될 만큼 충분히 일반적이라고 생각한다. 우리는 이것을 앞으로의 일을 위해 남겨 둔다.

 나머지 논문은 다음과 같이 정리되어 있다. 섹션 II는 CPU와 GPU 아키텍처에 필요한 배경을 제공한다. 섹션 III에서는 컴퓨팅 격리 및 메모리 대역폭 분리라는 용어를 정의한다. 섹션 IV와 V에서는 Nvidia GPU의 FGPU 사이에 이러한 격리 속성을 달성하는 방법을 설명한다. 섹션 VI는 실험 결과를 보고하고 관련 작업과 결론을 설명한다.

 

Background

 이 섹션에서는 Nvidia의 인기 있는 CUDA 플랫폼의 용어를 사용하는 GPU 아키텍처에 대해 몇 가지 필요한 배경을 제공한다. 우리는 일관성을 위해 나머지 논문들의 CUDA 용어를 사용한다.

 

A. GPU Compute Hierarchy

 

GPU는 수십 개의 스트리밍 멀티프로세서(SM)로 구성되어 있으며, 각각은 많은 수의 하드웨어 스레드를 가지고 있다. 그러나 SM은 프로그래머로부터 추상화된다. 대신, CUDA는 블록이라고 불리는 소프트웨어 추상화를 제공하며, 각각은 스레드의 집합이다. GPU에서 코드를 실행하기 위해, CUDA 프로그래머는 1) 커널(a set of instructions), 2) 사용할 블록의 수, 3) 블록 내의 스레드 수를 지정한다. 그런 다음 CUDA는 각각 동일한 수의 스레드 및 각 스레드가 동일한 커널(다른 데이터에서 작동하지만)을 실행하는 지정된 수의 블록을 발사한다. 하드웨어 GPU 스케줄러는 단일 SM이 여러 블록을 실행하면서 블록과 SM 간의 매핑을 결정한다. 그러나 SM은 한 번에 일정한 수의 블록만 실행할 수 있으므로, 필요한 하드웨어 리소스를 SM에서 사용할 수 있을 때까지 일부 블록이 대기열에 포함될 수 있다. Nvidia GPU 하드웨어 스케줄러에 대한 자세한 내용은 공개적으로 알려지지 않았다. CPU에 있어서 GPU는 CUDA API를 통해 CUDA 커널을 시작할 수 있는 co-프로세서처럼 보이지만 스케줄링 알고리즘을 제어할 수 없다.

 나머지 논문에서 커널은 CUDA 커널을, 장치 드라이버는 OS 커널 모듈을, 스레드는 GPU 하드웨어 스레드를, 프로세스는 OS 프로세스를 가리킨다.

 

B. GPU Memory Hierarchy

 

 오늘날 GPU는 메모리 계층 구조에서 세 가지 주요 수준을 가지고 있다. :
1) L1 캐시 2) L2 캐시, 3) D램. L1 캐시는 SM의 전용 캐시로 해당 SM의 모든 스레드가 공유한다. 모든 SM은 L2 캐시와 D램을 공유한다. Nvidia는 GPU의 메모리 아키텍처에 대한 세부사항을 밝히지 않는다. 따라서, 우리는 CPU에서 이러한 메모리의 전통적인 계층구조에 대해 간단히 논한다.
 캐시는 DRAM보다 접속 속도가 빠르지만 크기가 작고 자주 접속하는 데이터를 캐싱하는 데 사용되는 메모리 모듈이다. 캐시는 m 세트로 나뉘며, 각 세트는 캐시 라인이 연속적인 워드의 집합이 되는 n 캐시 라인이 있다. 메모리 트랜잭션과 관련된 주소는 잠재적으로 거짓말을 할 수 있는 캐시 세트를 지시한다. 요청된 메모리 블록이 이 세트의 캐시 라인 중 하나에 있으면 캐시 적중이다. 그렇지 않으면 캐시 미스에서 캐시 라인이 이 세트에서 제외되도록 선택되고 요청이 DRAM으로 전달된다. DRAM에서 데이터를 가져온 후에는 제거된 캐시 라인 대신 새 캐시 라인이 이 세트에 배치된다.

 D램은 a set of chips 로 구성되며, 각각은 banks으로 더 나누어져 있으며 각 뱅크는 여러 rows로 구성되어 있다. 각 뱅크에는 마지막 읽기 행을 캐시하는 데 사용되는 row buffer가 있다. 실제 주소를 기준으로 D램에 요청이 들어오면 특정 칩 안에 있는 특정 뱅크로 전달된다. 요청된 주소가 캐시된 행에 있는 경우, 데이터는 행 버퍼에서 가져온다. 그렇지 않으면 행 버퍼의 행이 제거되고, 액세스 중인 새 행이 버퍼로 가져오기 시작한다. 그런 다음, 요청은 행 버퍼의 데이터를 사용하여 이행된다. 행 버퍼 miss는 행 버퍼 제거의 추가 오버헤드로 인해 행 버퍼 hit보다 훨씬 느리다.

 

 주소에서 캐시 세트 또는 D램 chip/bank/row로의 매핑은 고려 중인 하드웨어에 고유한 별도의 해시 함수를 통해 수행된다. Nvidia는 메모리 주소가 DRAM에 분산시키기 위해 새로운 GPU[11]에 해시되어 있다고 보고하지만 이러한 해시함수는 보고하지 않는다.

 본 문서의 나머지 부분에서는 DRAM은 GPU DRAM, cache sets는 GPU cache sets를, banks은 GPU DRAM bank를 가리킨다.

 

Performance Isolation on GPU

단일 GPU 의 두 GPU 파티션 ,[P1 P2], 그리고 두 커널 세트, [K1 K2]가 있을때, for any k1 ∈ K1 and k2 ∈ K2, k1(k2)의 런타임이 다음 시나리오에서 변하지 않을 경우, 파티션 간에 완벽한 성능 격리가 된다고 한다. :

 1) k1(k2)는 P1(P2)에서 실행되고 P2(P1)는 비어있다.

 2) k1은 P1에서 실행되고 k2는 P2에서 병렬로 실행된다.

즉, k2는 k1의 런타임(k2는 k1의 자원을 빼앗지 않는다)에 아무런 영향을 주지 않고, 그 반대에도 해당된다.

자원은 크게 compute(thread/SMs)과 memory resources(caches/DRAM/memory buses)의 두 종류가 있다. compute isolation를 고려하면서, Section 1에 언급된 요인에 의한 메모리 계층의 충돌을 통한 k2의 메모리 트랜잭션(읽기/쓰기)이 k1의 런타임에 미치는 영향을 무시할 수 있다. P1과 P2도 그들 사이의 메모리 대역폭 격리(memory bandwidth isolation)를 달성하려면, k1과 k2의 메모리 트랜잭션이 서로 분리되는 것은 필수적이다.

 

Implementation of Compute Isolation on GPU

 

단일 GPU는 여러 파티션 Pi (1 ≤ i ≤ n, where n is the number of partitions)로 분할된다. 각 애플리케이션에 대해 해당 애플리케이션의 모든 커널은 단일 커널 세트 Ki에 할당된다. 모든 커널들 ki(∈ Ki)은 어플리케이션 내의 모든 커널이 단일 GPU 파티션에서 실행되도록 보장하면서 Pi에서 실행된다. 각 파티션 Pi에 SM의 분리 세트를 할당하여 각 파티션 쌍 간에 compute isolation을 달성할 수 있다. 예를 들어, GPU가 10개의 SM을 가지고 있을 경우, 1 ~ 5번의 SM을 P1에 할당하고 6 ~ 10번의 SM을 P2에 할당할 수 있다. k1이 실행되면 1~5번의 SM만 사용할 수 밖에 없게 된다. 같은 커널 세트의 다른 커널은 시간간격 방식으로 차례로 실행되지만 다른 세트의 커널은 병렬로 실행할 수 있다. 여러 애플리케이션을 동일한 파티션에 할당할 수 있으며, 이 경우 서로 다른 파티션에서 실행되는 애플리케이션은 서로 컴퓨팅 격리되지만 동일한 파티션에서 실행되는 애플리케이션은 서로 영향을 받는다. SM이 파티션에 할당된 컴퓨팅 유닛 중 가장 작기 때문에 최대 컴퓨팅 파티션 수는 SM의 총 수와 동일하다. 현재 Nvidia의 GPU는 커널에 특정 SM을 할당하기 위한 하드웨어 지원을 가지고 있지 않다. 따라서, 우리는 소프트웨어 기반 컴퓨팅 분리를 구현한다. 당사의 구현은 [34]에 기반을 두고 있으며, 성능을 개선하기 위한 몇 가지 확장 기능도 있다. 이제 SM affinity를 달성하기 위한 전반적인 메커니즘에 대해 간략하게 설명하겠다. :

 

 1) 각 애플리케이션은 초기화 중에 실행하고자 하는 파티션을 식별한다.

 2) 커널 런치에서는 프로그래머가 지정한 대로 각 블록에 nt 스레드가 있는 nb 블록을 론칭하는 대신, npb persistent block을 nt 스레드로 런치한다. npb는 npb 블록이 모든 SM의 모든 스레드를 차지하도록 계산된다. npb =(Threads in SM / nt ) * Number of SM1.

 3) 각 persistent block은 해당 ID를 가져와 실행 중인 SM을 식별한다. 이 SM이 애플리케이션의 파티션에 할당된 SM의 set에 있지 않으면 블록은 아무런 작업도 하지 않고 종료된다. 남은 persistent block(예: npb')은 모든 올바른 SM에서 실행되고 있다.

 4) 0부터 nb-1까지 블록 인덱스의 중앙집중식 큐는 원자 운용을 이용한 GPU에서 구현된다. npb' persistent block은 이 중앙 집중화된 대기열에서 블록 인덱스를 팝업하여 원래 nb 블록을 실행한다. 

 

Original Code
// CUDA Kernel
// Implements c [ i ] = a [ i ] + b [ i ]
__global__ void vectorAdd ( float *a , float *b , float *c ) {
	int i = blockIdx * blockDim + threadIdx ;
	c[i] = a[i] + b[i];
}
...
// CPU code
nt = 256; // Number of threads
nb = num_elements / nt; // Number of blocks
vectorAdd<<<nb , nt>>>(A, B, C); // Run nb blocks
....

Modified Code
__global__ void FGPU_DEFINE_KERNEL( vectorAdd , float *a,
	float *b , float *c ) {
	int blockIdx;
	FGPU_BLOCK_INIT ();
    
	FGPU_FOR_EACH_BLOCK( _blockIdx ) {
		int i = blockIdx * blockDim + threadIdx;
		c[i] = a[i] + b[i] ;
	}
}
...
fgpu_init( partition_id )
...
FGPU_LAUNCH_KERNEL( vectorAdd , nb , nt , A, B, C)
...

Fig. 1: Comparison of original vector addition code and modified code to support compute isolation using FGPU API. blockIdx, blockDim, threadIdx are CUDA-supported keywords

 

 Figure 1은 두 개의 벡터, A와 B를 추가하고 결과를 출력 벡터 C로 반환하는 예제 어플리케이션인 벡터 덧셈을 보여준다.우리는 원본 코드와 FGPU API를 사용하는 수정된 코드를 보여준다. 원래 코드에서 CPU는 벡터 각 블록에 nb 블록과 nt 스레드를 갖는 vectorAdd kernel을 런치한다. 커널 안에서 각 스레드는 단일 요소의 인덱스를 계산하고 입력 벡터의 해당 요소에 대한 합계를 계산한다. blockIdx와 threadIdx는 각각 현재 블록과 스레드의 인덱스를 반환하는 CUDA가 제공하는 기초 요소이다. blockDim은 블록의 스레드 수를 지정한다.

 수정된 코드에서는 FGPU API에서 제공하는 매크로를 사용하도록 커널 론칭 코드와 커널 코드를 수정한다. 애플리케이션 시작 시 CPU에서 fgpu_init()를 사용하여 프로그래머는 현재 애플리케이션과 연결할 GPU 파티션을 표시한다. 커널 론칭 시 FGPU_LAUNCH_KERNEL() 매크로는 커널, 론치 파라미터 및 커널 인수를 입력으로 한다. nb 블록을 런칭하는 대신 vectorAdd의 npb persistent block이 런칭된다. 각 블록 내에서 FGPU_BLOCK_INIT()는 두가지 기능을 수행한다. :

1) SM이 파티션에 없으면 현재 persistent block을 종료하고

2) 블록 인덱스가 포함된 큐를 초기화한다.

나머지 영구 블록은 FGPU_FOR_EACH_BLOCK()을 사용하여 중앙 집중식 큐에서 한 번에 하나의 블록 인덱스를 가져오며 이 블록을 실행한다

 현재의 구현에는 애플리케이션의 소스 코드의 수정이 필요하지만, 우리는 필요한 변경 사항이 미미하기 때문에 컴파일러 지원 코드 변환 기법을 사용하여 프로세스를 자동화할 수 있다고 믿는다. 우리는 이것을 장래의 일로 남겨 둔다.

 Nvidia는 일반적인 API를 통해 서로 다른 애플리케이션의 커널이 함께 실행되는 것을 허용하지 않는다. 따라서 [34]에서는 애플리케이션을 단일 애플리케이션으로 병합하도록 수정해야 한다. 구현 시, 우리는 Nvidia MPS를 사용하고 애플리케이션을 병합할 필요 없이 이러한 제약을 무시한다. 이것은 우리의 접근법을 사용하기 쉽게 만든다.

 MPS는 또한 프로그래머들이 각 어플리케이션에 사용할 수 있는 스레드 수에 대한 상한 값을 지정할 수 있는 QoS 기능을 가지고 있다. 그러나 FGPU는 compute isolation를 위해 MPS를 사용하지 않는다. MPS는 스레드 세분화에서 파티셔닝을 하기 때문에 동일한 SM에서 여러 애플리케이션을 실행할 수 있으므로 SM 내의 제한된 리소스(예: 스레드 레지스터)로 인해 애플리케이션이 충돌할 수 있다. 우리는 섹션 VI-A의 MPS QoS 기능보다 SM 선호도 기반 컴퓨팅 분리(SM affinity based compute isolation)가 더 낫다는 것을 보여준다.

 

Memory Bandwidth Isolation on GPU

A. Introduction

 

 Memory bandwidth isolation는 충돌을 피하기 위해 메모리 계층을 분할해야 한다. 이를 달성하기 위해서는 엔비디아 GPU 메모리 계층, 특히 SM 간에 공유되는 주요 메모리 구성요소인 L2 캐시와 D램을 이해해야 한다. 아쉽게도 이러한 세부사항은 공개되지 않는다. 다음 절에서는 GTX 1080을 예로 들어 Nvidia GPU의 L2 캐시와 D램 계층 구조를 역설계하기 위한 실험과 알고리즘을 공식화한다. 우리는 나중에 유사한 메모리 계층이 다른 Nvidia GPU 칩셋에도 존재하는지 검증한다.

 Nvidia GPU에는 애플리케이션에서 선택적으로 사용할 수 있는 다른 유형의 메모리[24], 즉 L1 cache, shared memory, constant memory 그리고 textured memory도 있다. L1 캐시와 공유 메모리는 각 SM에 대해 비공개적이므로 각 파티션에 SM의 분리 세트가 있으므로 메모리 대역폭 예약과 관련이 없다. 상수 메모리는 크기가 작다(GTX 1080의 경우 64KB). 따라서, 우리는 지속적인 기억력에 대해 큰 논쟁을 기대하지 않는다. 우리는 텍스처 메모리를 둘러싼 충돌의 영향을 향후 작업으로 탐구한다(평가에서 사용된 텍스처 메모리 적용 없음).

 

 

B. Reverse-Engineering

 

 캐시 및 DRAM 계층을 역설계하기 위한 이전 알고리즘[19]은 하드웨어에 대해 확실한 가정을 한다. 예를 들어 그림 2는 [30]에서 역설계한 Intel Core i7-2600의 매핑 기능을 보여준다. 이 매핑을 역설계하는 데 사용되는 알고리즘은 물리적 주소의 끝의 2비트에서 bank index를 도출하기 위해 XOR로 설정되었다고 가정했다. 이 알고리즘은 Intel Core i7-7700과 같은 이 제약 조건을 충족하지 못하는 새로운 CPU 아키텍처에서 실패한다. 더욱이, GPU의 메모리 아키텍처는 공개적으로 알려져 있지 않고 우리가 보여 있듯이, 기존의 CPU 메모리 계층 구조와는 상당히 다르기 때문에, 이러한 접근 방식은 GPU에서는 작동하지 않는다. 따라서 CPU와 GPU에서 모두 작동하는 일반 알고리즘을 제안한다.

주소가 액세스하는 DRAM bank/캐시 set의 주소를 매핑하는 해시 함수를 역설계하기 위한 알고리즘의 주요 원칙은 다음처럼 요약할 수 있다:

 

1) 하드웨어와 무관한 속성을 이용하여 동일한 DRAM 뱅크/캐시 세트에 놓여 있는 여러 쌍의 주소를 찾아라.
2) 가능한 모든 해시함수의 전체 목록을 작성한다.
3) Brute Force를 이용하여 모든 쌍의 주소를 동일한 DRAM 뱅크/캐시 세트에 매핑하는 유효한 해시 함수를 찾는다.

 

하드웨어에 대한 추정은 다음과 같다:

1) 해시함수는 물리적 주소를 입력으로 삼는다. 이것이 꼭 필요한 가정은 아니며 우리의 알고리즘도 가상 주소로도 쉽게 작동할 수 있지만, 우리는 이 가정이 우리가 테스트한 모든 GPU에 유효하다는 것을 발견했다.
2) 해시함수는 물리적 주소(AND, OR, XOR)의 비트에 대한 bitwise operation으로 제한된다. 이러한 해시함수는 하드웨어에서 구현되기 때문에 빠르고 간단할 필요가 있다. 이 가정은 가능한 고유 해시 함수의 최대 수를 제한한다.
3) GPU 캐시는 eviction policy로 LRU(최소 사용)를 사용한다. 이러한 가정에 대한 이유는 섹션 V-C 및 V-D에 설명되어 있다.

 

C. Reverse-Engineering of DRAM Bank Addressing

 

 이 섹션에서는 물리적 주소를 DRAM bank로 매핑하는 해시함수를 발견하기 위한 알고리즘을 설명한다. 알고리즘은 동일한 뱅크에 있지만 다른 행(행 버퍼 퇴출을 유발하는)에 있는 두 개의 주소에 액세스하는 것이 다른 뱅크에 있는 두 개의 주소에 액세스하는 것보다 눈에 띄게 느려진다는 사실을 이용한다. 행버퍼 eviction를 보려면 읽기/쓰기 트랜잭션이 D램에 도착해야 한다. 이런 일이 일어나려면 캐시를 우회해야 한다. L2 캐시를 비활성화하거나 플러시하는 CUDA 지침을 찾을 수 없었다. 따라서 GPU의 캐시가 LRU (Assumption (3))라고 가정할 경우, 우리는 충분한 모의 데이터를 읽음으로써 암묵적으로 캐시를 지울 수 있다. 가짜 데이터 읽기 양은 적어도 캐시 크기와 같아야 한다.이는 플러시 캐시(예: x86의 clflush)를 사용하는 명시적 지침이 존재하는 CPU와는 다르다. 우리가 테스트한 GPU에서는 이러한 가정이 강하게 유지되는 것 같다.

 

/* Code Executed by CPU */
function PrintBankMappingFunctions
	/* Find size of free memory on GPU */
	S ← GetGpuFreeMem();
	/* Allocate contiguous physical memory chunk */
	<V irtStart, PhyStart> ← AllocGpuPhyMem(S)
	Pairs ← ∅
	/* Collect all pairs */
	for Offset ← 1 to S do
		V addr ← Offset + V irtStart
		Paddr ← Offset + PhyStart
		/* Virtual Primary and Secondary Addresses */
		if IsGpuRowEvicted(V irtStart,V addr) then
			Pairs.append(PhyStart, Paddr)
		end if
	end for
	F ← GenerateAllHashF unctions()
	/* Test all permutations of mapping function */
	for each func ∈ F do
		if IsV alidMapping(func,Pairs) then
			Print(func)
			return
		end if
	end for
end function
/* Code Executed by GPU */
function IsGpuRowEvicted(PrimaryAddr, SecondaryAddr)
	/* Clear cache by reading enough spurious data */
	ClearGpuCache()
	StartTime ← GetClock()
	/* Access Primary and Secondary Addresses */
	data ← PrimaryAddr
	data ← SecondaryAddr
	AccessTime ← GetClock() − StartTime
	if AccessTime ≥ Threshold then
		return True
	else
		return False
	end if
end function

Algorithm 1. Hash function to reverse-engineer GPU DRAM bank bits

 

 Algorithm 1은 역설계용 유사코드를 bank index 비트에 대한 해시함수를 제공한다. Assumption (1)에 따라 해시함수를 찾기 위한 물리적 주소를 알아야 한다. CUDA는 지정된 가상 주소에서 물리적 주소를 쿼리할 수 있는 API를 제공하지 않기 때문에, (a) GPU 메모리의 연속적인 물리적 청크를 할당하고, (b) 이 청크의 물리적 주소 변환 매핑에 가상 주소를 생성하고, (c) 시작 및 물리적 가상 주소를 반환하도록,  new API(shown as AllocGpuPhyMem() in the pseudo code)를 추가하므로써 Nvidia device driver3를 수정하였다. 이를 통해 청크 내의 모든 가상 주소(line 1-11)에 대한 물리적 주소를 계산할 수 있다. 루프(line 9-16)에서는 기본 주소를 시작 주소에 고정하고 액세스 시간의 급증을 식별하여 행 버퍼 eviction을 유발하는 여러 쌍의 주소를 찾으면서 보조 주소를 변경한다. 그런 다음 가능한 모든 해시함수(based on Assumtion (2))를 검색하여 유효한 해시함수(즉, 모든 쌍 중에서 대해 양쪽 주소에 동일한 뱅크를 할당하는 해시함수)를 찾는다. 충분한 쌍을 테스트할 경우, 하나의 해시함수만 유효한 것으로 식별되어야 한다.

 

 우리는 한 가지 기능만 남았을 때 후보 funcfion을 없애고 멈추도록 테스트할 주소 쌍을 지능적으로 식별함으로써 이 brute-force 접근방식을 더욱 개선할 수 있다. 우리의 모든 역설계 실험은 우리가 실험한 모든 GPU에서 1시간 이내에 끝났다.

 

 Figure 3은 주소 쌍의 하위 집합에 대한 액세스 시간을 나타내고 Figure 4는 히스토그램을 보여준다. 대부분의 접속 시간은 495-520 사이클 사이에 있다. 임계값(552 사이클)보다 많은 액세스 시간을 갖는 페어는 거의 없다. 이 바이모달 배분은 행 버퍼 miss가 GPU에서 행 버퍼 hit과 차별화 하도록 상당한 액세스 시간 페널티를 발생시키는 것을 검증한다. Figure 6(b)는 GTX 1080에 대해 우리가 실험적으로 발견한 해시함수를 상세히 기술한다. 우리는 이 해시함수를 섹션 V-F에서 실험적으로 검증할 것이다.

 

 

 

 

/* Code Executed by CPU */
function PrintCachelineMappingFunctions
	/* Find size of free memory on GPU */
	S ← GetGpuFreeMem();
	/* Allocate contiguous physical */
	<V irtStart, PhyStart> ← AllocGpuPhyMem(S)
	Pairs ← ∅
	/* Set up p-chase from start to end */
	GpuSetP chase(V irtStart,V irtStart + S)
	CurV addr ← V irtStart
	PrevV addr ← CurV addr
	/* Collect all pairs */
	for Offset ← 1 to S do
		CurV addr ← Offset + V irtStart
		CurPaddr ← Offset + PhyStart
		TraverseP chase(V irtStart,CurV addr)
		/* Tests if cached data at VirtStart is evicted */
		if IsGpuCLEvicted(V irtStart) then
			Pair.append(PhyStart,CurPaddr)
			/* Remove CurVaddr from P-chase */
			∗PrevV addr ← CurV addr + 1
		else
			PrevV addr ← CurV addr
		end if
	end for
	F ← GenerateAllHashF unctions()
	/* Test all permutations of mapping function */
	for each func ∈ F do
		if IsV alidMapping(func,Pairs) then
			Print(func) return
		end if
	end for
end function

/* Code Executed by GPU */
/* Sets a pchase from BaseAddr to EndAddr */
function GpuSetPchase(BaseAddr, EndAddr)
	Addr ← BaseAddr
	/* Each address points to next address*/
	while Addr = EndAddr do
		*Addr ← Addr + 1
		Addr ← Addr + 1
	end while
	*Addr ← 0
end function
/* Traverse pchase from BaseAddr to EndAddr */
function TraversePchase(BaseAddr, EndAddr)
	Addr ← BaseAddr
	while Addr = 0 do
		Addr ← ∗Addr
	end while
end function
/* Check if data in BaseAddr is in cache */
function IsGpuCLEvicted(BaseAddr)
	/* Measure time to read Base address */
	StartTime ← GetClock()
	data ← BaseAddr
	AccessTime ← GetClock() − StartTime
	if AccessTime ≥ Threshold then
		return True
	end if
	return False
end function

Algorithm 2. Reverse-Engineering GPU Cache-line bits

 

D. Reverse-Engineering of L2 Cache set Addressing

 

물리적 주소를 캐시 세트에 매핑하는 해시함수를 역설계하는 우리의 알고리즘은 두 가지 사실에 의존한다:

1) 동일한 캐시 세트에 놓여 있는 캐시 라인은 서로를 쫓아낼 수 있고 2) 캐시된 워드에 액세스하는 것이 D램의 워드에 액세스하는 것보다 훨씬 빠르다. 알고리즘 2는 수도 코드를 개략적으로 설명한다. 핵심 아이디어는 알고리즘 1과 동일하다. 우리는 같은 L2 캐시 세트에 놓여 있는 것으로 알고 있는 충분한 쌍의 주소를 수집한 후, 유효한 해시 함수를 찾는다. 쌍을 찾기 위해 P-chase(pointer chase) [27], [29]를 사용한다. P-chase는 기본적으로 linked-list traversal이다. 핵심은 링크된 리스트의 요소들이 우리가 읽기를 원하는 주소에 배치된다는 것이다. P-chase list traversal(다음 요소의 주소는 현재 요소를 읽음으로써 발견됨)은 데이터에 의존하므로, 한 번에 하나의 읽기는 고장난 실행 코어에도 대기한다. 따라서 P-chase list은 Memory Barriers를 사용하지 않고 읽기 실행 순서에 대한 보증을 제공한다. (Nvidia GPU는 글로벌 메모리 배리어가 없다)

 GpuSetPchase() 함수는 연속된 주소에 요소가 놓여 있는 P-chase 목록을 구현한다. TraversePchase() 함수는 이 목록을 통과하는 데 사용된다. 12-25행에서, 우리는 P-chase list를 여러 번 트래버스하는데, 매번 끝 주소를 증가시키되 시작 주소는 고정된 상태로 유지한다. P-chase list를 통과할 때 시작 주소를 포함하여 모든 요소가 액세스될 때(LRU 캐시4로 가정) 캐시된다. 각 트래버설이 끝날 때, 우리는 시작 주소 VirtStart에서 데이터에 액세스하는 데 걸리는 시간을 측정한다. 두 가지 가능성이 있다:

1) 액세스 시간이 쓰레시홀드보다 작거나 같다. VirtStart는 제거되지 않았다. 우리는 다음 반복을 계속한다.
2) 쓰레시홀드보다 액세스 시간이 더 크다. 이는 VirtStart가 캐시에서 제거되었음을 나타낸다. 이전 트래버설이 VirtStart를 제거하지 않았고 이전과 현재 트래버설 사이에 액세스한 유일한 추가 요소가 현재 트래버설의 끝 주소이기 때문에 현재 트래버설의 시작 주소와 끝 주소가 동일한 캐시 세트에 있다고 결론짓는다. P-chase(라인 21)에서 이 엔드 주소를 제거하여 향후 P-chase 트래버스에서 VirtStart를 제거하지 않도록 한다.

 섹션 V-C에서와 같이, 우리는 여러 쌍을 수집한 후 유효한 매핑(line 26-32)을 찾는다. 그림 6(a)에는 GTX 1080에 대해 실험적으로 발견한 매핑 함수가 자세히 나와 있다. 비교를 위해 Figure 5는 CPU의 일반적인 캐시 세트 주소 지정 메커니즘을 보여준다. GPU L2 캐시에 대한 캐시 세트 어드레싱은 기존 CPU 캐시에 비해 훨씬 복잡하고 뱅크 어드레싱과 유사하다는 것을 알 수 있다.

 해시함수를 검증하기 위해 우리는 L2 캐시 hit 횟수와 miss 횟수를 나열할 수 있는 Nvidia가 제공하는 GPU 프로파일러인 nvprof[10]를 사용했다. 우리는 해시함수를 사용하여 같고 다른 캐시 세트에 있는 주소 세트를 생성한다. 우리의 실험에 따르면 nvprof는 동일한 캐시 세트의 주소가 액세스될 때 높은 캐시 누락률을 보고하고 다른 캐시 세트의 주소가 액세스될 때 낮은 캐시 누락률을 보고한다(일부 캐시 누락은 콜드 캐시 누락으로 인해 발생한다).

 

 

 

E. GPU Memory Hierarchy

 GTX 1070, GTX 1080, Tesla V100 GPU를 대상으로 실험(알고리즘 1, 2)을 실시했더니 해시함수가 그림 6(a) 및 (b)와 유사한 것으로 나타났다. 이러한 모든 GPU의 경우 뱅크 인덱스와 캐시 인덱스는 d LSB 비트를 공통으로 가지고 있으며, 여기서 d는 GPU 아키텍처에 의존한다. GTX 1070과 GTX 1080의 경우 d가 3이고 Tesla V100의 경우 d가 5이다. 이러한 측정에 기초하여 전체적인 GPU 아키텍처를 그림 6(e)와 같이 설계한다고 결론짓는다.

 

우리는 메모리 모듈을 GPU 메모리 계층의 파티션으로 정의하는데, GPU 메모리 계층의 파티션은 캐시 세트나 뱅크를 2개의 메모리 모듈이 공유하지 않도록 한다. 인덱스에 동일한 공통 d비트를 가진 모든 캐시 세트와 뱅크는 동일한 메모리 모듈에 놓여 있다. 따라서 총 메모리 모듈 수는 GTX 1070/GTX 1080에서는 2^3 = 8, 테슬라 V100에서는 2^5 = 32이다. 우리의 연구 결과는 Nvidia GPU GTX 970의 아키텍처를 보여주는 Figure 7과 일치한다. 이것은 우리의 결과를 더욱 입증한다. Figure 7에 근거해, 우리는 또한 각 메모리 모듈에 독립적인 메모리 컨트롤러와 D램 버스가 있다고 추측한다.

 GTX 1080에 대한 우리의 마이크로 벤치마크는 또한 우리가 부록 A에 기술한 L2 캐시(캐시 라인 크기, 연관성 설정, 캐시 교체 정책 등)와 D램의 다른 흥미로운 특성들을 밝혀냈다.

 

 

F. Page Coloring

 

페이지 컬러링의 숨은 아이디어는 두 개의 GPU 파티션 사이의 메모리 간섭을 제한하기 위해 서로 다른 GPU 파티션에 특정 물리적 페이지를 할당하는 것이다. 물리적 페이지는 두 개의 GPU 파티션이 동일한 캐시 세트에 액세스할 수 없도록 할당될 수 있으며 따라서 서로의 캐시 라인을 제거할 수 없다. 마찬가지로 GPU 파티션이 D램 뱅크를 공유하지 않도록 할당하여 행 버퍼 충돌을 피할 수 있다. 섹션 V-E에서 설명한 것처럼 Nvidia GPU의 경우, 다른 메모리 모듈에 있는 페이지를 다른 파티션에 할당하여 캐시 충돌과 DRAM 뱅크 충돌을 함께 피할 수도 있다.

 

어떤 접근법이 더 나은 isolation를 제공하는지를 평가하기 위해, 우리는 일련의 실험을 한다. 각 실험에서, 우리는 GTX 1080에서 각각 하나의 쓰레드를 갖는 n개의 CUDA 블록을 실행하는데, 여기서 n은 1에서 50까지 있다. 우리는 첫 번째 블록의 스레드를 Primary Thread라고 하고 다른 모든 스레드를 Secondary Threads라고 부른다. 모든 스레드는 루프에서 읽은 고유한 주소 집합에 할당된다. 그러나 단일 스레드에 할당된 모든 주소는 동일한 캐시 세트와 동일한 뱅크에 있다. Primary Thread가 액세스하는 주소는 일정하게 유지되며 Secondary Threads가 액세스하는 주소는 5가지 다른 경우에 따라 변경된다 :

1) primary thread에서 액세스한 주소와 동일한 캐시 세트 및 동일한 뱅크에 있는 주소. (SCSB)

 Addresses lying on the same cache set and the same bank as addresses accessed by the primary thread (SCSB).
2) 동일한 뱅크에 위치하지만 다른 캐시 세트에 놓여 있는 주소. (DCSB)

 Addresses lying on a different cache set but the same bank (DCSB).
3) 동일한 캐시 세트에 위치하지만 다른 뱅크에 위치하는 주소. (SCDB)

 Addresses lying on the same cache set but a different bank (SCDB).
4) 동일한 메모리 모듈에 위치하지만 다른 캐시 세트와 다른 뱅크에 놓여 있는 주소. (DCDB)

 Addresses lying on a different cache set and a different bank but the same memory module (DCDB).
5) 다른 메모리 모듈에 위치한 주소 (따라서 암시적으로 다른 캐시 세트와 다른 뱅크에 놓여 있는 주소, DCDB) (DM).

 Address lying on a different memory module (and hence, implicitly, on a different cache set and a different bank) (DM).

 

이러한 구성에서 우리는 primary thread가 단일 워드에 액세스하는 데 걸리는 평균 시간을 측정한다. Secondary Threads에 의한 간섭이 많아지면 primary thread에 걸리는 시간이 비례적으로 더 크게 증가할 것이다. 실험의 목적은 primary thread가 Secondary Threads로 인해 간섭으로부터 가장 격리된 경우를 찾는 것이다. 우리는 실험을 구성할 때 다음과 같은 예방 조치를 취한다 :

1) 섹션 IV에서 논의한 computation isolation implementation을 사용하여 메모리 간섭 이외의 간섭(즉, 컴퓨팅 간섭)을 방지하기 위해, Secondary Threads가 primary thread와 동일한 SM에서 실행되지 않도록 한다.
2) 뱅크 충돌을 확인하기 위해, 각 스레드는 모든 액세스가 D램에 도달하기 위해서 L2 캐시를 흘릴 수 있는 충분한 주소를 액세스한다. (i.e. accessing more cache lines than set associativity)

 

 

 Figure 8은 모든 5개의 시나리오에서 primary thread가 걸린 평균 시간을 보여준다. 예상대로 the primary thread and the secondary threads가 모두 동일한 캐시 세트 및 뱅크(SCSB)에 액세스할 때 최악의 간섭이 나타난다. 기본 스레드와 보조 스레드가 다른 캐시 세트 및 동일한 뱅크(DCSB)에 액세스하는 경우가 이에 가깝다. 이 두 경우 모두 뱅크 갈등에서 주로 간섭이 발생하고 있다. SCDB와 DCDB의 경우, L2 캐시의 MSHR(Miss Status Holding Register)을 둘러싼 충돌에서 간섭이 발생하는 것으로 본다. MSHR은 보류 중인 캐시 miss들을 추적하는 제한된 레지스터다. MSHR은 캐시 누락의 원인이 되는 각 request에 대해 할당되고 데이터를 D램에서 가져온 후에만 해제된다. 캐시 누락의 원인이 되는 추가 request은 MSHRs 중 하나가 해제될 때까지 차단된다. DCDB 곡선의 굴곡부에 기초하여, 단일 메모리 모듈 내의 L2 캐시 파티션에 대한 MSHR의 수는 16으로 나타난다. 따라서 the primary and secondary threads가 서로 다른 캐시 세트와 뱅크에 액세스하는 경우에도 동일한 메모리 모듈에 액세스하기 때문에 간섭이 발생하는 것을 볼 수 있다. primary thread가 간섭을 받지 않는 DM이 isolation에 관한 Best Case이다. 최악의 경우와 최선의 경우를 비교해 보면, 메모리 대역폭 파티셔닝은 잠재적으로 10배 이상의 인수로 간섭을 줄일 수 있다는 것을 알 수 있다.

 이러한 실험에서는 섹션 V-C에서 발견한 해시 함수를 사용하여 기본 및 보조 스레드에 주소를 할당했다. 우리는 그들이 같은 뱅크에 접근할 때 스레드들 사이에 날카로운 간섭이 있는 것을 본다. 따라서 우리는 알고리즘 1이 정확한 해시함수를 생성했다고 확신한다.

 

G. Implementing Page Coloring

 

섹션 V-F의 논의를 바탕으로 메모리 대역폭 분리를 달성하기 위해, 최상의 방법론은 다른 메모리 모듈의 주소를 다른 GPU 파티션에 할당하는 것이다. 또한 Figure 6(e)와 Figure 7을 보면 각 파티션에 자체 메모리 모듈을 할당하면 L2 캐시 세트와 뱅크의 분리 세트를 제외하고 각 파티션에 독립적인 메모리 컨트롤러와 D램 버스를 가질 수 있으며, 이는 높은 수준의 isolation을 제공할 것임을 알 수 있다. 8개의 메모리 모듈이 있는 GTX 1080의 예를 들면, 우리는 각각 고유한 메모리 모듈이나 컬러가 할당된 최대 8개의 GPU 파티션을 가질 수 있다. GPU 파티션에서 실행되는 모든 응용 프로그램에는 해당 메모리 모듈에서만 주소가 할당된다. 애플리케이션에 물리적 메모리를 할당하는 가장 세부적인 것은 페이지다. 따라서 페이지 컬러링의 경우, 한 페이지 내의 모든 주소가 동일한 메모리 모듈에 있어야 한다. Nvidia 장치 드라이버의 코드를 검사하여 최신 Nvidia GPU 아키텍처(Pascal [5] 및 Volta [9])가  다중 페이지 크기를 지원한다는 것을 알아냈다 : 4KB, 64KB, 2MB(모든 아키텍처는 4KB 페이지를 지원한다).

 

Figure 6(c)를 보면 GTX 1080의 경우 페이지 컬러링에 사용할 수 있는 페이지 크기는 4KB밖에 없다는 것을 알 수 있다. 4KB 페이지 내의 모든 주소는 M0-M3 또는 M4-M7에 위치하며, 여기서 Mi는 i번째 메모리 모듈이다. 메모리 모듈 인덱스(mbit2)의 세 번째 비트에는 12(및 2^12 = 4KB) 미만의 물리적 주소가 전혀 사용되지 않기 때문이다. 64KB 또는 2MB 크기의 페이지에는 모든 메모리 모듈에 주소가 배치된다. 그래서 GTX 1080의 경우 8개의 메모리 모듈을 가지고 있음에도 불구하고 2개의 메모리 색상만 가능하다(M0 - M3 및 M4 - M7). 마찬가지로 테슬라 V100의 경우 32개의 메모리 모듈에서 8개의 메모리 색상이 가능하다.

 

 

우리는 다음에 엔비디아 장치 드라이버에 페이지 컬러링 코드를 구현했다. 각 GPU는 해당 컬러의 모든 free page를 포함하는 메모리 컬러당 하나의 free-page list를 유지한다. GPU 초기화 시, 모든 페이지는 일치하는 free-page list에 배치된다. GPU 메모리를 할당하기 전에 애플리케이션은 사용할 컬러를 장치 드라이버에 표시한다. 해당 애플리케이션에 대한 GPU 메모리의 향후 모든 할당은 적절한 free-page list을 사용하여 이행된다.

 

EVALUATION

이번 섹션에서는 다양한 애플리케이션을 이용하여 compute and memory bandwidth isolation의 효과를 평가한다.

A. Micro-benchmark Experiments

 

Table 1에는 마이크로 벤치마크 연구에 사용하는 애플리케이션이 나열되어 있다. 이 세트에는 CUDA SDK[3]와 Rodinia benchmark suite[12]에서 가져온 계산 및 메모리 집약적 애플리케이션이 혼합되어 있다. 우리는 GTX 1080과 Tesla V100이라는 두 개의 Nvidia GPU를 대상으로 평가를 실시한다. Table 2에는 시험한 플랫폼의 사양이 수록되어 있다. 우리는 다음과 같은 접근방식을 사용하여 각 GPU를 두 개의 동일한 파티션인 P1과 P2로 분할한다 :

1) Compute Partitioning only(CP) - 각 파티션에는 SM의 분리 세트가 할당된다.
2) Both Compute and Memory Partitioning(CMP) - 각 파티션에는 다른 메모리 컬러와 SM의 분리 세트가 할당된다.
3) Nvidia MPS QoS feature5 (MPS) - 각 파티션은 GPU의 전체 쓰레드의 50%까지만 사용할 수 있다. Volta 아키텍처를 가진 GPU만이 이 QoS 기능을 가지고 있기 때문에 GTX 1080 평가에서 이러한 접근방식을 생략한다.

 

각 애플리케이션 A에 대해 다음과 같은 여러 시나리오를 실행한다 :
1) P1에서 실행되는 A. P2 is idle

2) P1에서 실행되는 A와 P2에서 실행되는 kernels of compute-intensive Matrix Multiplication (MM)
3) P1에서 실행되는 A와 P2에서 실행되는 kernels of memory-intensive Fast Walsh Transform (FWT)

4) P1에서 실행되는 A와 P2에서 실행되는 kernels of memory-intensive Vector Addition (VA)

 

  모든 다른 파티셔닝 접근에 대해 위 시나리오에서 시나리오의  A의 커널 평균 런타임을 측정한다. 완벽한 파티셔닝 기법으로 모든 시나리오의 실행 시간은 같아야 한다(즉, corunn인 커널의 간섭은 없을 것이다). 런타임의 변동량은 간섭의 양을 정의한다. 따라서 우리는 격리를 측정하기 위한 다음 측정 기준을 정의한다.

 여기서 T_A,B는 다른 파티션에서 실행되는 B의 커널이 A의 커널의 평균 런타임을 의미한다(그리고 T_A,-는 P2가 idle 상태임을 나타낸다). variation이 작을수록 더욱 예측 가능한 런타임인 것을 의미한다.

 우리는 기준선으로서 A의 커널이 파티셔닝 없이 GPU에서 단독으로 실행될 때 그 런타임을 측정한다. 우리는 (GPU의 50% 파티션에서 A의 런타임은 2배 증가할 것으로 예상되므로) 2 ∗T_baseline 인수로 모든 런타임을 노멀라이즈한다. Figure 9(a)는 GTX 1080의 결과를, 9(b)는 테슬라 V100의 결과를 나타낸다. Table 3는 Table 1에 나열된 모든 애플리케이션의 평균 및 최대 변동을 나타낸다.

 

Key observations from the evaluation on GTX 1080:

1) 예측 가능 성능에 있어 CP보다 CMP가 훨씬 좋다. Computation isolation은 예측 가능 성능에 충분치 않다. memory isolation이 없으면 memory-intensive 애플리케이션이 높은 간섭을 일으킬 수 있다. 이것이 vector addition(VA)이 memoryintensive가 높아 가장 큰 저하를 유발하는 이유이다. memory bandwidth isolation를 추가하면 평균 130.4%에서 7.5%로 변동이 줄어든다.
2) 예측가능성과 성능 사이에 tradeoff가 있다. CP_<None>의 경우 CMP_<None> 에 비해 모든 애플리케이션의 실행 시간이 짧다. 이 효과는 중요하며, memory bandwidth partitioning이 GPU 파티션 간에 bandwidth를 분할하여 다른 파티션이 idle 상태일 때에도 단일 파티션에서 사용할 수 있는 메모리 대역폭을 감소시키기 때문이다. 이러한 영향은 당연히 사용 가능한 memory bandwidth에 민감한 메모리 집약적 애플리케이션에 더 큰 영향을 미친다.

3) GPU 자원의 활용도가 낮은 애플리케이션도 있다. 표준화된 실행 시간이 1 미만인 모든 애플리케이션은 GPU를 충분히 활용하지 못하고 있음을 나타낸다.예를 들어 vector addition에 대한 CP_<None>의 경우, 기준선에 비해 컴퓨팅 리소스의 절반을 얻고 있음에도 표준화 런타임은 0.5이다. 이는 전체 GPU에서 실행될 때와 단일 파티션에서 실행될 때의 런타임은 동일하다는 것을 의미한다(런타임은 2 ∗ T_baseline으로 표준화된다). 그 이유는 VA가 메모리 집약적인 애플리케이션이기 때문에 대부분의 시간을 읽기/쓰기에 정체되어 있기 때문이다. 따라서 컴퓨팅 리소스의 감소는 런타임의 증가로 이어지지 않는다.

 

Tesla V100의 결과는 GTX 1080의 결과와 유사하며 다음과 같은 몇 가지 차이점이 있다 :

1) 예측 가능 성능을 대해 CP가 엔비디아 MPS보다 좋다. MPS의 평균 변동률은 135.8%인 반면 CP의 평균 변동률은 48.1%로 상당히 낮다. 이는 스레드 세분화에서는 MPS 파티션이, SM 세분화에서는 CP 파티션이 분할되기 때문이다. 따라서 MPS 파티셔닝에서 두 개의 애플리케이션이 동일한 SM에서 실행되어 SM 리소스를 공유하게 되어 충돌이 발생할 수 있다.
2) GTX 1080보다 Tesla V100에서 CP가 더 좋은 성능을 발휘한다. 테슬라 V100의 메모리 대역폭이 더 높기 때문이다. DRAM bandwidth는 3배, L2 캐시는 3배, D램 뱅크는 4배 더 많다. 이러한 추가 자원은 L2 캐시/DRAM에서 충돌 가능성을 감소시킨다. 그러나 CMP는 평균 및 최대 변동에 있어 여전히 더 좋다. 테슬라 V100은 현재 가장 큰 [8]의 엔비디아 GPU이기 때문에, 다른 GPU(GTX 1080의 경우 등)의 경우 CMP와 CP/MP의 격차가 더 클 것으로 본다.

3) CP와 CMP는 SN과 CFD에 대해 높은 부하를 갖는다. 이는 SN 및 CFD에 대한 MPS_<None>와 CP_<None>을 비교할 때 명백하다. compute isolation의 구현은 커널당 고정 오버헤드를 가지며, 이는 작은 커널의 경우 상대적으로 커진다. Tesla V100은 GTX 1080보다 더 강력하기 때문에 모든 커널 실행 시간이 단축된다. SN과 CFD 모두 여러 개의 작은 커널을 가지고 있다. 따라서 SN 및 CFD의 경우 컴퓨팅 파티셔닝에 대한 Tesla V100의 오버헤드가 상대적으로 크다. 다른 애플리케이션의 경우 MPS_<None>과 CP_<None>이 비교 가능하므로, 컴퓨팅 파티셔닝에 대한 오버헤드가 상대적으로 작음을 나타낸다.

 테슬라 V100은 메모리 색상이 최대 8개까지 가능하기 때문에 메모리 계층을 2, 4, 8개 파티션으로 나눌 수 있다는 것을 검증했다. Appendix B는 2가 아닌 4개의 메모리 파티션을 사용하여 Tesla V100의 마이크로벤치마크 결과를 제시한다. 결과는 그림 9(b)와 유사하다.

 

 

B. Caffe Framework Experiments

 

 우리는 FGPU API를 사용하기 위해 Caffe[16]을 포팅했다. Caffe는 Closed 소스인 BLAS(기본 선형 대수 서브루틴)의 CUDA 구현인 cuBLAS 라이브러리[2]를 사용한다. 우리의 현재 FGPU 구현은 커널을 약간 수정해야 하기 때문에, 우리는 cuBLAS 루틴을 기본 CUDA 커널로 대체했다. Caffe의 포팅은 수많은 기존 애플리케이션이 수정 없이 FGPU의 혜택을 누릴 수 있도록 해준다.

 Caffe 성능을 평가하기 위해, 우리는 ImageNet[18]에서 훈련된 AlexNet[14] 모델을 사용하는 Caffe의 이미지 분류(IC) 애플리케이션을 실행했다. 우리는 첫 번째 파티션에 IC가 실행되는 섹션 VI-A에서와 같은 실험을 했다. 또한 IC도 두 번째 파티션에 실행되어 간섭 작업으로 작용하는 또 다른 시나리오를 추가한다. 커널 실행 시간만 측정하는 것이 아니라 IC가 취한 총 시간을 측정한다. IC는 GPU 집약적이며 CPU는 주로 CPU와 GPU D램 간의 데이터 전송에 사용된다. Figure 10은 GTX 1080의 결과를 보여준다. 컴퓨팅 파티셔닝의 variation_IC 경우 104.6%인 반면 메모리 파티셔닝을 추가하면 19.5%로 줄어든다. FGPU의 컴퓨팅과 메모리 파티셔닝은 그들의 부재에 비해 극적으로 변동을 감소시키지만, 여전히 이해할 수 없는 간섭이 존재한다. 안전에 치명적인 시스템은 이러한 제한을 인식하고 적절한 안전 여유도를 사용해야 한다.

 

C. Related work

 

 

 이전 연구[17] [21] [22] [32] [33]는 여러 아키텍처에 걸쳐 다양한 Nvidia GPU의 메모리 계층 구조를 해부하려고 시도했다. [32], [33]의 작업에서는 CPU와 유사한 L2 캐시 구조를 가진 초기 단순 Nvidia GPU(2010년 이전)를 평가했다. [21]에서 저자들은 새로운 아키텍처에서 L2 캐시의 구조를 식별할 수 없었으며, 평가한 GPU 중 전통적인 메모리 주소를 사용한 GPU는 없다고 기술하고 있다. [17]에서는 GPU의 L2 캐시 집합과 볼타 아키텍처의 연관성을 찾기 위해 [21]에 소개된 기법을 사용하였으나, 캐시 구조도 찾지 못하였다. [22]의 저자들은 C2070 GPU를 마이크로 벤치마킹하려고 시도했으나 잡음이 많은 데이터를 얻었으며 C2070의 L2 캐시가 평균 13.7의 연관성을 가지고 있다고 진술했다. 우리가 아는 한, GPU에서 D램의 구조를 해부하려는 사전 작업은 없었다. 이 작품에서는 새로운 Nvidia GPU에 초점을 맞추고 L2 캐시와 D램의 구조를 모두 역설계하여 페이지 컬러링을 구현했다.

 페이지 컬러링은 CPU에 대해 잘 알려진 기법이다. [20]은 Real-Time System CPU에 캐시 컬러링을 구현했고 [19]는 D램 뱅크 컬러링을 구현했다. [30]은 멀티 코어 CPU 시스템의 캐시 색상과 뱅크 색상을 모두 결합했다. 페이지 컬러링은 CPU에 대한 D램 버스 경합 문제를 해결하지 못하므로, 다른 연구 [13] [26] [35]는 그 영향을 분석하거나 제한하려고 시도했다. [31] 캐시 컬러링이 공유 캐시의 MSHR을 둘러싼 갈등 문제를 해결하지 못함을 보여주었다. 그들의 작업에서 시뮬레이터에서 사용자 지정 하드웨어 확장을 구현하는 데 필요한 CPU의 MSR 충돌을 제거했다. FGPU는 캐시 세트, D램 뱅크, D램 버스, 메모리 컨트롤러, MSHR 레지스터 등 GPU의 공유 메모리 자원을 하드웨어 수정 없이 분할하기 위해 페이지 컬러링 기법을 사용한다. 우리는 GPU 메모리 계층 구조가 페이지 컬러링을 공동 실행 작업 간의 간섭을 줄이기 위한 간단하고 효과적인 솔루션으로 만든다는 것을 보여준다.

 

 GPU 컴퓨팅 리소스의 파티셔닝과 관련하여, 우리의 구현은 [34]를 기반으로 한다. 이들의 작업은 SM 중심의 커널 변환을 통해 GPU에 대한 스케줄을 제어하고 이 제어장치를 활용해 시스템 처리량을 높이고 평균 턴어라운드 시간을 줄이는 데 초점을 맞추고 있다. [15]은 SM 중심의 커널 변환을 사용하여 Chollesky Factorization의 작업 처리량을 개선한다. [28]의 저자들은 SM을 SM 중심 변환을 사용한 조정 작업 사이에 분할하여 시스템에서 스케줄링할 수 있는 총 작업 수를 늘리고 스케줄링 가능성 분석을 검토하기 위한 경험담을 생각해 냈다. 이 작업에서는 GPU에서 다중 애플리케이션을 실행하기 위한 알고리즘 스케줄링에 초점을 맞추지 않고, 서로 다른 파티션에서 실행되는 애플리케이션을 가장 잘 분리하는 방법에 초점을 맞춘다. 컴퓨팅 파티셔닝과는 별도로 이를 위해 메모리 대역폭 격리가 필요하다는 것을 보여준다. [25]의 저자들은 GPU에서 공동 실행 워크로드 간의 리소스 경합도 관찰했지만 경합을 유발하는 리소스는 설명하지 않았다.

 

CONCLUSIONS

 

본 논문에서 이러한 공동 실행 작업 중 컴퓨팅 및 메모리 리소스를 분할하여 격리 상태를 유지하면서 여러 애플리케이션을 GPU에서 병렬로 실행할 수 있는 소프트웨어 기반 메커니즘을 제시한다. 하나의 대형 GPU를 더 작은 부분 GPU로 분할하면 실시간 시스템 설계자에게 GPU 리소스 스케줄링과 관련된 더 많은 옵션이 제공된다. 우리의 평가는 FGPU를 사용하는 애플리케이션이 병렬로 실행되는 다른 애플리케이션과 관계없이 훨씬 더 예측 가능한 실행 시간을 가지고 있다는 것을 보여준다. 우리는 또한 이전에 공공 문헌에서 알려지지 않았던 엔비디아 GPU에 대한 다양한 세부사항들을 제시한다. 우리는 이러한 통찰력이 프로그래머들이 GPU 아키텍처를 더 잘 이해하고 프로그램의 성능을 최적화하는 데 도움이 되기를 바란다.

 

 

REFERENCES

[1] Drive PX-series. https://en.wikipedia.org/wiki/Drive PX-series.
[2] NVIDIA cuBLAS. https://developer.nvidia.com/cublas.
[3] NVIDIA CUDA Toolkit. http://www.nvidia.com/object/embedded-systems-dev-kits-modules.html.
[4] NVIDIA Discloses Full Memory Structure and Limitations of GTX 970. https://www.pcper.com/reviews/Graphics-Cards/NVIDIA-Discloses-Full-Memory-Structure-and-Limitations-GTX-970.
[5] NVIDIA GP100 Pascal Whitepaper. http://www.nvidia.com.
[6] NVIDIA GTX 1080. https://international.download.nvidia.com/
geforce-com/international/pdfs/GeForce GTX 1080 Whitepaper FINAL.pdf.
[7] NVIDIA MPS. https://docs.nvidia.com/deploy/pdf/CUDA Multi Process Service Overview.pdf.
[8] NVIDIA Tesla V100 Tensor Core. https://www.nvidia.com/en-us/data-center/tesla-v100/.
[9] NVIDIA V100 Volta Whitepaper. http://www.nvidia.com.
[10] NVIDIA Visual Profiler User’s Guide. http://docs.nvidia.com/cuda/profiler-users-guide/index.html.
[11] Optimizing Matrix Transpose in CUDA. http://developer.download.
nvidia.com/compute/DevZone/C/html x64/6 Advanced/transpose/doc/MatrixTranspose.pdf.
[12] S. Che, M. Boyer, J. Meng, D. Tarjan, J. W. Sheaffer, S.-H. Lee, and K. Skadron. Rodinia: A benchmark suite for heterogeneous computing. In Workload Characterization, 2009. IISWC 2009. IEEE International Symposium on, pages 44–54. Ieee, 2009.
[13] D. Dasari, B. Andersson, V. Nelis, S. M. Petters, A. Easwaran, and J. Lee. Response time analysis of cots-based multicores considering the contention on the shared memory bus. In 8th IEEE International Conference on Embedded Software and Systems, pages 1068–1075. IEEE, 2011.
[14] J. Deng, W. Dong, R. Socher, L.-J. Li, K. Li, and L. Fei-Fei. Imagenet: A large-scale hierarchical image database. In Computer Vision and Pattern Recognition, 2009. CVPR 2009. IEEE Conference on, pages 248–255. Ieee, 2009.
[15] J. Janz´en, D. Black-Schaffer, and A. Hugo. Partitioning gpus for improved scalability. In Computer Architecture and High Performance Computing (SBAC-PAD), 2016 28th International Symposium on, pages 42–49. IEEE, 2016.
[16] Y. Jia, E. Shelhamer, J. Donahue, S. Karayev, J. Long, R. Girshick, S. Guadarrama, and T. Darrell. Caffe: Convolutional architecture for fast feature embedding. In Proceedings of the 22nd ACM international conference on Multimedia, pages 675–678. ACM, 2014.
[17] Z. Jia, M. Maggioni, B. Staiger, and D. P. Scarpazza. Dissecting the nvidia volta gpu architecture via microbenchmarking. arXiv preprint arXiv:1804.06826, 2018.
[18] A. Krizhevsky, I. Sutskever, and G. E. Hinton. Imagenet classification with deep convolutional neural networks. In Advances in neural information processing systems, pages 1097–1105, 2012.
[19] L. Liu, Z. Cui, M. Xing, Y. Bao, M. Chen, and C. Wu. A software memory partition approach for eliminating bank-level interference in multicore systems. In Proceedings of the 21st international conference on Parallel architectures and compilation techniques, pages 367–376. ACM, 2012.
[20] R. Mancuso, R. Dudko, E. Betti, M. Cesati, M. Caccamo, and R. Pellizzoni. Real-time cache management framework for multi-core architectures. In Real-Time and Embedded Technology and Applications Symposium (RTAS), 2013 IEEE 19th, pages 45–54. IEEE, 2013.
[21] X. Mei and X. Chu. Dissecting gpu memory hierarchy through microbenchmarking. IEEE Transactions on Parallel and Distributed Systems, 28(1):72–86, 2017.
[22] R. Meltzer, C. Zeng, and C. Cecka. Micro-benchmarking the c2070. In GPU Technology Conference. Citeseer, 2013.
[23] J. Nickolls, I. Buck, M. Garland, and K. Skadron. Scalable parallel programming with CUDA. ACM Queue, 6(2):40–53, 2008.
[24] C. Nvidia. Nvidia cuda c programming guide. Nvidia Corporation, 120(18):8, 2011.
[25] N. Otterness, M. Yang, S. Rust, E. Park, J. H. Anderson, F. D. Smith, A. Berg, and S. Wang. An evaluation of the nvidia tx1 for supporting real-time computer-vision workloads. In Real-Time and Embedded Technology and Applications Symposium (RTAS), 2017 IEEE, pages 353–364. IEEE, 2017.
[26] R. Pellizzoni, A. Schranzhofer, J.-J. Chen, M. Caccamo, and L. Thiele. Worst case delay analysis for memory interference in multicore systems. In Design, Automation & Test in Europe Conference & Exhibition (DATE), 2010, pages 741–746. IEEE, 2010.

[27] R. H. Saavedra-Barrera. CPU performance evaluation and execution time prediction using narrow spectrum benchmarking. PhD thesis, University of California, Berkeley, 1992.
[28] S. K. Saha. Spatio-Temporal GPU Management for Real-Time Cyber- Physical Systems. PhD thesis, UC Riverside, 2018.
[29] A. J. Smith and R. H. Saavedra. Measuring cache and tlb performance and their effect on benchmark runtimes. IEEE Transactions on Computers, (10):1223–1235, 1995.
[30] N. Suzuki, H. Kim, D. De Niz, B. Andersson, L. Wrage, M. Klein, and R. Rajkumar. Coordinated bank and cache coloring for temporal protection of memory accesses. In Computational Science and Engineering (CSE), 2013 IEEE 16th International Conference on, pages 685–692. IEEE, 2013.
[31] P. K. Valsan, H. Yun, and F. Farshchi. Taming non-blocking caches to improve isolation in multicore real-time systems. In Real-Time and Embedded Technology and Applications Symposium (RTAS), 2016 IEEE, pages 1–12. IEEE, 2016.
[32] V. Volkov and J. W. Demmel. Benchmarking gpus to tune dense linear algebra. In High Performance Computing, Networking, Storage and Analysis, 2008. SC 2008. International Conference for, pages 1–11. IEEE, 2008.
[33] H. Wong, M.-M. Papadopoulou, M. Sadooghi-Alvandi, and A. Moshovos. Demystifying gpu microarchitecture through microbenchmarking. In Performance Analysis of Systems & Software (ISPASS), 2010 IEEE International Symposium on, pages 235–246. IEEE, 2010.
[34] B. Wu, G. Chen, D. Li, X. Shen, and J. Vetter. Enabling and exploiting flexible task assignment on gpu through sm-centric program transformations. In Proceedings of the 29th ACM on International Conference on Supercomputing, pages 119–130. ACM, 2015.
[35] H. Yun, G. Yao, R. Pellizzoni, M. Caccamo, and L. Sha. Memory access control in multiprocessor for real-time systems with mixed criticality. In Real-Time Systems (ECRTS), 2012 24th Euromicro Conference on, pages 299–308. IEEE, 2012.