새소식

Perception/Gaussian Splatting

3DGS 소스 코드 분석 : def forward()와 CUDA 프로그래밍

  • -

최근 3DGS 연구를 시작하면서 지도해주시는 선배님이 한 줄 한 줄 설명해주셔서 3DGS 원본 소스코드를 자세히 살펴 볼 기회가 있었다. 방법론만 새로운 것이 아니라 아름다운 수준의 코드로 말도 안 되게 구현되어 있는 것 같아 놀라웠다.
하드웨어 수준의 섬세한 병렬화와 최적화가 간결하게 구현되어 있었다. 이때까지 주먹구구식으로 "이거 입력하면 보이네?" 하고 넘어갔어서 마음이 불편했는데, 좋은 이론을 잘 배울 수 있어서 오랜만에 배우는 기쁨을 느낄 수 있었다. 특히 CUDA 프로그래밍은 예전에 책만 사놓고 잘 안 봤었는데 이참에 입문할 수 있어서 좋았다. 혼자였으면 코드를 제대로 볼 생각도 못 했을 것 같은데 감사한 일... 

GS가 주목받는 이유는 속도가 빠르다는 점 때문인데, 이 비결이 GPU 병렬처리를 효율적으로 활용할 수 있는 문제 모델링과 파이프라인 아키텍쳐 때문이라는 점도 신기했다. 

많은 수의 연산 코어를 집적한 GPU는 CPU에 비해 FLOPS 관점에서 매우 높은 처리 성능을 갖는다. 하지만 이러한 높은 성능은 모든 연산 및 메모리 접근이 완벽히 규칙적일 때 얻을 수 있는 이론적인 수치다. GPU는 제어 장치 및 캐시에 적은 공간을 할애한 만큼 스레드 사이의 불규칙한 연산 흐름 및 임의적인 메모리 접근 형태에 취약하다. 한 스레드 그룹 내 if 문 등의 분기가 있으면 다른 분기를 따르는 스레드의 동작은 직렬화되며 병렬 처리 성능은 크게 떨ㄹ어진다. 불규칙한 메모리 접근 또한 비슷한 이유로 GPU의 성능을 크게 떨어뜨릴 수 있다. 

[CUDA 기반 GPU 병렬 처리 프로그래밍] p.27

배우면서도 이건 클로드나 지피티가 알기 쉽게 설명해주기 어려운 소재이고, 그래서 접근 장벽이 꽤 높을 것 같다는 생각이 들었다. 컴구, 컴그, AI를 다 잘 알아야 🤯 
심지어 로봇을 위한 3DGS를 하려면 로보틱스도 잘 알고 3DGS도 잘 알아야 하니 진입장벽이 미쳤다 

아하 그래서 그런 거구나! 하는 모먼트들

https://github.com/graphdeco-inria/gaussian-splatting

 

GitHub - graphdeco-inria/gaussian-splatting: Original reference implementation of "3D Gaussian Splatting for Real-Time Radiance

Original reference implementation of "3D Gaussian Splatting for Real-Time Radiance Field Rendering" - graphdeco-inria/gaussian-splatting

github.com

Filetree

gs_analysis/
├── train.py                    # 3DGS 모델 학습 메인 스크립트
├── render.py                   # 학습된 모델로 이미지 렌더링
├── convert.py                  # COLMAP을 이용한 데이터 전처리
├── metrics.py                  # PSNR, SSIM, LPIPS 메트릭 계산
├── full_eval.py                # 전체 벤치마크 평가 자동화 스크립트
├── environment.yml             # Conda 환경 설정 파일
├── LICENSE.md                  
├── README.md                   
├── results.md                  
│
├── arguments/
│   └── __init__.py             # CLI 인자 파싱 (ModelParams, PipelineParams, OptimizationParams)
│
├── gaussian_renderer/
│   ├── __init__.py             # Gaussian Splatting 렌더링 핵심 함수 (render)
│   └── network_gui.py          # 실시간 뷰어 GUI 네트워크 통신
│
├── scene/
│   ├── __init__.py             # Scene 클래스 (데이터 로딩 및 관리)
│   ├── cameras.py              # Camera, MiniCam 클래스 (카메라 파라미터 관리)
│   ├── colmap_loader.py        # COLMAP 데이터 파싱 (intrinsics/extrinsics/points3D)
│   ├── dataset_readers.py      # 데이터셋 로더 (COLMAP, Blender 형식 지원)
│   └── gaussian_model.py       # GaussianModel 클래스 (3D Gaussian 파라미터 및 최적화)
│
├── utils/
│   ├── camera_utils.py         # 카메라 로딩 및 JSON 변환 유틸리티
│   ├── general_utils.py        # 일반 유틸리티 (학습률 스케줄러, 회전 행렬 등)
│   ├── graphics_utils.py       # 그래픽스 유틸리티 (투영 행렬, FOV 변환 등)
│   ├── image_utils.py          # 이미지 메트릭 (MSE, PSNR)
│   ├── loss_utils.py           # 손실 함수 (L1, L2, SSIM)
│   ├── sh_utils.py             # Spherical Harmonics 연산 (eval_sh, RGB2SH, SH2RGB)
│   ├── system_utils.py         # 시스템 유틸리티 (디렉토리 생성, iteration 검색)
│   ├── make_depth_scale.py     # Depth 스케일 계산 유틸리티
│   └── read_write_model.py     # COLMAP 모델 읽기/쓰기
│
├── lpipsPyTorch/
│   ├── __init__.py             # LPIPS 메트릭 인터페이스
│   └── modules/
│       ├── lpips.py            # LPIPS 모델 정의
│       ├── networks.py         # 백본 네트워크 (AlexNet, VGG, SqueezeNet)
│       └── utils.py            # LPIPS 유틸리티
│
├── assets/            
├── SIBR_viewers/               # 실시간 뷰어 (C++ 기반)
└── submodules/                 # 외부 서브모듈 (diff-gaussian-rasterization 등)
submodules/
├── diff-gaussian-rasterization/          # 핵심: Differentiable Gaussian Rasterizer
│   ├── setup.py                          # PyTorch CUDA Extension 빌드 설정
│   ├── ext.cpp                           # PyBind11 바인딩 (Python ↔ C++)
│   ├── rasterize_points.cu               # CUDA 래퍼 함수 (PyTorch Tensor ↔ CUDA)
│   ├── rasterize_points.h                # 헤더 파일
│   ├── diff_gaussian_rasterization/
│   │   └── __init__.py                   # Python 인터페이스 (GaussianRasterizer 클래스)
│   ├── cuda_rasterizer/
│   │   ├── forward.cu                    # Forward pass CUDA 커널
│   │   ├── backward.cu                   # Backward pass CUDA 커널 (gradient 계산)
│   │   ├── rasterizer_impl.cu            # Rasterizer 구현체
│   │   ├── auxiliary.h                   # 보조 함수
│   │   └── config.h                      # 설정 (NUM_CHANNELS 등)
│   └── third_party/glm/                  # GLM 수학 라이브러리
│
├── simple-knn/                           # KNN 기반 거리 계산
│   ├── setup.py                          # 빌드 설정
│   ├── ext.cpp                           # PyBind11 바인딩
│   ├── simple_knn.cu                     # KNN CUDA 커널
│   └── spatial.cu                        # distCUDA2 함수 (초기 scale 계산용)
│
└── fused-ssim/                           # Fused SSIM 손실 함수
    ├── setup.py                          # 빌드 설정
    ├── ext.cpp                           # PyBind11 바인딩
    └── ssim.cu                           # SSIM CUDA 커널

여기서는 forward.cu 와 같은 파일에서 병렬 렌더링 코드를 작성하고,  CUDA의 C++ 래퍼 함수들이 (.cu) PyTorch 텐서를 받아서 CUDA 커널에 전달한다. 그리고 그 C++ 함수를 Pybind하여 Python에서 호출 가능하도록 노출시킨다. 

따라서 전체 호출의 흐름은 다음과 같다. 

Python (train.py)
    ↓
gaussian_renderer/__init__.py → render()
    ↓
GaussianRasterizer.forward()
    ↓
_RasterizeGaussians.apply()  (torch.autograd.Function)
    ↓
_C.rasterize_gaussians()  (PyBind11)
    ↓
RasterizeGaussiansCUDA()  (C++ wrapper in rasterize_points.cu)
    ↓
CudaRasterizer::Rasterizer::forward()  (cuda_rasterizer/rasterizer_impl.cu)
    ↓
CUDA Kernels (forward.cu)  ← GPU에서 실행

train.py

루트 폴더에 있는 train.py는 3DGS 모델을 학습하는 코드이다. Densification, Pruning, L1+SSIM 손실, Depth regularization 등을 포함한다. def training은

    lp = ModelParams(parser)
    op = OptimizationParams(parser)
    pp = PipelineParams(parser)
    
	gaussians = GaussianModel(dataset.sh_degree, opt.optimizer_type)
    scene = Scene(dataset, gaussians)

이와 같이 GaussianModel과 Scene 클래스를 초기화한다. 

Class GaussianModel

    def __init__(self, sh_degree, optimizer_type="default"):
        self.active_sh_degree = 0
        self.optimizer_type = optimizer_type
        self.max_sh_degree = sh_degree  
        self._xyz = torch.empty(0)
        self._features_dc = torch.empty(0)
        self._features_rest = torch.empty(0)
        self._scaling = torch.empty(0)
        self._rotation = torch.empty(0)
        self._opacity = torch.empty(0)
        self.max_radii2D = torch.empty(0)
        self.xyz_gradient_accum = torch.empty(0)
        self.denom = torch.empty(0)
        self.optimizer = None
        self.percent_dense = 0
        self.spatial_lr_scale = 0
        self.setup_functions()
  • features_dc : RGB
  • features_rest : SH coefficient
  • scale : log scale임 (GS map scaling)
  • quaternion : normalized quaternion임 - SE(3)여서 Lie group임
  • max_radii2D : adaptive density control
    • 이 변수는 각 Gaussian이 이미지 공간에서 가지는 최대 반지름을 추적함
    • densify_and_prune()에서 사용
      • 화면에서 너무 큰 Gaussian을 제거하는데 사용
        • max_raddi2D > max_screen_size : 이미지 공간에서 반지름이 너무 큰 가우시안 제거
        • get_scaling.max : 월드 공간에서 너무 큰 가우시안 제거
        • opacity < min_opacity : 투명도가 너무 낮은 가우시안 제거
  • max_sh_degree : 학습 초기에는 낮은 차수로 안정적 수렴을 하고, 이후 고차수로 디테일한 view-dependent 효과 학습
    • 0 : Diffuse only (시점 무관)
    • 1 : 기본 방향성
    • 2 : 중간 반사
    • 3 : 높은 반사
  • optimizer : 차원 너무 커짐 이슈로 인해 Adam같은 first-order optimizer (기울기만 사용)를 주로 쓰긴 하지만 요새는 2차 미분기도 사용하는 추세 (Hessian)

멤버 함수들

Property

  • get_scaling
  • get_rotation
  • get_xyz
  • get_features : DC (0차 SH) + SH 계수 합쳐서 반환
  • get_features_dc : DC만 반환
  • get_features_rest : 고차 SH 계수만 반환
  • get_opacity : 0~1 범위 Opacity 반환
  • get_exposure 
  • get_covariance : scale + rotation -> 3D covariance 행렬 계산

모델 생성 및 로드/저장

  • create_from_pcd
    • point cloud에서 Gaussian 초기화
    • RGB는 SH로 변환하고 KNN으로 초기 Scale 계산
    • Rotation은 identity
    • Opacity는 0.1로 초기화
  • save_ply
  • load_ply
  • capture : 체크포인트용 모델 상태 캡쳐
  • restore : 체크포인트에서 모델 복원

학습 설정

  • training_setup
  • update_learning_rate
  • oneupSHdegree

Adaptive Density Control

  • densify_and_prune
    • clone : 작은 가우시안 복제
    • split : 큰 가우시안 분할
    • prune : 불필요한 가우시안 제거
  • densify_and_clone : gradient 크고 scale 작은 가우시안 복제
  • densify_and_split : gradient 크고 scale 큰 가우시안 분할
  • prune_points : 마스크에 해당하는 가우시안 제거
  • add_densification_stats
  • reset_opacity

render.py

3D 가우시안을 2D로 렌더링하는 핵심이다. 논문에서도 나왔듯이 이 렌더링의 핵심은 "Depth 기준으로 정렬된 Gaussian을 가까운 순서대로 알파블렌딩"하는 것이다.

  • 파라미터
    • viewpoint_camera : 렌더링할 카메라
    • pc : GaussianModel 인스턴스
    • bg_color : 배경색
    • scaling_modifier
    • seperate_sh : DC와 Rest SH를 분리할 것인지
    raster_settings = GaussianRasterizationSettings(
        image_height=int(viewpoint_camera.image_height),
        image_width=int(viewpoint_camera.image_width),
        tanfovx=tanfovx,
        tanfovy=tanfovy,
        bg=bg_color,
        scale_modifier=scaling_modifier,
        viewmatrix=viewpoint_camera.world_view_transform,
        projmatrix=viewpoint_camera.full_proj_transform,
        sh_degree=pc.active_sh_degree,
        campos=viewpoint_camera.camera_center,
        prefiltered=False,
        debug=pipe.debug,
        antialiasing=pipe.antialiasing
    )

    rasterizer = GaussianRasterizer(raster_settings=raster_settings)

    means3D = pc.get_xyz
    means2D = screenspace_points
    opacity = pc.get_opacity

여기서는 GaussianRasterizer를 세팅한다. 

나는 처음에는 Rasterize랑 알파블렌딩이 약간 헷갈렸는데 (둘 다 렌더링한다는 거아녀 ?)

  • Rasterize
    • 연속적인 가우시안들을 픽셀 격자 이미지로 바꾸는 전체 과정
    • 3D -> 2D로 투영
    • 해당 픽셀에서 가우시안 값 평가
    • 깊이 정렬 및 타일별 후보 리스트 구성
  • Alpha blending
    • 여러 가우시안이 한 픽셀에 기여할 때 어떻게 섞을지를 정하는 규칙
    • GS에서는 Front-to-back
means3D = pc.get_xyz          # 3D 위치
means2D = screenspace_points  # 2D gradient용
opacity = pc.get_opacity      # 투명도

# Covariance 계산 방식 선택
if pipe.compute_cov3D_python:
    cov3D_precomp = pc.get_covariance(scaling_modifier)  # Python에서 계산
else:
    scales = pc.get_scaling      # CUDA에서 계산하도록
    rotations = pc.get_rotation
if override_color is None:
    if pipe.convert_SHs_python:
        # Python에서 SH → RGB 변환
        dir_pp = pc.get_xyz - viewpoint_camera.camera_center
        dir_pp_normalized = dir_pp / dir_pp.norm(dim=1, keepdim=True)
        sh2rgb = eval_sh(pc.active_sh_degree, shs_view, dir_pp_normalized)
        colors_precomp = torch.clamp_min(sh2rgb + 0.5, 0.0)
    else:
        # CUDA에서 SH → RGB 변환
        shs = pc.get_features

CUDA Rasterization 호출은 아래서 한다. 

rendered_image, radii, depth_image = rasterizer(
    means3D=means3D,
    means2D=means2D,
    shs=shs,
    colors_precomp=colors_precomp,
    opacities=opacity,
    scales=scales,
    rotations=rotations,
    cov3D_precomp=cov3D_precomp
)

이를 통해 실제 렌더링은 CUDA 커널에서 수행하게 된다. 이때 Depth_image도 나온다. 

out = {
    "render": rendered_image,           # 렌더링된 이미지 [3, H, W]
    "viewspace_points": screenspace_points,  # 2D gradient용
    "visibility_filter": (radii > 0).nonzero(),  # 보이는 Gaussian 인덱스
    "radii": radii,                      # 화면 반지름 (densification용)
    "depth": depth_image                 # Inverse depth map
}

CUDA

CUDA(Compute Unified Device Architecture)는 엔비디아에서 GPU를 GPGPU 목적으로 사용할 수 있게 제공하는 프로그래밍 인터페이스이다. CUDA는 C/C++의 확장 언어로 파이썬 등의 언어에서도 CUDA로 작성된 모듈을 호출해서 사용할 수 있다. 하지만 CUDA를 통해 GPU를 직접 제어하기 위해서는 CUDA C/C++로 코드를 작성해야 한다. 

[CUDA 기반 GPU 병렬 처리 프로그래밍] p.36

여기서 많이 쓰이는 용어는 블록 스레드이다. 

4.1.1 CUDA 스레드 계층

CUDA의 스레드 계층은 스레드, 워프, 블록, 그리드라는 네 개의 계층으로 이루어져 있다. 

가. 스레드

CUDA의 계층 구조에서 가장 작은 단위는 스레드이다. CUDA에서 연산을 수행하거나 CUDA 코어를 사용하는 기본 단위이고, 작성한 커널 코드는 모든 스레드에 공유되며 각 스레드가 독립적으로 커널 코드를 수행한다. 

나. 워프

워프틑 32개의 스레드를 하나로 묶은 것을 말하며, CUDA의 기본 수행 단위이기도 하다. 한 워프에 속한 스레드들은 하나의 제어 장치에 의해 제어된다. 하나의 명령에 따라 32개의 스레드가 동시에 움직이며 이는 CUDA 프로그램에서 중요한 개념이다. 

다. 블록

워프들의 집합이다. 하나의 블록에 포함된 각 스레드는 고유한 스레드 번호를 가진다. 

라. 그리드

CUDA 계층 구조에서 가장 상위 단계는 그리드이다. 그리드는 여러 개의 블록을 포함하는 블록들의 그룹이다. 하나의 그리드에 포함된 블록들은 서로 다른 자신만의 고유한 블록 번호를 가진다. 커널이 호출되면 그리드가 생성된다. 하나의 그리드는 하나의 커널 호출되 1:1 대응되면 해당 커널을 수행할 스레드를 생성한다. 

4.2.1 스레드 레이아웃 설정 및 커널 호출

스레드 레이아웃은 스레드의 배치 형태를 지칭하는 말로서 그리드와 블록의 형태로 정의된다. 스레드 레이아웃은 커널 호출 시 설정하며 우리가 커널을 호출할 때 사용한 <<<>>> 실행 구성 문법을 통해 전달한다. 문법은 다음과 같다. 

Kernel<<<그리드의 형태, 블록의 형태>>>()

[CUDA 기반 GPU 병렬 처리 프로그래밍] 

또한 공유메모리라는 개념도 나온다. 

공유 메모리 공간을 사용하는 첫 번째 방법은 블록 내 스레드들이 고유하는 데이터를 보관하는 것이다. 이는 공유 메모리의 특성 중 블록 내 스레드들이 모두 접근 가능하다는 특성에 기반을 둔 사용 방법이다. 

class GaussianRasterizer

이제 submodules에 있는 본격 CUDA 코드이다. diff_gaussian_rasterization 폴더 내의 Init 파일은 

from . import _C

이렇게 되어 있는데, setup.py에서 

CUDAExtension(
    name="diff_gaussian_rasterization._C",  # ← 모듈 이름
    sources=[
    "cuda_rasterizer/rasterizer_impl.cu",
    "cuda_rasterizer/forward.cu",
    "cuda_rasterizer/backward.cu",
    "rasterize_points.cu",
    "ext.cpp"],

이렇게 정의하고

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
  m.def("rasterize_gaussians", &RasterizeGaussiansCUDA);
  m.def("rasterize_gaussians_backward", &RasterizeGaussiansBackwardCUDA);
  m.def("mark_visible", &markVisible);
}

빌드를 하면

pip install submodules/diff-gaussian-rasterization
    ↓
site-packages/diff_gaussian_rasterization/
    ├── __init__.py
    └── _C.cpython-3x-linux-gnu.so  ← 컴파일된 바이너리

이렇게 쓸 수 있다. 따라서 from . import _C는 diff_gaussian_rasterization내의 so 파일에서 CUDA 커널을 파이썬에서 호출할 수 있게 해주는 바이너리 모듈을 불러오는 과정이다. 

 

 

Contents

포스팅 주소를 복사했습니다

이 글이 도움이 되었다면 공감 부탁드립니다.