BEVPoolV2

Computer Vision 2026. 7. 24. 13:46

Bevpoolv2 연산에서 rank랑 interval 연산이 하는 역할이 뭐야? 그리고 카메라 calibration이 동적으로 들어가야 돼? 정적으로 들어가야돼?

BEVPoolv2는 LSS(Lift-Splat-Shoot) 기반의 2D→3D 뷰 변환(View Transformation) 과정을 병렬 연산에 최적화하여 speed와 memory 사용량을 파격적으로 줄인 연산입니다.

질문해주신 1) Rank와 Interval의 역할과 2) 카메라 Calibration의 동적/정적 처리를 명확하게 정리해 드리겠습니다.

1. BEVPoolv2에서 Rank와 Interval의 역할

2D 특징 맵(Feature Map)과 예측된 Depth 확률을 곱하면 3D 공간상의 시깔(Frustum) 포인트들이 만들어집니다. 이 수많은 포인트들을 BEV 격자(Grid Voxel)에 투영할 때 "어떤 포인트가 어떤 BEV 그리드 셀로 들어갈 것인가?"를 관리하고 빠르게 합산(Pooling)하기 위해 Rank와 Interval이 쓰입니다.

[Frustum Points]  --->  [Sort by BEV Grid Index]  --->  [Interval Reduction]  --->  [BEV Feature Map]
(ranks_bev, etc.)                                      (interval_starts/lengths)

1) Rank (ranks_depth, ranks_feat, ranks_bev)

Rank는 쉽게 말해 "주소(인덱스) 매핑 표"입니다.

ranks_depth & ranks_feat: 원래 2D 이미지 특징 및 depth 텐서에서 몇 번째 원소를 가져와야 하는지 1D flattened 인덱스로 가리킵니다. BEVPoolv2에서는 거대한 3D Frustum feature를 메모리에 미리 만들어두지 않고, 이 인덱스를 통해 원본 2D feature/depth에서 필요한 순간에만 On-the-fly로 직접 조회를 수행(Gathering)합니다.
CSDN博客
ranks_bev: 각 Frustum 포인트가 3D 공간상에서 "몇 번째 BEV 그리드 셀(Voxel)"에 떨어지는지를 나타내는 BEV 공간의 인덱스입니다.
CSDN博客

2) Interval (interval_starts, interval_lengths)

CSDN

Rank를 기준으로 포인트를 정렬하면, 동일한 BEV 셀에 떨어지는 포인트들이 연속된 구간(Interval)으로 뭉치게 됩니다. Interval은 이 연속된 구간의 시작과 길이를 정의합니다.

interval_starts: 특정 BEV 셀에 들어갈 포인트들의 시작 위치
CSDN
interval_lengths: 해당 BEV 셀에 모인 포인트들의 총 개수
CSDN

핵심 역할: CUDA 스레드 하나가 **BEV 셀 하나(하나의 Interval)**를 담당하게 만듭니다. 스레드는 interval_starts부터 interval_lengths만큼의 포인트를 슬라이딩하며 단순 Add(Sum Pooling)만 수행하면 되므로, Atomics(동기화 병목) 없이 엄청난 속도로 GPU 병렬 가속을 할 수 있습니다.

2. Camera Calibration은 동적(Dynamic)이어야 할까, 정적(Static)이어야 할까?

결론부터 말씀드리면, 실제 차량 양산 및 배포(Deployment) 환경에서는 "동적(Dynamic) Calibration 업데이트 + Look-up Table(Rank/Interval) 재계산" 구조로 가야 합니다.

CSDN

구분정적 (Static Calibration)동적 (Dynamic Calibration)

개념	공장 출고 시 측정한 내/외란 파라미터(Extrinsics/Intrinsics)를 모델에 고정.	차량 주행 중 센서 변형(피칭, 노면 진동, 타이어 공기압, 카메라 재정렬)에 따라 카메라 파라미터가 실시간 업데이트됨.
장점	Rank/Interval을 최초 1회만 계산해두고 constant memory 등에 고정하므로 연산 오버헤드가 거의 없음.	실시간 자세 변화나 카메라 틀어짐에 완벽히 대응하여 3D Bounding Box / Occupancy 투영 오차(Misalignment) 방지.
단점	차량 가속/제동 시 피칭(Pitching), 짐 실림으로 인한 차고 변화, 장기적인 카메라 비틀림이 반영되지 않아 거리 오차 누적.	Extrinsics가 바뀔 때마다 Rank와 Interval을 다시 계산(Pre-computation 업데이트)하는 연산 소모 발생.

엔지니어링 실무에서의 해결책

학습(Training) 단계나 가벼운 벤치마크에서는 정적 calibration으로 둔 채 Rank/Interval을 오프라인 pre-compute하여 사용해도 무방합니다.

하지만 실제 자율주행 차량에 탑재할 때는 다음 중 하나의 전략을 선택합니다.

이벤트 기반 동적 업데이트 (추천): 매 프레임 Rank/Interval을 재계산하는 것은 불필요합니다. Dynamic Calibration / Online Extrinsic Calibration 모듈이 카메라 틀어짐이나 차량 틸팅(Pitch/Roll) 임계값 초과를 감지했을 때만 비동기(Asynchronous) CPU/GPU 스레드에서 Rank와 Interval을 재계산하여 Swap합니다.
Homography / Ego-motion compensation 포함: 3D 투영 좌표계 자체를 Dynamic Extrinsics 기반으로 매 프레임 업데이트하되, BEVPoolv2에 들어가는 Pre-computation 인덱스 계산 로직(C++/CUDA 확장 모듈)을 매우 가볍게 설계하여 1ms 이내로 처리되도록 내장합니다.

2번 방법을 실제 구현하는 방향으로 자세하게 설명해줘

매 프레임 카메라의 동적 파라미터(Dynamic Extrinsics)를 반영하면서도 BEVPoolv2의 핵심 장점인 연산 속도를 유지하기 위해서는, 기존에 오프라인(CPU/Python)에서 수행되던 Pre-computation(좌표 변환, 정렬, Interval 생성) 전체를 하나의 가벼운 Custom CUDA 파이프라인으로 포팅해야 합니다.

실제 양산 수준의 자율주행 파이프라인에서 이 2번 방법을 구현하기 위한 구체적인 아키텍처와 CUDA/C++ 레벨의 최적화 전략을 정리해 드립니다.

1. 전체 구현 파이프라인 아키텍처

Python(PyTorch) 레벨에서 torch.matmul이나 torch.argsort를 사용해 매 프레임 좌표를 변환하고 정렬하면 Python 커널 런칭 오버헤드와 메모리 단편화 때문에 절대 1ms 이내로 들어올 수 없습니다.

따라서 다음의 3단계를 수행하는 C++/CUDA Extension (Custom Operator)을 작성하여 병목을 제거해야 합니다.

Coordinate Transformation Kernel: 동적 Extrinsic을 받아 3D Frustum을 BEV Grid 인덱스로 매핑.
Radix Sort (CUB): 매핑된 인덱스를 고속 정렬.
Interval Generation Kernel: 정렬된 배열에서 Interval 추출.

2. 단계별 CUDA 커널 구현 디테일

Step 1: 동적 좌표 변환 및 필터링 (Coordinate Transformation Kernel)

카메라 이미지 픽셀 $(u, v)$와 깊이(Depth) 를 받아 차량 중심 3D 좌표계로 변환합니다. 매 프레임 업데이트되는 4x4 Extrinsic 행렬 $E_{cam \to ego}$를 커널의 인자로 전달합니다.

연산:
구현 전략:
- 2D Image의 (H, W)와 Depth_bin 차원을 1D 블록/스레드로 펼쳐서(Flatten) 할당합니다.
- 각 스레드는 자신이 담당하는 포인트의 $P_{ego}$를 계산한 뒤, 설정된 BEV Grid 범위(X, Y, Z limit)를 벗어나면 Early Exit(Drop) 처리합니다.
- 유효한 포인트에 대해서만 ranks_bev(BEV 1D 인덱스), ranks_feat(Feature 1D 인덱스), ranks_depth를 글로벌 메모리에 기록합니다.

Step 2: CUB를 이용한 초고속 Key-Value 정렬 (Fast Radix Sort)

동적으로 생성된 ranks_bev를 기준으로 포인트들을 정렬해야 Interval을 만들 수 있습니다.

구현 전략:
- PyTorch의 torch.argsort는 내부적으로 다소 무겁기 때문에, NVIDIA의 Thrust 또는 CUB (cub::DeviceRadixSort::SortPairs)를 직접 호출하는 것이 훨씬 빠릅니다.
- Key는 ranks_bev로, Value는 원본 배열의 인덱스(또는 ranks_feat 자체)로 두고 SortPairs를 수행합니다.
- 정렬할 데이터의 최대 개수(Maximum Frustum Points)는 정해져 있으므로, 정렬용 Temp Storage 공간은 프로그램 초기화 시점에 미리 할당(Pre-allocation)해 두어 cudaMalloc으로 인한 프레임 지연을 방지합니다.

Step 3: Interval Boundary 감지 (Interval Generation Kernel)

정렬된 ranks_bev 배열을 스캔하여, 값이 달라지는 경계(Boundary)를 찾아 interval_starts와 interval_lengths를 기록합니다.

구현 전략:
- 1D 스레드를 띄워 sorted_ranks_bev[i]와 sorted_ranks_bev[i-1]을 비교합니다.
- 값이 다르다면 인덱스 i가 새로운 BEV 셀의 시작점입니다. 해당 위치의 정보를 interval_starts 배열에 원자적 연산(atomicAdd) 없이 직접 기록할 수 있도록 설계합니다.

C++

// Interval 생성 커널 (슈도코드 예시)
__global__ void generate_intervals_kernel(
    const int* sorted_ranks_bev, 
    int* interval_starts, 
    int* interval_lengths, 
    int num_valid_points) 
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid >= num_valid_points) return;

    if (tid == 0) {
        interval_starts[0] = 0;
    } else {
        // 인접한 두 rank가 다르면 새로운 구간의 시작
        if (sorted_ranks_bev[tid] != sorted_ranks_bev[tid - 1]) {
            // 현재 BEV 인덱스를 기반으로 저장 위치 결정
            int current_bev_idx = sorted_ranks_bev[tid];
            int prev_bev_idx = sorted_ranks_bev[tid - 1];
            
            interval_starts[current_bev_idx] = tid;
            // 이전 구간의 길이 계산
            interval_lengths[prev_bev_idx] = tid - interval_starts[prev_bev_idx]; 
        }
    }
}

3. 하드웨어/메모리 최적화 포인트 (Performance Tuning)

커스텀 어텐션이나 병렬 처리 커널을 최적화해 본 경험이 있으시다면, 이 파이프라인에서 가장 중요한 것이 메모리 I/O 최소화라는 것을 잘 아실 것입니다.

Memory Pool Pre-allocation (Zero-Allocation Pipeline):
- 프레임마다 유효한 Frustum 포인트의 개수가 달라지지만 (보통 전체의 40~60%), 메모리 버퍼는 이론상 최대 크기(Batch * Cam * Depth * H * W)로 모델 로드 시점에 딱 한 번만 잡아둡니다.
- 매 프레임마다 버퍼의 일부분만 덮어쓰며 재사용(In-place update)해야 합니다.
Kernel Fusion:
- Step 1(좌표 변환) 커널이 메모리에 결과값을 쓰기 전에, 각 블록 내의 Shared Memory를 활용해 아예 불필요한 포인트를 1차적으로 Culling(압축)한 뒤 글로벌 메모리에 쓰면 이후 Step 2의 Sort 연산량을 크게 줄일 수 있습니다.
반정밀도(FP16 / BF16) 행렬 연산:
- Extrinsic 행렬 곱셈 시 FP32가 아닌 FP16/BF16을 사용하고 Tensor Core 매핑을 유도하면 연산량을 더 줄일 수 있습니다. (다만 거리 오차에 민감한 Z축 값의 정밀도는 확인이 필요합니다.)

이렇게 3단계 로직을 하나의 C++ Extension으로 묶어버리면, 매 프레임 차량의 서스펜션 변화나 피치(Pitch)에 의해 Extrinsic 행렬이 업데이트되더라도 1ms 내외의 오버헤드만으로 정확한 3D 매핑(Rank/Interval)을 실시간으로 갱신할 수 있습니다. 이후 수행되는 BEVPool 연산 자체는 완전히 동일하게 진행됩니다.

그럼 이 과정에서 입력과 결과물이 뭐고 이게 어떻게 bevpoolv2에 들어갈 수 있어?

앞서 설명해 드린 동적 Pre-computation 파이프라인(C++/CUDA)은 실제 무거운 특징(Feature) 벡터를 연산하는 것이 아니라, "어떤 데이터를 어디서 가져와서 어디에 합칠 것인가"에 대한 인덱스(주소) 맵을 실시간으로 만들어내는 과정입니다.

이 파이프라인의 입출력(I/O)과, 이것이 실제 BEVPoolv2 커널과 어떻게 결합(Integration)되는지 데이터 흐름을 명확히 짚어드리겠습니다.

1. 동적 Pre-computation 파이프라인의 입출력

이 과정은 딥러닝 네트워크의 Forward 연산(Convolution 등)과는 독립적으로, 기하학적 투영 정보만 계산합니다.

[입력 (Inputs)]

Dynamic Extrinsics (동적 외란 파라미터): (N_cams, 4, 4)
- 차량의 서스펜션(Pitch, Roll, Z-bounce) 센서나 실시간 캘리브레이션 모듈에서 올라오는 매 프레임 업데이트된 카메라 위치/자세 행렬입니다.
Intrinsics (내란 파라미터): (N_cams, 3, 3)
Grid Configurations (정적 설정값):
- 원본 2D 이미지 크기 (H, W), Depth Bin 개수 (D)
- BEV 공간의 최소/최대 좌표 및 Voxel 해상도 (X_bound, Y_bound, Z_bound, dx, dy, dz)

[출력 (Outputs)]

이 파이프라인의 결과물은 순수하게 1D 인덱스 텐서들입니다. (정렬 및 필터링 완료 상태)

ranks_bev: (N_valid_points,) - 유효한 각 3D 포인트가 속할 최종 BEV 그리드 인덱스
ranks_depth: (N_valid_points,) - 해당 포인트의 Depth 확률값을 가져올 인덱스
ranks_feat: (N_valid_points,) - 해당 포인트의 2D 이미지 특징(Feature)을 가져올 인덱스
interval_starts: (N_unique_bev_cells,) - 각 BEV 셀의 연산 시작점
interval_lengths: (N_unique_bev_cells,) - 각 BEV 셀에 더해져야 할 포인트의 개수

2. 이 결과물이 BEVPoolv2 연산에 들어가는 과정

위에서 만든 5개의 출력 텐서가 이제 진짜 무거운 데이터(Feature, Depth)를 처리하는 BEVPoolv2 CUDA 커널의 입력(Arguments)으로 주입됩니다.

BEVPoolv2 커널의 전체 입력은 다음과 같아집니다.

데이터 텐서 (from Network):
- x (2D Feature Map): (B, N_cams, C, H, W)
- depth (Predicted Depth): (B, N_cams, D, H, W)
인덱스 텐서 (from Dynamic Pipeline):
- ranks_bev, ranks_depth, ranks_feat, interval_starts, interval_lengths

BEVPoolv2 CUDA 커널 내부의 실제 동작 (Gather-Multiply-Reduce)

BEVPoolv2 커널은 GPU 스레드를 띄울 때, 전체 포인트 단위가 아니라 interval 단위(즉, 타겟 BEV 셀 단위)로 블록/스레드를 할당합니다. 이것이 BEVPoolv2가 빠른 핵심 이유입니다.

스레드 할당: GPU의 스레드 하나(혹은 스레드 블록 하나)가 특정 BEV 인덱스 를 담당합니다.
구간 탐색 (Interval Fetch): 이 스레드는 interval_starts[i]부터 interval_lengths[i]만큼의 횟수 동안 for 루프를 돕니다.
데이터 수집 (Gather) 및 곱셈 (Multiply): 루프 안에서 ranks_feat와 ranks_depth를 사용해 원래 배열에서 데이터를 On-the-fly로 읽어와서 곱합니다.
합산 (Reduce / Sum Pooling): 루프를 돌면서 계산된 값들을 누적해서 더합니다.
결과 저장 (Scatter): 루프가 끝나면 누적된 최종 Feature 벡터를 출력 BEV 맵의 ranks_bev[i] 위치에 한 번만 씁니다.

요약: 전체 파이프라인의 그림

Plaintext

[Camera / IMU Sensors] ---> [Dynamic Extrinsics]
                                   |
                                   v
+-------------------------------------------------------------+
| 1. Dynamic Pre-computation (C++/CUDA Custom Kernel)         |
|    - 3D Projection -> Culling -> Radix Sort -> Interval Gen |
+-------------------------------------------------------------+
                                   |
         (ranks_feat, ranks_depth, ranks_bev, intervals)
                                   |
[Image] -> [Backbone 2D] ----> (Features, Depth)
                                   |
                                   v
+-------------------------------------------------------------+
| 2. BEVPoolv2 Forward Kernel (기존 LSS 대체)                 |
|    - Interval 단위 병렬 Thread 실행                         |
|    - Gather(Feature * Depth) -> Sum Reduce -> Scatter       |
+-------------------------------------------------------------+
                                   |
                           [BEV Feature Map] -> (3D Detection 헤드로 전달)

이처럼 기하학적 인덱스 계산(1번)과 실제 텐서 연산(2번)이 분리되어 있기 때문에, Extrinsic이 매 프레임 변하더라도 무거운 Feature Map 복사나 3D Frustum 생성 없이 1번 단계만 가볍게 재수행하여 BEVPoolv2에 그대로 밀어 넣을 수 있습니다.

ABOUT ME

CAMEL CAMEL