Cuda에서의 워프 스케줄링과 워프 컨텍스트 스위칭 간의 관계

Question

내가 이해하는 한, 준비 워프는 워프 스케줄링에서 실행할 수있는 워프입니다. 대기중인 워프는 소스 피연산자를 가져 오거나 계산할 수 없도록 실행되기 위해 기다리고 있습니다. 워프 스케쥴러는 "워프 스케쥴링"을 위해 실행할 워프를 선택합니다.

반면에 워프가 파이프 라인 스톨 또는 긴 전역 메모리 대기 시간을 갖는 경우 대기 시간을 숨기기 위해 다른 워프가 실행됩니다. 이것은 cuda에서 "warp context switching"의 기본 개념입니다.

제 질문은 : Cuda에서의 워프 스케줄링과 워프 컨텍스트 전환 간의 관계는 무엇입니까? 내 질문을 정교하게하기 위해 아래 예가있다.

예. 워프 A가 정지되고 워프 A가 전역 메모리를 가져올 대기중인 워프이면, 요소가 페치되면 워프 A가 준비되거나 워프 풀로 전환됩니다. 이를 기반으로 워프 컨텍스트 스위칭은 워프 스케줄링의 일부입니다. 맞습니까?

누구나 쿠다의 워프 컨텍스트 스위칭 및 워프 스케쥴링에 대한 참조를 제공 할 수 있습니까? Nvidia는이 문서를 공개적으로 사용할 수 없게 만듭니다.

미리 답변 해 주셔서 감사합니다.

Answer 1

준비 준비가 된 준비는 다음주기에 계획 할 수 있습니다. 실속 된 날실은 예약 할 수 없습니다.

매우 간단한 예제를 사용하여 대기 시간에 대한 질문에 대답하려면 주기억 장치에 대한 대기 시간이 8 실행주기라고 가정하고 컴퓨터가 파이프 라인이라는 사실을 무시합시다. 데이터가 준비되면 모든 명령을 한 주기로 실행할 수 있다고 가정합니다. 코드가 완료 될 때,

int idx = threadIdx.x+blockDim.x*blockIdx.x; 

int myval = global_data[idx]*global_data[idx]; 

이 myval는 글로벌 메모리에있는 항목의 제곱이 포함되어 있어야합니다

지금 나는이 같은 C 코드가 있다고 가정합니다. 이것은 일련의 어셈블리 언어 명령으로 분해됩니다. 다음과 같이 가정 해 봅시다 :

I0: R0 = global_data[idx]; 
I1: R1 = R0 * R0; 
I2: ... 

모든 스레드가 첫 번째 코드 행을 실행할 수 있습니다 (처음에는 스톨 없음). 아직 의존성이 없으며 독서 자체가 실속을 일으키지 않습니다. 그러나 모든 스레드는 두 번째 코드 행으로 이동할 수 있으며 이제는 R0의 값이 정확해야하므로 읽기가 검색 될 때까지 대기가 발생합니다. 앞서 언급했듯이 대기 시간이 8 사이클이고 32의 워프와 512의 스레드 블럭 크기를 사용하면 총 16 개의 워프가 있습니다. 단순함을 위해서 우리는 오직 32 단위의 실행 단위를 가진 Fermi SM을 가지고 있다고 가정 해 봅시다. 순서는 다음과 같이 보일 것이다 :

cycle:  ready warps: executing warp:  instruction executed:  Latency: 
    0   1-16     0   I0 -> I1 (stall) -- 
    1   2-16     1   I0 -> I1 (stall)  | -- 
    2   3-16     2   I0 -> I1 (stall)  | | 
    3   4-16     3   I0 -> I1 (stall)  | | 
    4   5-16     4   I0 -> I1 (stall)  | | 
    5   6-16     5   I0 -> I1 (stall)  | | 
    6   7-16     6   I0 -> I1 (stall)  | | 
    7   8-16     7   I0 -> I1 (stall)  | | 
    8   0,9-16     8   I0 -> I1 (stall) <- | 
    9   1,9-16     0   I1 -> I2   <---- 

우리가 참조하는 대기 시간이 다른 날실의 명령어를 실행하여 성취 한 후, 이전에 "정체"워프가 준비 워프 풀을 다시 입력하는 것입니다, 그리고 그것은 가능 (즉, I1에 포함 된 곱셈 연산을 수행하기 위해) 스톨 조건이 제거 된 바로 다음주기에서 스케쥴러가 그 워프를 다시 스케쥴하기 위해서.

대기 시간 숨기기와 워프 일정 사이에는 모순이 없습니다. 이들은 충분한 작업이 필요한 코드를 위해 함께 작업하여 전역 메모리에서 읽는 것과 같은 다양한 작업과 관련된 대기 시간을 숨 깁니다.

위 예제는 실제 동작과 비교하여 단순화되었지만 "충분한 작업"이있는 상태에서 워프 일정을 통해 대기 시간을 숨길 수있는 방법을 보여주기 위해 대기 시간 숨기기 및 워프 일정 계획의 개념을 적절하게 나타냅니다.