nms_kernel.cu 파일 분석 -2

nms_kernel 코드

• 실제 병렬 비교 kernel

---

함수 정의

template <typename scalar_t>
__global__ void nms_kernel(
    const int64_t n_boxes,
    const scalar_t nms_overlap_thresh,
    const scalar_t *dev_boxes,
    const int64_t *idx,
    int64_t *dev_mask) {

• template은 boxes 관련 dtype을 대응하기 위한 것
• __global__은 CUDA 커널(CPU 코드에서 호출되지만 실제 계산은 GPU에서 실행되는 함수) 임을 정의할 때 붙이는 표시

n_boxes 차선 후보 개수 N

nms_overlap_thresh 차선이 비슷한지 판단하는 threshold

dev_boxes GPU 메모리에 있는 boxes 데이터 포인터

idx 점수 높은 순서로 정렬된 index 배열

dev_mask 어떤 후보가 어떤 후보를 제거해야 하는지 저장하는 bit mask 배열

---

row_start, col_start

const int64_t row_start = blockIdx.y;
const int64_t col_start = blockIdx.x;

• 차선 후보들을 64개씩 묶어서 비교하기 위해 이 kernel은 2차원 grid를 사용함
• row_start는 현재 CUDA block이 담당하는 row 후보 묶음 번호이고,
• col_start는 현재 CUDA block이 담당하는 column 후보 묶음 번호이다.
• 따라서 실제 row 후보 시작 index는 row_start * 64, 실제 column 후보 시작 index는 col_start * 64로 계산됨

---

중복 비교 제거

if (row_start > col_start) return;

• A와 B를 비교하는 것과 B와 A를 비교하는 것은 같은 일임 (왜냐하면 절대값으로 비교하기 때문)
• 따라서 비교량을 절반으로 줄이기 위해 다음과 같이 필터 코드를 작성함

---

row_size와 col_size

const int row_size = min(n_boxes - row_start * threadsPerBlock, threadsPerBlock);

const int col_size = min(n_boxes - col_start * threadsPerBlock, threadsPerBlock);

• n_boxes는 전체 차선 후보 수임
• row_start * threadsPerBlock은 현재 row block의 시작 후보 index임
• 따라서 n_boxes - row_start * threadsPerBlock은 현재 block 시작점부터 남아 있는 후보 개수임
• 대부분의 block은 64개 후보를 처리하지만,
• 마지막 block은 64개보다 적을 수 있으므로
• min(..., threadsPerBlock)을 사용해 실제 처리할 후보 개수를 구함

---

shared memory 선언

__shared__ scalar_t block_boxes[threadsPerBlock * PROP_SIZE];

• __shared__는  CUDA 커널(__global__)안에서 같은 block 안의 thread들이 공유하는 빠른 메모리 선언임
• 즉 블록에서 쓸 차선 후보 데이터를 저장할 shared 메모리 선언
• __shared__는 다음 조건일 때 유용
  • 같은 block 안의 여러 thread가 같은 데이터를 반복해서 읽을 때
  • global memory에서 매번 읽으면 비쌀 때
  • block 안에서 데이터를 한 번 가져와서 여러 번 재사용할 수 있을 때
  • shared memory 크기가 너무 크지 않을 때
• threadsPerBlock * PROP_SIZE 은 블럭 내 thread 수 * 차선 후보 1개당 내부 값임
• 쓸 때는 block_boxes[0 ~ 76] 0번째 차선 후보 ... block_boxes[63 * 77 ~ 64 * 77 - 1] 63번째 차선 후보 이런 식으로 하면 됨

---

col block의 후보를 shared memory로 복사

if (threadIdx.x < col_size) {
    for (int i = 0; i <  PROP_SIZE; ++i) {
        block_boxes[threadIdx.x * PROP_SIZE + i] =
            dev_boxes[idx[(threadsPerBlock * col_start + threadIdx.x)] * PROP_SIZE + i];
    }
}

• global memory에 있는 column block의 차선 후보들을 shared memory로 복사하는 부분
• 현재 CUDA block이 비교할 column 후보 묶음을 shared memory에 올리기 위한 역할

if (threadIdx.x < col_size)

• threadIdx.x는 현재 block 안에서의 thread 번호임
• threadIdx.x = 0, 1, 2, ..., 63
• 이때 마지막 block은 후보가 64개보다 적을 수 있으므로 col_size가 64보다 작아질 수 있음
• 하지만 CUDA는 그럼에서 thread 64개를 실행하므로 범위 밖 데이터를 읽지 않도록 제한한 조건절임

for (int i = 0; i < PROP_SIZE; ++i)

• 차선 후보 하나는 PROP_SIZE개의 값(77개)를 가짐
• 즉 앞에서 선언한 shared 매모리에 0부터 73, 74 ~ 74+(77-1),... n ~ n + (77-1) 이런 식으로 연속되게 데이터를 저장해야 함
• 따라서 77 즉 PROP_SIZE개의 값만큼 반복문을 돌림

block_boxes[threadIdx.x * PROP_SIZE + i] = 
	dev_boxes[idx[(threadsPerBlock * col_start + threadIdx.x)] * PROP_SIZE + i];

• block_boxes[threadIdx.x * PROP_SIZE + i]
  • block_boxes는 위에서 정의한 shared memory 배열
  • column block의 후보들이 연속으로 저장됨
  • 이때 각 thread는 자신이 담당한 차선의 정보 77개를 연속적으로 저장해야 함
  • 따라서 자신 담당인 구역을 알기 위해 (threadIdx.x) * (차선 정보 개수)를 더해 저장 위치를 보정함
    • 0번째 thread는 0~76, 1번째 thread는 77~153 = 1*77, 1*77+77 
• dev_boxes[idx[(threadsPerBlock * col_start + threadIdx.x)] * PROP_SIZE + i];
  • global memory에서 읽는 위치
  • 이때 NMS는 원래 boxes 순서가 아니라 점수 높은 순서으로 저장이 되어야 함
    • (순차적으로 저장되는게 아닌 결과적으로 저장이 점수 높은 순서로 된다는 의미)

1. threadsPerBlock * col_start + threadIdx.x
   → column block 안에서 몇 번째 정렬 위치를 담당하는지 값

2. idx[...]
   → 그 점수순 후보가 원래 boxes 배열에서 몇 번 후보인지 구함

3. idx[...] * PROP_SIZE
   → 그 후보의 시작 위치로 이동

4. + i
   → 그 후보 안에서 i번째 값에 접근

---

thread 싱크 맞추기

__syncthreads();

• 같은 block 안의 thread들이 특정 지점까지 모두 도착할 때까지 기다리게 하는 베리어
  • 해당 코드는 같은 block 안의 모든 thread가 shared memory 복사를 끝낼 때까지 기다리는 목적

---

row 후보 비교 시작

if (threadIdx.x < row_size) {

• 각 thread는 row block 안의 후보 하나를 담당함
• 마지막 row block은 후보가 64개보다 적을 수 있으므로 범위를 체크 목적

---

현재 후보 index 계산

const int cur_box_idx = threadsPerBlock * row_start + threadIdx.x;

• 현재 thread가 row block 안에서 몇 번째 정렬 위치를 담당하는지 값

---

현재 후보 포인터

const scalar_t *cur_box = dev_boxes + idx[cur_box_idx] * PROP_SIZE;

• 현재 thread가 담당하는 row 차선 후보 하나의 시작 주소를 구하는 코드
• dev_boxes는 GPU 메모리에 있는 전체 차선 후보 배열

  • dev_boxes[0]  ~ dev_boxes[76]   → boxes[0]
    dev_boxes[77] ~ dev_boxes[153]  → boxes[1]
    dev_boxes[154]~ dev_boxes[230]  → boxes[2]

• cur_box_idx는 현재 thread가 담당하는 후보의 점수순 위치
• idx[cur_box_idx]는 현재 thread가 담당하는 후보의 원래 boxes index
• idx[cur_box_idx] * PROP_SIZE는 현재 후보의 시작 위치를 의미
• dev_boxes + 는 C/C++에서 포인터에 숫자를 더하면, 그 숫자만큼 뒤의 원소 위치로 이동함 
  • 즉 cur_box가 현재 후보의 첫 번째 값 주소를 가리키게 함

---

bit mask 변수