Same Workload Baseline 에 이은 후속 실험 설계 : Scheduler Policy Probing

배경

동일 Task Baseline 실험에서는 모든 Warp 가 동일한 independent ALu workload 를 수행하도록 구성

이 실험은 장기 실행 상황에서 Warp 간 progress 편향이 발생하는지 확인하는 데 유효하지만, 모든 Warp 가 같은 조건에 놓이기 때문에 스케줄러의 내부 선택 기준을 명확히 분리하기 어렵다.

즉 관찰된 결과가 다음 중 무엇 때문인지 구분하기 어렵다. 

• Warp Id 기반 고정 우선순위
• Round-Robin 계열 순환 발행
• Ready 상태 Warp 에 대한 Greedy 선택
• 파이프라인 상태나 scoreboard 상태에 따른 동적 선택
• 측정 타이밍 또는 retire 순서에 따른 외향적 편향

따라서 다음 단계에서는 Warp 간 조건을 의도적으로 비대칭화하거나, 특정 자원에 대한 경합을 강제로 유도하여 스케줄러의 선택 법칙을 더 명확히 드러내는 실험이 필요하다.

 

핵심 방향

후속 실험의 목적은 단순히 Warp 별 완료 순서를 보는 것이 아니라, 다음 질문에 답하는 것이다.

여러 Warp 가 동시에 또는 거의 동시에 실행 가능한 상태가 되었을 때, GPU 스케줄러는 어떤 Warp 를 먼저 issue 하는가?

이를 위해 다음 세 가지 조건을 만드낟.

1. Warp 별 작업량 차이
2. Barrier 이후 동시 Ready 상태
3. Memory 또는 Shared Memory 자원 경합

 

1. Mixed Workload Probe

목적

Warp 별로 서로 다른 작업랴응ㄹ 부여했을 때, 가벼운 Warp 가 무거운 Warp 를 추우러할 수 있는지 확인한다. 

동일 Task Baseline 에서는 모든 Warp 가 같은 작업을 수행했기 때문에, 완료 순서만으로는 스커줄러가 고정 우선순위인지 동적 선택인지 판단하기 어렵다.

Warp 마다 다른 부하를 부여하며, 스케줄러가 실제 실행 가능 상태와 workload 차이에 반응하는지 확인하기 위한 실험이다.

 

실험 설계

Warp ID 에 따라ㅏ 서로 다른 연산량을 부여한다.

.

관찰 항목

• Warp별 시작 cycle
• Warp별 종료 cycle
• Warp별 실행 duration
• 완료 순서
• Warp ID와 완료 순서의 상관관계
• workload 크기와 완료 순서의 상관관계

기대되는 해석

Case A: 가벼운 Warp가 먼저 종료됨

작업량이 작은 Warp가 무거운 Warp를 추월하여 먼저 끝난다면, 스케줄러가 고정 ID 순서만 따르는 것이 아니라 Ready 상태와 실행 가능성을 기반으로 동적으로 issue하고 있을 가능성이 높다.

Case B: 작업량과 무관하게 Warp ID 순서가 유지됨

작업량 차이가 있음에도 항상 낮은 Warp ID 또는 특정 Warp ID 순서대로 진행된다면, 고정 우선순위 또는 deterministic issue policy가 개입되어 있을 가능성이 있다.

Case C: 실행마다 순서가 미세하게 달라짐

동일한 커널을 반복 실행했을 때 완료 순서가 흔들린다면, 스케줄러 선택에 비결정적 요소 또는 다른 하드웨어 상태 의존성이 포함되어 있을 가능성이 있다.

---

2. Barrier Release Probe

목적

모든 Warp를 동기화 지점에 모은 뒤, Barrier 해제 직후 어떤 Warp가 가장 먼저 실행을 재개하는지 확인한다.

Mixed Workload가 작업량 차이를 이용한 실험이라면, Barrier Release Probe는 모든 Warp가 동시에 Ready 상태가 되었을 때의 순수 issue 우선순위를 관찰하는 실험이다.

실험 설계

모든 Warp가 __syncthreads()에 도달하도록 한 뒤, Barrier 직후 각 Warp의 lane 0이 timestamp를 기록한다.

int warp_id = threadIdx.x / 32;
int lane_id = threadIdx.x & 31;

__syncthreads();

uint64_t t = clock64();

if (lane_id == 0) {
    out[warp_id] = t;
}

보다 명확한 차이를 보기 위해 Barrier 이후에 짧은 ALU loop를 추가할 수 있다.

__syncthreads();

uint64_t start = clock64();

for (int i = 0; i < FIXED_ITER; i++) {
    acc = acc * 1664525 + 1013904223;
}

uint64_t end = clock64();

if (lane_id == 0) {
    out[warp_id].start = start;
    out[warp_id].end = end;
}

관찰 항목

• Barrier 직후 start timestamp
• Warp ID별 start 순서
• 반복 실행 간 start 순서 안정성
• 같은 SM 내 block 단위 결과
• block 반복 수 증가 시 패턴 유지 여부

기대되는 해석

Case A: 항상 동일한 Warp ID가 먼저 시작

Barrier 해제 직후에도 특정 Warp ID가 항상 먼저 실행된다면, Ready queue 또는 warp scheduler 내부에 deterministic priority가 있을 가능성이 있다.

Case B: Round-Robin 형태의 순환 패턴

실행마다 시작 Warp가 순환한다면, scheduler가 마지막 issue 위치를 기억하거나 round-robin pointer를 유지할 가능성이 있다.

Case C: 실행마다 불규칙

Barrier 이후 모든 Warp가 동시에 Ready 상태가 되었는데도 순서가 흔들린다면, issue selection이 완전히 고정된 구조가 아니라 scoreboard, dispatch 상태, scheduler partition 등의 영향을 받을 가능성이 있다.

---

3. Resource Contention Probe

목적

특정 하드웨어 자원에 경합을 유도하여, Warp가 stall 상태에 빠졌다가 해제될 때 스케줄러가 어떤 Warp를 먼저 재개하는지 확인한다.

Baseline과 Mixed Workload는 주로 ALU 중심의 실험이다. 하지만 실제 GPU 스케줄링은 memory dependency, scoreboard, shared memory bank conflict, register pressure 등의 영향을 크게 받는다.

따라서 자원 경합 상황을 만들어야 스케줄러의 arbitration 성향을 더 강하게 드러낼 수 있다.

---

3.1 Shared Memory Bank Conflict

실험 설계

일부 Warp는 shared memory bank conflict를 강하게 일으키고, 다른 Warp는 conflict가 적은 접근 패턴을 사용한다.

extern __shared__ int smem[];

int warp_id = threadIdx.x / 32;
int lane_id = threadIdx.x & 31;

uint64_t start = clock64();

int idx;

if (warp_id % 2 == 0) {
    // conflict 유도
    idx = 0;
} else {
    // conflict 완화
    idx = lane_id;
}

for (int i = 0; i < ITER; i++) {
    acc += smem[idx];
}

uint64_t end = clock64();

if (lane_id == 0) {
    out[warp_id].start = start;
    out[warp_id].end = end;
}

관찰 항목

• Conflict Warp와 Non-conflict Warp의 duration 차이
• Non-conflict Warp가 먼저 retire하는지 여부
• Conflict 해제 후 Warp ID별 재진입 순서
• 반복 실행 간 재진입 순서 안정성

---

3.2 Memory Latency / Global Load Stall

실험 설계

특정 Warp만 global memory dependent load chain을 수행하게 하여 scoreboard stall을 유도한다.

int warp_id = threadIdx.x / 32;
int lane_id = threadIdx.x & 31;

uint64_t start = clock64();

if (warp_id % 2 == 0) {
    for (int i = 0; i < ITER; i++) {
        idx = global_buf[idx];
    }
} else {
    for (int i = 0; i < ITER; i++) {
        acc = acc * 1664525 + 1013904223;
    }
}

uint64_t end = clock64();

if (lane_id == 0) {
    out[warp_id].start = start;
    out[warp_id].end = end;
}

관찰 항목

• Memory-bound Warp의 stall duration
• ALU-bound Warp가 memory-bound Warp를 추월하는지 여부
• Memory dependency가 풀린 Warp가 즉시 재개되는지 여부
• Warp ID보다 dependency 해제 시점이 더 중요한지 여부

---

4. Divergence / Active Mask Probe

목적

Warp 내부에서 lane들이 서로 다른 branch path를 수행할 때, 하드웨어가 어떤 path를 먼저 실행하는지 확인한다.

이 실험은 Warp 간 scheduling이 아니라, Warp 내부 divergence 처리 순서와 reconvergence 동작을 보기 위한 보조 실험이다.

실험 설계

Warp 내부 lane을 둘로 나누어 서로 다른 branch를 수행하게 한다.

int lane_id = threadIdx.x & 31;

uint64_t t0 = clock64();

if (lane_id < 16) {
    for (int i = 0; i < HEAVY_ITER; i++) {
        acc = acc * 1664525 + 1013904223;
    }
    path_out[lane_id] = clock64();
} else {
    for (int i = 0; i < LIGHT_ITER; i++) {
        acc = acc * 1664525 + 1013904223;
    }
    path_out[lane_id] = clock64();
}

uint64_t t1 = clock64();

관찰 항목

• if path와 else path 중 어느 쪽 timestamp가 먼저 찍히는지
• branch 조건을 뒤집었을 때 순서가 바뀌는지
• PC가 낮은 경로가 먼저 실행되는지
• lane mask 크기 또는 workload 크기가 path 선택에 영향을 주는지

해석 주의점

Divergence 실험은 Warp scheduler 자체의 Warp 간 issue policy를 직접 보여주지는 않는다.

다만 다음을 확인하는 데 유효하다.

• branch path 실행 순서
• reconvergence 방식
• active mask 변화
• Warp 내부 path serialization 특성

---

추천 실험 순서

순서실험목적

1	Mixed Workload Probe	Warp 간 workload 차이가 있을 때 추월 현상이 발생하는지 확인
2	Barrier Release Probe	모든 Warp가 동시에 Ready 상태가 되었을 때 순수 issue 우선순위 확인
3	Resource Contention Probe	자원 경합과 stall 해제 이후 scheduler arbitration 확인
4	Divergence / Active Mask Probe	Warp 내부 branch path 실행 순서 확인

---

측정 방식

단순히 결과 배열의 write 순서만 보는 것은 부족하다.

Global memory write 순서는 다음 요인의 영향을 받을 수 있다.

• memory coalescing
• store buffer
• L1/L2 cache 상태
• write combining
• pipeline retire timing
• compiler optimization

따라서 각 Warp의 lane 0에서 다음 값을 함께 기록한다.

struct WarpTrace {
    uint64_t start;
    uint64_t end;
    int warp_id;
    int workload;
    int mode;
};

기본 기록 항목은 다음과 같다.

필드의미

warp_id	block 내부 Warp ID
start	관찰 구간 시작 cycle
end	관찰 구간 종료 cycle
workload	부여된 반복 횟수 또는 부하 강도
mode	ALU, memory, conflict 등 실험 조건

---

분석 기준

1. Start Order

Barrier 이후 또는 실험 구간 진입 직후 어떤 Warp가 먼저 timestamp를 기록했는지 본다.

이는 issue priority를 추정하는 데 사용한다.

2. End Order

어떤 Warp가 먼저 완료되었는지 본다.

이는 workload 차이와 stall 차이가 실제 완료 순서에 반영되는지 확인하는 데 사용한다.

3. Duration

duration = end - start

Warp별 실행 시간이 workload 또는 contention 강도에 비례하는지 확인한다.

4. Stability Across Runs

동일한 커널을 여러 번 반복 실행했을 때 순서가 안정적인지 확인한다.

• 항상 같은 순서: deterministic scheduler 가능성
• 부분적으로 같은 순서: scheduler partition 또는 warp ID group 가능성
• 자주 바뀜: dynamic scheduling 또는 외부 상태 의존성 가능성

---

실험 결론 프레임

후속 실험의 목표는 하나의 결과로 스케줄러를 단정하는 것이 아니다.

목표는 다음 가설들을 단계적으로 제거하는 것이다.

가설확인 방법

Warp ID 고정 우선순위	Barrier Release에서 항상 같은 ID가 먼저 시작되는지 확인
Round-Robin	반복 실행에서 시작 Warp가 순환하는지 확인
Greedy / Ready-first	Mixed Workload에서 가벼운 Warp가 무거운 Warp를 추월하는지 확인
Scoreboard 의존	Memory Stall 해제 후 재개 순서 확인
Resource arbitration 영향	Shared Memory conflict 상황에서 재진입 순서 확인

---

다음 구현 대상

가장 먼저 구현할 실험은 Mixed Workload Probe이다.

이 실험은 기존 Same Workload Baseline에서 반복 횟수만 Warp ID별로 차등화하면 되므로 구현 비용이 낮고, baseline과 직접 비교하기 쉽다.

연결 구조

Same Workload Baseline
        ↓
Mixed Workload Probe
        ↓
Barrier Release Probe
        ↓
Resource Contention Probe

현재 단계의 결론

동일 Task Baseline은 스케줄러 편향을 관찰하기 위한 출발점으로는 유효하지만, 스케줄링 정책을 확정하기에는 조건이 너무 대칭적이다.

따라서 다음 단계에서는 Warp별 workload를 비대칭화하여, 스케줄러가 단순 ID 순서로 진행하는지, Ready 상태와 workload 차이에 따라 동적으로 선택하는지 확인한다.