GPU 친화적 데이터 구조 - 해시 테이블과 메모리 접근 최적화

GPU 친화적 데이터 구조 - 해시 테이블과 메모리 접근 최적화

로봇이 미지의 환경을 탐사하며 실시간으로 3D 지도를 생성하는 SLAM(Simultaneous Localization and Mapping) 및 내비게이션 작업에 있어, 데이터를 저장하고 접근하는 아키텍처는 전체 시스템의 성능을 결정짓는 가장 중요한 요소이다. 특히 NvBlox가 표방하는 ’GPU 가속’의 이점을 온전히 누리기 위해서는 GPU의 하드웨어적 특성—수천 개의 코어가 동시에 데이터를 처리하는 SIMT(Single Instruction, Multiple Threads) 아키텍처—에 완벽하게 부합하는 데이터 구조가 필수적이다. 본 절에서는 NvBlox가 채택한 GPU 친화적 데이터 구조의 핵심인 ‘블록 해싱(Block Hashing)’ 기반의 희소(Sparse) 복셀 그리드와, 이를 뒷받침하는 해시 테이블의 구현 원리, 그리고 GPU 메모리 대역폭을 극한까지 활용하기 위한 메모리 접근 최적화(Memory Access Optimization) 전략에 대해 심층적으로 분석한다.

1. GPU 컴퓨팅 패러다임과 데이터 구조의 불일치성 극복

NvBlox의 데이터 구조를 이해하기 위해서는 먼저 기존의 CPU 기반 매핑 시스템들이 GPU로 포팅될 때 겪는 근본적인 한계와 이를 극복하기 위한 패러다임의 전환을 직시해야 한다. OctoMap이나 Voxblox와 같은 CPU 중심의 라이브러리들은 주로 옥트리(Octree)나 포인터 기반의 해시 맵을 사용한다. 이러한 구조는 CPU의 거대하고 복잡한 캐시 계층, 정교한 분기 예측(Branch Prediction), 그리고 비순차적 명령어 처리(Out-of-Order Execution) 기능에는 적합할지 모르나, GPU 환경에서는 심각한 성능 병목을 초래한다.

1.1 포인터 추적(Pointer Chasing)의 비효율성

기존의 트리(Tree) 기반 구조, 특히 옥트리(Octree)는 부모 노드에서 자식 노드로 이동하기 위해 메모리 주소를 참조하여 점프하는 과정이 반복된다. 이를 ’포인터 추적(Pointer Chasing)’이라 한다. CPU는 이러한 랜덤 메모리 접근(Random Memory Access)을 처리하는 데 비교적 능숙하며, 깊은 캐시 계층을 통해 지연 시간을 완화할 수 있다. 그러나 GPU는 수천 개의 스레드가 동시에 메모리를 요청할 때 최고의 성능을 발휘하도록 설계되었다.

만약 GPU 상에서 트리 순회를 구현하게 되면, 각 스레드가 서로 다른 메모리 주소를 가리키는 포인터를 따라 개별적으로 이동하게 된다. 이는 GPU 메모리 컨트롤러에 엄청난 과부하를 주며, 메모리 트랜잭션 효율을 급격히 떨어뜨린다. GPU 메모리 접근의 지연 시간(Latency)은 수백 클럭 사이클에 달하는데, 포인터를 따라갈 때마다 이 지연 시간을 온전히 감내해야 한다. CPU와 달리 GPU는 개별 스레드의 지연 시간을 줄이기보다는, 수만 개의 스레드를 동시에 실행하여 지연 시간을 ’은폐(Hide)’하는 전략을 취한다.1 그러나 포인터 추적은 스레드들이 각기 다른 시점에 메모리 대기를 하게 만들어, 워프(Warp) 스케줄러가 효율적으로 스레드를 교체하는 것을 방해한다.

1.2 스레드 발산(Thread Divergence)과 직렬화

트리 구조의 또 다른 문제는 탐색 깊이가 복셀의 위치마다 달라질 수 있다는 점이다. 동일한 워프(NVIDIA GPU의 기본 실행 단위인 32개 스레드 그룹) 내의 스레드들이 서로 다른 실행 경로를 타게 되면 ’스레드 발산(Thread Divergence)’이 발생한다. 예를 들어, 어떤 스레드는 리프(Leaf) 노드에 도달하여 연산을 수행하는 반면, 다른 스레드는 여전히 트리의 중간 노드를 순회하고 있다면, GPU는 이들의 실행을 병렬로 처리하지 못하고 직렬화(Serialize)해야 한다. 이는 GPU의 병렬 처리 능력을 심각하게 저하시키는 요인이다.

따라서 NvBlox는 이러한 트리 구조 대신, O(1)의 접근 시간을 보장하며 메모리 접근 패턴을 예측 가능하게 만드는 공간 해싱(Spatial Hashing) 기법을 채택하였다. 이는 3차원 공간을 균일한 격자로 나누되, 실제 데이터가 존재하는 영역(표면 근처)에만 메모리를 할당하고, 이를 해시 테이블로 관리하는 방식이다. 이는 희소성(Sparsity)을 유지하면서도 GPU의 병렬성을 극대화하는 핵심 전략이다.3

2. NvBlox의 2단계 계층 구조 (Two-Level Hierarchy)

NvBlox는 GPU 메모리의 효율적 사용과 빠른 접근 속도라는 두 마리 토끼를 잡기 위해 ’2단계 계층 구조’를 사용한다. 이는 전체 맵을 관리하는 희소한(Sparse) 해시 테이블 레이어와, 실제 복셀 데이터를 담고 있는 밀집된(Dense) 블록 레이어로 나뉜다. 이 구조는 Voxblox의 설계를 계승하면서도 GPU 아키텍처에 맞게 최적화된 형태이다.

[표 2-1] NvBlox 2단계 계층 구조의 역할 및 특징

계층 (Hierarchy Level)구성 요소 (Component)데이터 특성 (Data Characteristic)메모리 위치 (Memory Location)주요 역할 (Primary Role)
Level 1해시 테이블 (Global Index)희소 (Sparse)GPU Global Memory전체 3D 공간의 인덱싱, VoxelBlock의 위치 추적
Level 2복셀 블록 (Local Block)밀집 (Dense)GPU Global Memory (Contiguous)실제 센서 데이터(TSDF, Color, ESDF) 저장 및 연산
  1. 1단계: 해시 테이블 (Global Sparse Indexing)
  • 전체 3D 공간을 무한한 복셀 그리드로 가정할 때, 유의미한 데이터(표면 근처)가 존재하는 블록의 인덱스(Block Index)만을 키(Key)로 하여 관리한다.
  • 이 해시 테이블은 GPU 글로벌 메모리(Global Memory)에 상주하며, CUDA 커널(Kernel) 함수 내에서 수만 개의 스레드가 병렬적으로 조회(Lookup) 및 삽입(Insertion)을 수행한다.
  • 키(Key)는 3차원 정수 좌표 (x, y, z)이며, 값(Value)은 해당 좌표에 대응하는 ‘VoxelBlock’ 구조체의 포인터(또는 인덱스)이다.
  1. 2단계: 복셀 블록 (Local Dense Storage)
  • 하나의 VoxelBlock은 8 \times 8 \times 8 크기의 복셀 집합을 포함한다. 즉, 512개의 복셀이 하나의 단위로 묶여 있다.3
  • 이 블록 내부는 빽빽하게(Dense) 채워진 3차원 배열(실제로는 1차원 선형 메모리)로 구성되어 있어, 블록 내부의 특정 복셀에 접근할 때는 해시 조회 없이 단순한 오프셋(Offset) 계산만으로 접근이 가능하다.

이러한 계층 구조는 ‘희소성은 해시로, 지역성은 블록으로’ 해결한다는 철학을 담고 있다. 3D 공간의 대부분은 빈 공간(Free Space)이므로 해시 테이블을 통해 불필요한 메모리 할당을 방지하고, 표면 근처의 데이터는 공간적 지역성(Spatial Locality)이 높으므로 512개의 복셀을 뭉쳐서 처리함으로써 메모리 접근 효율을 극대화한다.

3. GPU 기반 해시 테이블의 구현: 개방 주소법과 선형 탐사

NvBlox의 성능을 지탱하는 허리는 바로 GPU 상에서 동작하는 해시 테이블이다. CPU에서의 해시 테이블 구현(예: C++ 표준 라이브러리의 std::unordered_map)은 주로 체이닝(Chaining, 연결 리스트) 방식을 사용하여 해시 충돌(Collision)을 해결한다. 그러나 GPU에서 연결 리스트를 동적으로 할당하고 순회하는 것은 앞서 언급한 포인터 추적 문제를 야기하므로 금기시된다. 따라서 NvBlox는 개방 주소법(Open Addressing) 기반의 해시 테이블을 사용한다.6

3.1 개방 주소법 (Open Addressing)과 선형 탐사 (Linear Probing)

개방 주소법은 해시 충돌이 발생했을 때 별도의 메모리 공간(연결 리스트)을 찾지 않고, 해시 테이블 배열 내의 다른 빈 슬롯(Bucket)을 찾아 데이터를 저장하는 방식이다. NvBlox가 활용하는 stdgpu 라이브러리 또는 이에 기반한 자체 해시 구현체는 주로 선형 탐사(Linear Probing) 기법을 활용한다.7

  • 동작 메커니즘: 어떤 키 K에 대한 해시 값 h(K)가 가리키는 슬롯이 이미 다른 데이터에 의해 점유되어 있다면, h(K)+1, h(K)+2, \dots 순서로 다음 슬롯을 순차적으로 확인한다.
  • GPU 적합성: 선형 탐사는 메모리를 연속적으로 읽기 때문에 캐시 적중률(Cache Hit Rate)이 매우 높다. GPU의 L2 캐시는 일반적으로 128바이트 라인 단위로 데이터를 가져오는데, 선형 탐사는 인접한 메모리 슬롯을 확인하므로 한 번의 메모리 페치(Fetch)로 여러 번의 탐사(Probe)를 수행할 수 있는 확률이 높다. 이는 메모리 대역폭이 성능의 병목이 되는 GPU 환경에서 매우 유리한 특성이다.

비록 선형 탐사는 데이터가 특정 영역에 뭉치는 클러스터링(Clustering) 현상에 취약하다는 이론적 단점이 있으나, GPU의 막대한 메모리 대역폭과 병렬성은 이를 상쇄한다. 또한, 쿠쿠 해싱(Cuckoo Hashing)과 같은 복잡한 충돌 해결 방식은 다중 해시 함수 계산과 빈번한 데이터 이동을 유발하여 GPU의 SIMT 구조에서 제어 흐름의 복잡성을 증가시킬 수 있다. 반면 선형 탐사는 제어 흐름이 단순하여 워프 발산을 최소화할 수 있다.7

3.2 병렬 삽입과 원자적 연산 (Parallel Insertion & Atomics)

수만 개의 스레드가 동시에 새로운 맵 데이터를 해시 테이블에 삽입하려 할 때, 경쟁 상태(Race Condition)가 발생할 수 있다. 예를 들어, 서로 다른 두 스레드가 동시에 동일한 빈 슬롯에 데이터를 쓰려고 시도할 수 있다. NvBlox는 이를 해결하기 위해 CUDA의 원자적 연산(Atomic Operations), 특히 atomicCAS (Compare-And-Swap)를 사용하여 Lock-Free 구조를 구현한다.

  1. 스레드는 해시 함수를 통해 슬롯 위치를 계산한다.
  2. 해당 슬롯이 비어있다면 atomicCAS를 통해 자신의 키를 기록한다. atomicCAS는 하드웨어 수준에서 원자성을 보장하므로, 동시에 여러 스레드가 접근하더라도 단 하나의 스레드만이 성공적으로 값을 쓸 수 있다.
  3. 만약 다른 스레드가 먼저 선점했다면(충돌 발생), 선형 탐사 규칙에 따라 다음 슬롯으로 이동하여 재시도한다.

이 과정은 명시적인 락(Lock)을 걸지 않기 때문에, 대규모 병렬 환경에서도 데드락(Deadlock) 위험 없이 높은 처리량(Throughput)을 유지할 수 있게 한다. 특히 최신 NVIDIA GPU 아키텍처(Ampere, Ada Lovelace 등)에서는 원자적 연산의 성능이 비약적으로 향상되어, 이러한 Lock-Free 해시 테이블의 오버헤드는 무시할 수 있는 수준이다.10

3.3 로드 팩터(Load Factor) 관리와 메모리 할당 전략

개방 주소법의 치명적인 단점은 테이블이 꽉 찰수록 충돌 횟수가 기하급수적으로 증가하여 성능이 저하된다는 점이다. 일반적으로 로드 팩터(전체 슬롯 대비 사용 중인 슬롯의 비율)가 0.7~0.8을 넘어서면 성능 저하가 뚜렷해진다. 따라서 NvBlox는 해시 테이블의 로드 팩터를 엄격하게 관리한다.

  • 초기 할당: NvBlox는 초기화 시점에 GPU 메모리에 충분히 큰 해시 테이블 버퍼를 미리 할당한다. 이는 런타임 중에 메모리를 재할당(Reallocation)하는 비용을 피하기 위함이다. GPU에서의 메모리 할당(cudaMalloc)은 CPU와의 동기화를 유발하는 무거운 작업이다.
  • 리사이징(Resizing): 만약 로드 팩터가 임계값을 초과하면, NvBlox는 더 큰 버퍼를 할당하고 기존 데이터를 모두 이동시키는 리사이징(Rehashing) 과정을 거친다. stdgpu와 같은 라이브러리는 이를 효율적으로 처리하기 위한 커널을 제공한다.11 하지만 실시간 로봇 애플리케이션에서는 이러한 “세상을 멈추는(Stop-the-world)” 리사이징이 프레임 드랍을 유발할 수 있으므로, 초기 용량을 넉넉하게 설정하거나 순환 버퍼(Ring Buffer) 개념을 도입하여 오래된 데이터를 삭제하는 전략이 병행되기도 한다.

4. VoxelBlock: 8 \times 8 \times 8 구조의 비밀과 메모리 병합 접근 (Memory Coalescing)

해시 테이블을 통해 특정 공간의 VoxelBlock을 찾았다면, 이제 실제 거리 값(TSDF)이나 점유 확률을 읽고 써야 한다. 여기서 NvBlox가 8 \times 8 \times 8 크기의 블록을 선택한 이유는 GPU의 메모리 병합 접근(Memory Coalescing) 특성과 밀접한 관련이 있다.

4.1 메모리 병합 접근의 원리

GPU의 메모리 컨트롤러는 개별 바이트 단위가 아니라, 32바이트, 64바이트, 혹은 128바이트 단위의 트랜잭션으로 데이터를 전송한다. 워프 내의 32개 스레드가 동시에 글로벌 메모리에 접근할 때, 이들이 요청하는 주소가 연속적인 메모리 공간에 위치해 있다면, GPU는 이를 하나의 트랜잭션으로 묶어서(Coalesce) 처리할 수 있다.3

반대로 스레드들이 서로 멀리 떨어진 주소에 접근하거나(Strided Access), 무작위 주소에 접근하면 여러 번의 트랜잭션이 발생하여 유효 대역폭(Effective Bandwidth)이 급격히 감소한다. 이는 고속도로에 32명의 승객이 각각 별도의 승용차를 타고 가는 것(비효율적)과, 버스 한 대에 타고 가는 것(효율적)의 차이로 비유할 수 있다.

4.2 블록 내부의 선형 메모리 배치와 인덱싱

NvBlox의 VoxelBlock 내 512개 복셀은 3차원 배열처럼 논리적으로 구성되지만, 물리적으로는 1차원 배열로 평탄화(Flatten)되어 연속적으로 저장된다. 인덱싱 공식은 일반적으로 다음과 같다:
\text{Index} = z \cdot (W \cdot H) + y \cdot W + x
여기서 W=8, H=8이다.

GPU 커널이 VoxelBlock을 처리할 때, 보통 블록의 x축(가로) 방향을 스레드 인덱스(threadIdx.x)와 매핑한다. 이렇게 하면 워프 내의 인접한 스레드(0번 스레드, 1번 스레드…)가 복셀 배열의 인접한 인덱스(0번 복셀, 1번 복셀…)에 접근하게 된다.

  • 결과 분석: 32개의 스레드가 4바이트(float) 데이터를 읽을 때, 총 32 \times 4 \text{bytes} = 128 \text{bytes}의 연속된 데이터를 요청하게 된다. 이는 최신 NVIDIA GPU의 L1/L2 캐시 라인 크기 또는 메모리 트랜잭션 단위(128바이트)와 정확히 일치한다. 결과적으로 단 한 번의 메모리 트랜잭션으로 워프 전체의 데이터 로딩이 완료된다. 이것이 바로 NvBlox가 Voxblox 대비 177배 빠른 표면 재구성 속도를 달성할 수 있는 핵심적인 하드웨어 레벨의 최적화 비결이다.

4.3 왜 8 \times 8 \times 8인가? 최적의 블록 크기 선정 이유

왜 하필 8인가? 16이나 4는 안 되는가? 이에 대한 답은 워프 사이즈와 희소성 간의 트레이드오프에서 찾을 수 있다.

  1. 워프 사이즈와의 조화: 8은 32(워프 사이즈)의 약수이다. 8 \times 8 \times 8 = 512개의 복셀은 정확히 16개의 워프(16 \times 32 = 512)로 처리될 수 있어 스레드 낭비가 없다. 만약 블록 크기가 워프 사이즈의 배수와 맞지 않는다면, ’잔여 스레드(Remainder Threads)’가 발생하여 마스킹 처리(Masking)로 인한 연산력 낭비가 발생한다.14
  2. 표면 대 부피 비(Surface-to-Volume Ratio)와 희소성: 블록이 너무 작으면(예: 4 \times 4 \times 4), 해시 테이블의 항목 수가 늘어나 인덱싱 오버헤드가 커진다. 반면 블록이 너무 크면(예: 16 \times 16 \times 16), 표면을 포함하지 않는 빈 복셀들까지 불필요하게 메모리를 차지하게 되어 희소성(Sparsity)의 이점이 줄어든다. 16 \times 16 \times 16 블록은 4096개의 복셀을 포함하는데, 얇은 표면(Surface)만 저장하는 TSDF의 특성상 대부분의 복셀이 낭비될 가능성이 크다. 8 \times 8 \times 8은 이러한 메모리 효율성과 인덱싱 오버헤드 사이에서 경험적으로 검증된 최적의 균형점(Sweet Spot)이다.

5. 데이터 레이아웃: SoA (Structure of Arrays) vs AoS (Array of Structures)

데이터 구조 설계의 또 다른 중요한 축은 데이터를 구조체의 배열(AoS)로 저장할지, 배열의 구조체(SoA)로 저장할지 결정하는 것이다. NvBlox는 GPU 캐시 효율을 극대화하기 위해 SoA (Structure of Arrays) 방식을 적극적으로 채용한다.6

5.1 AoS와 SoA의 비교

  • AoS (Array of Structures):
struct Voxel {
float distance;
float weight;
uchar color;
};
Voxel block;

AoS 방식은 객체지향적 관점에서는 직관적이다. 하나의 복셀에 대한 모든 정보가 인접해 있다. 그러나 GPU 커널이 TSDF 통합을 수행할 때, 오직 distance와 weight 값만 필요하고 color는 필요하지 않은 경우가 많다. AoS 방식에서는 distance를 읽을 때 인접한 color 데이터까지 캐시 라인에 딸려오게 된다. 이는 사용하지 않을 데이터를 메모리 버스에 태우는 것으로, 대역폭 낭비(Bandwidth Wastage)를 초래한다.

  • SoA (Structure of Arrays):

C++

struct VoxelBlock {
float distances;
float weights;
uchar colors;
};

SoA 방식은 속성별로 배열을 분리한다. distance 배열은 distance끼리 뭉쳐 있다. 따라서 거리 업데이트 커널은 distances 배열만 연속적으로 읽어오면 되므로, 메모리 버스의 모든 비트를 유효한 데이터 전송에 사용할 수 있다.

5.2 NvBlox의 하이브리드 접근 및 성능 이점

NvBlox 및 이를 기반으로 한 최신 연구(coVoxSLAM 등)는 SoA 레이아웃을 통해 TSDF 통합 속도를 비약적으로 향상시켰다. 연구 결과에 따르면 SoA 방식은 AoS 대비 메모리 읽기/쓰기 처리량을 크게 높여주며, 특히 데이터 필드가 많아질수록(예: 시맨틱 레이블, 신뢰도 점수 등 추가) 그 격차는 벌어진다. NvBlox의 VoxelBlock 구조는 이러한 SoA 설계를 통해 특정 연산(예: ESDF 계산을 위한 거리 값 전파)에서 불필요한 데이터 로딩을 방지하고 커널 실행 속도를 높인다. 이는 제한된 메모리 대역폭을 가진 젯슨(Jetson)과 같은 임베디드 GPU 환경에서 더욱 결정적인 성능 차이를 만든다.6

6. 해시 충돌과 최악의 경우(Worst-Case) 시나리오 분석

현실 세계의 로봇 매핑에서 해시 충돌은 피할 수 없는 문제이다. NvBlox는 이를 어떻게 견디는가?

6.1 공간적 지역성과 해시 함수의 설계

로봇이 스캔하는 환경은 연속적이다. 즉, 새로운 블록이 생성될 때 그 블록의 좌표는 이전 블록의 좌표와 인접해 있을 확률이 높다. 만약 단순한 해시 함수를 사용하면 인접한 좌표들이 비슷한 해시 값을 가져 충돌이 빈번해질 수 있다. NvBlox는 이를 방지하기 위해 XOR 기반의 공간 해싱 함수나 Z-order curve (Morton Code) 등을 활용하여 공간적 인접성이 해시 테이블 상에서는 무작위로 분산되도록 설계한다.

일반적인 해시 함수의 형태는 다음과 같다:
H(x, y, z) = ((x \cdot p_1) \oplus (y \cdot p_2) \oplus (z \cdot p_3)) \mod N
여기서 p_1, p_2, p_3는 큰 소수(Prime Number)이며, \oplus는 XOR 연산이다. XOR 연산은 GPU 하드웨어에서 매우 적은 사이클로 처리되면서도 비트 레벨의 엔트로피를 효과적으로 섞어주는 특성이 있다.16

6.2 지연 시간 은폐 (Latency Hiding)

선형 탐사에서 충돌이 많아지면(탐사 길이가 길어지면) 원하는 데이터를 찾기 위해 여러 번 메모리를 읽어야 하므로 레이턴시가 증가한다. 하지만 GPU의 막대한 스레드 병렬성은 이러한 레이턴시를 ’은폐(Hide)’하는 데 탁월하다. 즉, 하나의 워프가 메모리를 읽느라 대기하는(Stall) 동안, 워프 스케줄러(Warp Scheduler)는 대기 상태가 아닌 다른 워프를 즉시 실행시킨다.2 NvBlox의 아키텍처는 이러한 GPU의 지연 시간 은폐(Latency Hiding) 능력을 믿고, 복잡한 충돌 해결 로직 대신 단순하고 빠른 선형 탐사를 선택한 것이다. 이는 CPU와는 전혀 다른 접근법이며, GPU 하드웨어에 대한 깊은 이해 없이는 내릴 수 없는 설계적 결단이다.

7. 결론: 하드웨어를 이해한 소프트웨어 설계의 승리

요약하자면, NvBlox의 “2.1.2 GPU 친화적 데이터 구조“는 단순히 데이터를 저장하는 방식의 변화가 아니다. 이는 CPU의 직렬 처리 패러다임을 완전히 버리고, GPU의 SIMT 아키텍처, 메모리 계층 구조, 트랜잭션 특성을 밑바닥부터 고려하여 재설계한 결과물이다.

  1. 희소 해시 테이블은 불필요한 메모리 사용을 억제하고 O(1) 접근을 보장한다.
  2. 개방 주소법과 선형 탐사는 포인터 추적을 제거하고 캐시 지역성을 높여 지연 시간을 최소화한다.
  3. 8 \times 8 \times 8 VoxelBlock은 워프 크기와 캐시 라인 크기에 딱 맞아떨어지는 메모리 병합 접근을 가능케 한다.
  4. SoA 레이아웃은 필요한 데이터만 선별적으로 로드하여 유효 메모리 대역폭(Effective Memory Bandwidth)을 극대화한다.

이러한 설계 덕분에 NvBlox는 Voxblox 대비 수십 배에서 백 배 이상의 성능 향상을 이뤄낼 수 있었으며, 이는 제한된 컴퓨팅 자원을 가진 엣지 디바이스(Jetson Orin 등)에서도 실시간으로 고해상도 3D 매핑을 가능하게 하는 원동력이 된다. 이는 로봇 소프트웨어 엔지니어들에게 하드웨어의 특성을 고려한 데이터 구조 설계가 얼마나 파괴적인 성능 혁신을 가져올 수 있는지를 보여주는 교과서적인 사례라 할 수 있다.

8. 참고 자료

  1. Understanding Latency Hiding on GPUs - Berkeley EECS, https://www2.eecs.berkeley.edu/Pubs/TechRpts/2016/EECS-2016-143.pdf
  2. How to understand the “hide latency” - CUDA Programming and Performance, https://forums.developer.nvidia.com/t/how-to-understand-the-hide-latency/258938
  3. nvblox: GPU-Accelerated Incremental Signed Distance Field Mapping - arXiv, https://arxiv.org/html/2311.00626v2
  4. Technical Details - NVIDIA Isaac ROS, https://nvidia-isaac-ros.github.io/concepts/scene_reconstruction/nvblox/technical_details.html
  5. Real-Time Metric-Semantic Mapping for Autonomous Navigation in Outdoor Environments, https://arxiv.org/html/2412.00291v1
  6. coVoxSLAM: GPU accelerated globally consistent dense SLAM - arXiv, https://arxiv.org/html/2410.21149v1
  7. Optimizing Open Addressing - Max Slater, https://thenumb.at/Hashtables/
  8. Maximizing Performance with Massively Parallel Hash Maps on GPUs - NVIDIA Developer, https://developer.nvidia.com/blog/maximizing-performance-with-massively-parallel-hash-maps-on-gpus/
  9. Hashtable Implementation Using Cuckoo Hashing - Warpzone Studios, https://warpzonestudios.com/hashtable-cuckoo/
  10. WarpCore: A Library for fast Hash Tables on GPUs | Request PDF - ResearchGate, https://www.researchgate.net/publication/351163278_WarpCore_A_Library_for_fast_Hash_Tables_on_GPUs
  11. Nvblox quickstart example crashing - Isaac ROS - NVIDIA Developer Forums, https://forums.developer.nvidia.com/t/nvblox-quickstart-example-crashing/333188
  12. Analyzing memory access coalescing of my CUDA kernel - Stack Overflow, https://stackoverflow.com/questions/13771538/analyzing-memory-access-coalescing-of-my-cuda-kernel
  13. definition - In CUDA, what is memory coalescing, and how is it achieved? - Stack Overflow, https://stackoverflow.com/questions/5041328/in-cuda-what-is-memory-coalescing-and-how-is-it-achieved
  14. Block size and occupancy - CUDA - NVIDIA Developer Forums, https://forums.developer.nvidia.com/t/block-size-and-occupancy/317352
  15. coVoxSLAM: GPU Accelerated Globally Consistent Dense SLAM | Request PDF, https://www.researchgate.net/publication/395224211_coVoxSLAM_GPU_Accelerated_Globally_Consistent_Dense_SLAM
  16. A Simple GPU Hash Table - Nosferalatu, https://nosferalatu.com/SimpleGPUHashTable.html