new_study

nms_kernel 코드실제 병렬 비교 kernel함수 정의template __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 차선 후보 개수 Nnms_overlap_thresh 차선이 비슷한지 판단하는 thresholddev_boxes GPU 메모리에 있는..

include 부분#include PyTorch extension 관련 기능PyTorch C++ Extension을 Python과 연결하기 위한 상위 헤더#include PyTorch 내부 tensor 라이브러리 ATen 사용PyTorch Tensor 계산의 핵심 C++ 라이브러리 헤더#include CUDA driver API 관련 헤더#include CUDA runtime API 관련 헤더#include std::vector 사용#include C++ 입출력상수 설명#define MAX_COL_BLOCKS 1000#define STRIDE 4#define N_OFFSETS 72#define N_STRIPS (N_OFFSETS - 1)#define PROP_SIZE (5 + N_OFFSETS)#defi..

해당 파일은 Python에서 넘어온 tensor를 확인하고, 점수순으로 정렬한 뒤 CUDA 함수로 넘김include 부분#include PyTorch C++ extension을 만들기 위한 헤더#include at::Tensor 같은 PyTorch tensor 타입을 쓰기 위한 헤더#include C++ 입출력용 헤더CUDA 함수 선언std::vector nms_cuda_forward( at::Tensor boxes, at::Tensor idx, float nms_overlap_thresh, unsigned long top_k);이 함수는 nms_kernel.cu에 실제 구현이 있음하지만 nms.cpp에서 이 함수를 호출하려면 먼저 이런 함수가 있다고 알..

코드 역할NMS 코드는 CLRNet이 예측한 192개(prior 개수)에서 서로 너무 비슷한 차선을 GPU에서 빠르게 제거하는 역할을 수행NMS는 Non-Maximum Suppression의 줄임말후보 A: 점수 0.95, x좌표들이 [100, 102, 105, ...]후보 B: 점수 0.90, x좌표들이 [101, 103, 106, ...]후보 C: 점수 0.70, x좌표들이 [300, 305, 310, ...]위 처럼 A, B, C가 있을 때 많이 겹치는 A, B에 대해 점수가 낮은 B를 제거겹치는 기준은 차선의 여러 y 위치에서 x좌표 차이가 얼마나 나는지로 측정함CUDA 프로그레밍 기초 지식커널 개념CUDA의 기본 실행 모델은 CPU가 GPU에게 일을 맡기고, GPU는 아주 많은 작은 일꾼들로 동..

목적빌드된 nms가 잘 동작되는지 확인과정import 잘 되는지 확인정상 동작 되는지 확인 1. import 태스트 코드 작성#!/usr/bin/env python3"""Check whether the official CLRNet CUDA NMS extension can be imported."""import sysfrom pathlib import PathWORKSPACE_DIR = Path(__file__).resolve().parents[2]CLRNET_DIR = WORKSPACE_DIR / "clrnet"def main() -> int: sys.path.insert(0, str(CLRNET_DIR)) try: from clrnet.ops.nms import nms ex..

목적jetson orin nano에 clrnet을 최적화하기환경OS/Kernel: Linux 5.15.185-tegra aarch64L4T: R36.5.0Python: 3.10.12Architecture: aarch64Device: Jetson Orin Nano / Orin GPUPyTorch: 2.11.0PyTorch CUDA: 12.6CUDA available: TrueGPU device: OrincuDNN: 90300NVCC: CUDA 12.6, V12.6.68TensorRT Python: 10.3.0trtexec: TensorRT v100300ONNX: 1.21.0Polygraphy: 0.50.3cuda-python: 12.6.2.post1OpenCV: 4.8.0NumPy: 1.26.4SciPy:..

ROIGatherMotivation기존 ROIAlign은 lane prior가 정의하는 차선 후보 위치를 기준으로 FPN feature map에서 lane prior feature를 추출할 수 있음ROIAlign은 lane prior 위의 Np개 point를 균일하게 샘플링하고, 각 point 위치에서 bilinear interpolation으로 feature 값을 가져오는 방식임즉, ROIAlign의 역할은 lane prior 위치에 해당하는 feature를 FPN feature map에서 추출하는 것임하지만 ROIAlign은 기본적으로 지정된 위치의 feature를 샘플링하는 연산이기 때문에, 그 자체만으로는 차선의 주변 맥락 정보를 충분히 학습하지 못함차선은 차량, 그림자, 강한 빛, 마모 등으로 ..

문제 의식차선은 high-level semantics를 가진 객체임왜냐하면 차선을 인식하려면 도로 상황과 구조를 이해해야 하기 때문차선은 픽셀 수준에서는 단순한 흰색/노란색이지만 해당 픽셀이 횡단보도인지 반사광인지, 노면 균열인지 구분하기 어려우며, 차선은 차량, 그림자 등으로 일부가 가려지는 경우가 많기 때문에 도로 구조와 주변 문맥을 바탕으로 이어지는 위치를 추론해야 함또한 차선은 semantic 정보와 함께 정확한 위치 localization이 중요함즉 차선 인식은 high-level feature와 low-level feature를 함께 활용해야 하지만 해당 연구가 미비함핵심 아이디어CLRNet은 high-level feature와 low-level feature를 모두 활용하여 차선 인식의 정확..

CLRNet 모델 내부 흐름1. Model input [1, 3, 320, 800]2. DLA34 Backbone 입력 이미지에서 multi-scale feature map 추출 output 0: [1, 64, 80, 200] output 1: [1, 128, 40, 100] output 2: [1, 256, 20, 50] output 3: [1, 512, 10, 25]3. FPN Neck DLA34 feature 중 뒤의 3개 scale 사용 채널을 모두 64로 맞추고 top-down fusion 수행 input: [1, 128, 40, 100] [1, 256, 20, 50] [1, 512, 10, 25] output..

prior 의미: 차선이 있을 법한 선 후보 핵심 파라미터는 시작점(start_y, start_x)과 각도(theta) 여기서 72개 고정 y 위치에 대한 x좌표를 계산할 수 있음학습 가능한 prior 값: start_y start_x thetaprior 초기화: 192개의 초기 차선 후보선을 왼쪽/아래/오른쪽 boundary에 분산 배치prior의 목적: 1. feature map에서 어디를 sampling할지 결정 2. 최종 lane prediction의 기준선 역할학습 후 prior의 의미: CULane 데이터에서 자주 유용한 차선 후보선(anchor lane/template)학습 가능한 이유: prior가 sampling 위치와 최종 차선 위치 계산에 직접 들어감 sampli..