# \#gpu-kernel **URL:** https://discuss.pytorch.kr/tag/gpu-kernel/6201.md [최신](https://discuss.pytorch.kr/latest.md) · [카테고리](https://discuss.pytorch.kr/categories.md) · [태그](https://discuss.pytorch.kr/tags.md) --- ## [tang: 연산자별 커널 없이 FlashAttention 순전파와 역전파를 유도하는 Rust 텐서 컴파일러](https://discuss.pytorch.kr/t/tang-flashattention-rust/12502)
**Author:** [@9bow](https://discuss.pytorch.kr/u/9bow)\ **댓글:** 0\ **마지막 업데이트:** [11월 13, 2026, 11:00오전 UTC](https://discuss.pytorch.kr/t/tang-flashattention-rust/12502 "2026-11-13T11:00:29Z")
핵심 요약 휴머노이드 Rose 등 오픈소스 로봇을 만드는 Municipal Robotics의 Cam Pedersen이 연구 보고서 "Don't Write Kernels, Derive Them"을 공개했습니다. Rust 라이브러리 tang의 텐서 IR(Intermediate Representation, 중간 표현)이 뷰와 모노이드 컴프리헨션 두 가지만으로 FlashAttention의 순전파와 역전파 커널을 손으로 쓰지 않고 유도한다는 내용입니다. 같은 그래프가 CUDA, Metal, WGSL, OpenCL, GLSL, C로 내려갑니다. 유휴 RTX 3090에서 트랜스포머 블록 순전파는 시퀀스 길이 512에서 CUDA 그래프로 캡처한 PyTorch보다 1.2배 빠르고 128에서는 1.2배 느립니다. 지는 구간도 저자가 함께 공개했습니다. 헤드 차원 64 순전파는 FlashAttention-2에 모든 크기에서 지고, H200의 헤드 차원 256 순전파는 FlashAttention-3 속도의 0.43~0.54배입니다. 80B MoE 엔진의 프리필은 vLLM보다 약 12배 느립니다. 재작성 규칙의 건전성과 트래픽 하한의 약한 쌍대성은 Lean 4로 증명했지만(실수 위에서, 부동소수점 제외), 이 하한으로 만든 시간 하한은 H200에서 실측 시간의 0.05~0.66배에 그쳐 아직 느슨합니다. arXiv나 학회 논문집이 아니라 Municipal Robotics 웹사이트에 직접 공개한 보고서이며, 수치는 모두 저자의 측정입니다. 코드와 원시 측정 기록은 저장소에 공개돼 있습니다. tang 소개 tang 은 모델을 수식으로 한 번만 적으면, 그 수식에서 GPU 커널을 컴파일러가 유도(derive)하도록 만든 Rust 텐서 컴파일러입니다. 손으로 작성한 연산자별 커널 라이브러리를 작은 대수 체계와 스케줄 탐색으로 대체할 수 있는지를 실측으로 검증한 것이 이번 보고서의 내용입니다. 이 질문이 나온 배경은 커널의 수입니다. 현대 모델의 서로 다른 연산자는 수십 개 수준이… --- ## [Jane Street 및 NVIDIA, Mamba-2 순전파를 Triton보다 최대 1.82배 빠르게 구현한 cuTile 커널 공개](https://discuss.pytorch.kr/t/jane-street-nvidia-mamba-2-triton-1-82-cutile/12261)
**Author:** [@9bow](https://discuss.pytorch.kr/u/9bow)\ **댓글:** 0\ **마지막 업데이트:** [10월 26, 2026, 11:00오전 UTC](https://discuss.pytorch.kr/t/jane-street-nvidia-mamba-2-triton-1-82-cutile/12261 "2026-10-26T11:00:32Z")
핵심 요약 Jane Street와 NVIDIA 엔지니어들이 NVIDIA 기술 블로그 글에서 Mamba-2 순전파를 cuTile Python 커널 3개(chunk\_cumsum, fwd\_pass\_state, compute\_out)로 구현한 과정을 수식과 코드로 설명했습니다. NVIDIA B300 GPU 1장에서 mamba\_ssm의 Triton 커널과 비교했을 때 1.14~1.82배 빨랐고, 상태 크기가 클수록 격차가 컸습니다. 측정 범위는 커널 단위 순전파뿐이며, 역전파와 학습 전체 속도는 측정하지 않았습니다. 구현은 TileGym 저장소의 tilegym.ops.mamba2\_chunk\_forward로 공개되어 있으며, 블로그에서 다루지 않은 역전파 커널 mamba2\_chunk\_backward도 함께 들어 있습니다. 블로그의 벤치마크 예제는 use\_cupti=True 인자를 쓰는데, 이 인자는 2026-10-09 TileGym 커밋에서 빠졌습니다. 현재 main에서는 이 인자를 지워야 합니다. cuTile로 다시 쓴 Mamba-2 순전파 커널 소개 Mamba-2는 상태 공간 모델(State Space Model, SSM) 계열의 시퀀스 모델입니다. 2026년 10월 19일 NVIDIA 기술 블로그에 올라온 이 글은 Mamba-2의 순전파(Forward Pass)를 NVIDIA CUDA Tile(cuTile)로 구현하고 최적화한 과정을 설명하는 튜토리얼입니다. 수식 유도부터 커널 코드, Blackwell GPU에 맞춘 최적화, Triton 구현과의 성능 비교까지 한 번에 다룹니다. GPU 커널을 직접 작성하거나 Mamba 계열 모델의 학습과 프리필(Prefill) 속도를 다루는 분들께 필요한 내용입니다. cuTile은 NVIDIA가 만든 Python 기반 GPU 프로그래밍 모델입니다. 스레드와 워프 단위 스케줄링을 직접 다루는 대신 타일(Tile, 텐서의 작은 블록) 단위로 연산을 적으면, 컴파일러가 텐서 코어(Tensor Core) 같은 하드웨어 기능을 … --- ## [Modular의 Mojo와 MAX: CUDA 종속 없이 여러 AI 하드웨어를 같은 코드로 다루려는 언어와 추론 프레임워크](https://discuss.pytorch.kr/t/modular-mojo-max-cuda-ai/12266)
**Author:** [@9bow](https://discuss.pytorch.kr/u/9bow)\ **댓글:** 0\ **마지막 업데이트:** [10월 23, 2026, 11:26오전 UTC](https://discuss.pytorch.kr/t/modular-mojo-max-cuda-ai/12266 "2026-10-23T11:26:49Z")
핵심 요약 Modular는 2022년 초 Chris Lattner와 Tim Davis가 세운 회사입니다. Democratizing AI Compute 연재에서 밝힌 목표는 CUDA에 묶인 AI 소프트웨어 스택을 여러 하드웨어가 함께 쓰는 공용 계층으로 다시 만드는 것입니다. 그 결과물이 시스템 프로그래밍 언어 Mojo 와 추론(Inference) 프레임워크 MAX 이며, 두 프로젝트의 코드는 modular/modular GitHub 저장소에 함께 공개되어 있습니다. Mojo는 2026년 8월 11일 1.0에 도달했고, 같은 달 18일에 컴파일러까지 Apache 2.0(LLVM 예외 포함)으로 공개되었습니다. 10월 22일 나온 1.2에는 패턴 매칭이 추가되었습니다. 현지 시각 10월 22일에는 MAX 26.7과 함께 그래프 컴파일러, 런타임, 하드웨어 백엔드까지 소스 코드 전체를 공개하였습니다. 다만 이 핵심부와 배포 바이너리에는 Apache 2.0이 아닌 MAX License가 적용되어, 수정한 소스는 기여 목적의 포크에서만 공개할 수 있는 소스 공개(Source-Available) 방식입니다. 2026년 7월 Qualcomm이 Modular 인수를 마쳤고, Modular는 Mojo와 MAX를 계속 유지하며 Qualcomm과 경쟁하는 하드웨어도 지원하겠다고 밝혔습니다. Modular 프로젝트 소개 Modular는 AI 모델을 여러 종류의 하드웨어에서 같은 코드로 실행할 수 있도록, 프로그래밍 언어부터 모델 서빙까지의 소프트웨어 스택을 처음부터 다시 만들고 있습니다. 지금의 AI 소프트웨어는 대부분 NVIDIA의 CUDA를 전제로 작성되어 있습니다. 그래서 AMD GPU나 Apple 실리콘, 클라우드 업체의 자체 가속기로 옮기려면 커널과 서빙 코드를 상당 부분 다시 작성해야 합니다. Modular는 하드웨어마다 소프트웨어 스택을 따로 만드는 대신, 모든 하드웨어가 공유하는 하나의 계층으로 이 문제를 풀겠다는 입장입니다. 이번에 소개할 Modular의… --- ## [\[GN⁺\] CUDA 커널을 실행하면 내부에서 벌어지는 일](https://discuss.pytorch.kr/t/gn-cuda/11006)
**Author:** [@9bow](https://discuss.pytorch.kr/u/9bow)\ **댓글:** 0\ **마지막 업데이트:** [7월 1, 2026, 9:30오전 UTC](https://discuss.pytorch.kr/t/gn-cuda/11006 "2026-07-01T09:30:29Z")
CUDA 커널을 실행하면 내부에서 벌어지는 일 글 소개 단순한 벡터 덧셈 CUDA 프로그램도 결과 2.000000을 얻기까지 컴파일 파이프라인, 드라이버 호출, GPU 명령 큐, 워프 스케줄링, 메모리 계층, 완료 세마포어를 거침 nvcc는 호스트 코드와 디바이스 코드를 나눠 cicc로 PTX, ptxas로 SASS를 만들고, cubin과 PTX를 fatbin에 묶어 Linux 실행 파일 안에 넣음 vadd\<\<\<4096, 256\>\>\> launch 구문은 호스트 launch stub으로 바뀌며, 인자 da, db, dc, n은 CUDA 런타임과 libcuda.so.1을 거쳐 드라이버에 전달됨 GPU 실행은 QMD, pushbuffer, GPFIFO, GP\_PUT, doorbell MMIO 쓰기로 시작되고, RTX 4090의 128개 SM이 4096개 블록과 256개 스레드 구성을 워프 단위로 실행함 이 커널은 float 덧셈 1회당 12바이트 전송이 필요한 낮은 산술 집약도 때문에 Nsight Compute에서 10.78μs, DRAM 피크의 79.65%, warp issue 5.17%로 메모리 대역폭에 좌우됨 예제 커널과 관찰 범위 예제 프로그램은 vadd CUDA 커널로 두 float 배열을 더해 세 번째 배열에 저장함 n = 1 \<\< 20으로 1,048,576개 float를 처리함 launch 구성은 vadd\<\<\<4096, 256\>\>\>(da, db, dc, n)이며 4096 \* 256 = n개 스레드를 사용함 RTX 4090 대상으로 nvcc -arch=sm\_89로 컴파일해 실행하면 c\[0\]=2.000000 c\[n-1\]=2.000000이 출력됨 이 한 줄의 결과에도 CPU 명령 수천만 개, device file, 약 900개의 ioctl, 메모리 매핑된 doorbell 레지스터가 관여함 nvcc가 실행 파일을 만드는 과정 nvcc --keep를 사용하면 컴파일 파이프라인 산출물을 직접 확인할 수 있음 vadd.ptx: cicc가 만든… --- ## [TileKernels: DeepSeek AI가 공개한, LLM 연산을 위한 고성능 TileLang 기반 GPU 커널 라이브러리](https://discuss.pytorch.kr/t/tilekernels-deepseek-ai-llm-tilelang-gpu/9878)
**Author:** [@9bow](https://discuss.pytorch.kr/u/9bow)\ **댓글:** 0\ **마지막 업데이트:** [6월 24, 2026, 11:30오전 UTC](https://discuss.pytorch.kr/t/tilekernels-deepseek-ai-llm-tilelang-gpu/9878 "2026-06-24T11:30:20Z")
TileKernels 소개 현대 대규모 언어 모델(LLM)의 학습과 추론에서 GPU 커널 최적화는 연산 효율을 좌우하는 핵심 요소입니다. 특히 Mixture of Experts(MoE) 아키텍처가 널리 쓰이면서 게이팅(Gating), 전문가 라우팅(Expert Routing), 양자화(Quantization) 등 다양한 연산 커널의 하드웨어 수준 최적화 필요성이 높아졌습니다. 기존에는 커스텀 CUDA 커널을 작성하기 위해 C++와 CUDA 내부를 깊이 이해해야 했기 때문에 진입 장벽이 높았고, 새로운 아키텍처가 등장할 때마다 커널을 처음부터 다시 작성해야 하는 문제가 있었습니다. TileKernels는 DeepSeek AI가 공개한 오픈소스 GPU 커널 라이브러리로, LLM 연산에 특화된 고성능 커널 모음입니다. 이 프로젝트는 TileLang을 기반으로 구축되었습니다. TileLang은 Python으로 고성능 GPU 커널을 표현할 수 있도록 설계된 도메인 특화 언어(Domain-Specific Language, DSL)로, 쉬운 마이그레이션, 빠른 개발, 자동 최적화를 특징으로 합니다. 기존 CUDA C++ 대신 Python 수준에서 커널을 작성할 수 있어 개발 생산성이 크게 향상됩니다. 프로젝트의 커널 대부분은 연산 집약도(Compute Intensity)와 메모리 대역폭(Memory Bandwidth) 측면에서 하드웨어 성능 한계에 근접하도록 설계되었습니다. 일부 커널은 이미 DeepSeek 내부의 학습 및 추론 시나리오에서 사용되고 있으며, MoE 기반 모델의 핵심 연산들을 담고 있습니다. 다만 개발팀은 현재 코드가 베스트 프랙티스를 대표하지는 않으며, 코드 품질과 문서화를 지속적으로 개선하는 중이라고 밝히고 있습니다. TileKernels의 주요 커널 구성 TileKernels는 LLM 연산에 필요한 다양한 커널을 체계적으로 제공합니다. 각 커널 모음은 저장소의 tile\_kernels/ 디렉토리 아래 영역별로 정리되어 있습니다. 게이팅(Ga… --- ## [FlashKDA: MoonshotAI가 공개한, Kimi Delta Attention을 위한 고성능 CUDA 커널 라이브러리](https://discuss.pytorch.kr/t/flashkda-moonshotai-kimi-delta-attention-cuda/9881)
**Author:** [@9bow](https://discuss.pytorch.kr/u/9bow)\ **댓글:** 0\ **마지막 업데이트:** [4월 26, 2026, 3:30오전 UTC](https://discuss.pytorch.kr/t/flashkda-moonshotai-kimi-delta-attention-cuda/9881 "2026-04-26T03:30:05Z")
FlashKDA 소개 대규모 언어 모델의 추론 효율을 높이기 위한 선형 어텐션(Linear Attention) 연구가 빠르게 발전하고 있습니다. Transformer의 표준 소프트맥스 어텐션은 시퀀스 길이 T 에 대해 O(T^2) 의 연산 복잡도를 가지기 때문에, 긴 컨텍스트 처리 시 메모리와 연산 비용이 크게 증가합니다. 이를 해결하기 위해 Mamba, RWKV, RetNet 등 선형 복잡도 아키텍처들이 등장했으며, 최근에는 이런 모델들의 핵심 연산을 CUDA 수준에서 하드웨어에 최적화하는 작업이 중요해졌습니다. 특히 Triton으로 작성된 선형 어텐션 커널은 구현이 유연하지만, CUTLASS 기반의 고도로 최적화된 CUDA 커널에 비해 성능 차이가 존재합니다. FlashKDA(Flash Kimi Delta Attention)는 MoonshotAI(Kimi 개발사)가 공개한, Kimi 모델에서 사용하는 Delta Attention(KDA, Kimi Delta Attention) 연산을 위한 고성능 CUDA 커널 라이브러리입니다. NVIDIA CUTLASS를 기반으로 구축되었으며, Hopper(H100) 및 그 이상의 GPU 아키텍처(SM90+)를 대상으로 최적화되었습니다. Flash Attention이 소프트맥스 어텐션을 하드웨어 수준에서 최적화한 것처럼, FlashKDA는 KDA 연산을 동일한 방식으로 가속화합니다. 2026년 4월 22일에는 FlashKDA v1의 설계 결정 사항을 상세히 설명하는 딥다이브 블로그 포스트도 공개되었습니다. FlashKDA는 flash-linear-attention 라이브러리와 통합되어, chunk\_kda 함수 호출 시 자동으로 FlashKDA 커널로 디스패치됩니다. 기존 Triton 구현 경로를 대체하여 H100 등 최신 GPU에서 더 높은 성능을 제공합니다. FLA\_FLASH\_KDA=0 환경 변수를 설정하면 언제든지 Triton 경로로 되돌릴 수 있어, 디버깅이나 하위 호환성이 필요한 경우에 유용합니다. Fla… --- ## [cuLA(CUDA Linear Attention): NVIDIA Blackwell, Hopper GPU를 위한 선형 어텐션 고성능 CUDA 커널 라이브러리](https://discuss.pytorch.kr/t/cula-cuda-linear-attention-nvidia-blackwell-hopper-gpu-cuda/9658)
**Author:** [@9bow](https://discuss.pytorch.kr/u/9bow)\ **댓글:** 0\ **마지막 업데이트:** [4월 12, 2026, 9:30오전 UTC](https://discuss.pytorch.kr/t/cula-cuda-linear-attention-nvidia-blackwell-hopper-gpu-cuda/9658 "2026-04-12T09:30:30Z")
cuLA 소개 트랜스포머(Transformer) 기반 언어 모델이 점점 더 긴 컨텍스트를 다루게 되면서, 기존의 소프트맥스 어텐션(Softmax Attention)이 가진 O(n^2) 계산 복잡도 문제는 핵심적인 병목이 되었습니다. 이를 해결하기 위해 등장한 선형 어텐션(Linear Attention)은 상태 업데이트를 선형 시간으로 수행함으로써 긴 컨텍스트 처리에 훨씬 유리한 구조를 제공합니다. 그러나 이론적인 효율성을 실제 하드웨어 성능으로 연결하려면, GPU의 특성에 맞게 세심하게 최적화된 커널 구현이 반드시 필요합니다. cuLA는 선형 어텐션 변형들을 위한 고성능 CUDA 커널 라이브러리입니다. inclusionAI가 개발한 이 프로젝트는 CuTe DSL과 CUTLASS C++를 활용하여 NVIDIA Blackwell(SM10X) 및 Hopper(SM90) GPU 아키텍처에 특화된 핸드튜닝(hand-tuned) 구현을 제공합니다. cuLA는 flash-linear-attention(FLA)의 서브모듈로 설계되어 있으며, 기존 FLA 사용자라면 단 한 줄의 임포트 변경만으로 cuLA의 가속화된 커널을 활용할 수 있습니다. 현재 v0.1.0(2026년 4월 2일 출시)으로 초기 개발 단계에 있으며, 커뮤니티의 기여를 적극적으로 환영하고 있습니다. cuLA가 지원하는 선형 어텐션 알고리즘 cuLA는 최근 연구에서 제안된 여러 선형 어텐션 변형들을 구현합니다. 각 알고리즘은 표준 어텐션의 한계를 다양한 방식으로 보완하며, cuLA는 이들을 Blackwell 및 Hopper GPU 위에서 최대한 효율적으로 동작하도록 최적화합니다. KDA(Kimi Delta Attention) 는 Moonshot AI의 Kimi 모델에서 사용되는 어텐션 메커니즘으로, 델타 업데이트 방식으로 선형 상태를 갱신합니다. cuLA는 모듈식 KDA 포워드(Modular KDA Forward)와 퓨전 KDA 포워드(Fused KDA Forward) 두 가지 변형을 지원합니다. 관… --- ## [AutoKernel: AI 에이전트가 GPU 커널을 자동으로 최적화하는 프로젝트 (= GPU 커널용 Auto Research)](https://discuss.pytorch.kr/t/autokernel-ai-gpu-gpu-auto-research/9200)
**Author:** [@9bow](https://discuss.pytorch.kr/u/9bow)\ **댓글:** 0\ **마지막 업데이트:** [3월 12, 2026, 9:30오후 UTC](https://discuss.pytorch.kr/t/autokernel-ai-gpu-gpu-auto-research/9200 "2026-03-12T21:30:06Z")
AutoKernel 소개 현대의 딥러닝 모델, 특히 대규모 언어 모델(LLM)을 학습하거나 추론할 때 GPU 연산의 병목 현상을 해결하는 것은 매우 중요합니다. 이를 위해 CUDA나 Triton과 같은 저수준 언어로 하드웨어에 맞게 GPU 커널을 최적화하는 작업이 필수적이지만, 이는 고도의 전문 지식과 끝없는 디버깅 시간을 요구하는 까다로운 영역입니다. 이러한 문제를 해결하기 위해 등장한 오픈소스 프레임워크가 바로 AutoKernel입니다. 이 프로젝트는 유명 AI 연구자 안드레이 카파시(Andrej Karpathy)가 선보였던 LLM 학습을 위한 자율 AI 연구 에이전트인 Auto Research 프로젝트에서 직접적인 영감을 받아 탄생했습니다. 일반적으로 GPU 커널 최적화는 사람이 직접 프로파일러를 보며 코드를 한 줄 한 줄 수정하고 컴파일하여 성능을 테스트하는 지루한 반복 작업(Manual Iteration)을 거칩니다. 반면, AutoKernel은 Auto Research의 자율 연구 철학을 GPU 커널 최적화 영역으로 가져왔습니다. 개발자가 PyTorch 모델을 제공하고 코딩 에이전트(Claude, Codex 등)를 실행해두면, 에이전트가 단일 커널 파일을 수정하고 벤치마크를 실행한 뒤 결과를 평가하여 코드를 유지할지 되돌릴지 끝없이 반복합니다. 즉, 개발자가 잠든 사이 AI가 병목 지점을 찾아내고 스스로 코드를 깎아내어 최적화된 Triton 커널을 완성해 내는 혁신적인 도구입니다. 기존에는 수작업으로 인해 병목이 큰 몇 개의 커널에만 최적화 시도가 국한되었다면, AutoKernel은 1회 실험에 약 90초만을 소모하며, 밤새 320회 이상의 실험을 모든 병목 커널에 걸쳐 지치지 않고 자동으로 수행할 수 있다는 점에서 압도적인 생산성 차이를 보여줍니다. AutoKernel의 구조 및 동작 파이프라인 AutoKernel은 여러 단계의 스크립트를 통해 모델 분석부터 최종 검증까지의 전체 과정을 매끄럽게 연결합니다: prof… --- ## [CUDA Tile: NVIDIA가 제시하는 차세대 타일(Tile) 기반 GPU 프로그래밍 (feat. cuTile Python, Tile IR, TileGym)](https://discuss.pytorch.kr/t/cuda-tile-nvidia-tile-gpu-feat-cutile-python-tile-ir-tilegym/8535)
**Author:** [@9bow](https://discuss.pytorch.kr/u/9bow)\ **댓글:** 0\ **마지막 업데이트:** [12월 28, 2025, 3:30오전 UTC](https://discuss.pytorch.kr/t/cuda-tile-nvidia-tile-gpu-feat-cutile-python-tile-ir-tilegym/8535 "2025-12-28T03:30:16Z")
CUDA Tile 소개 CUDA Tile은 NVIDIA가 최근 CUDA 13.1과 함께 발표한 차세대 GPU 프로그래밍 모델로, 인공지능(AI)과 고성능 컴퓨팅(HPC) 분야에서 필수적인 고성능 커널 개발의 패러다임을 근본적으로 변화시키기 위해 등장했습니다. 이 기술은 단순히 새로운 라이브러리를 추가하는 차원이 아니라, GPU 연산의 기본 단위를 스레드(Thread)에서 데이터 덩어리인 타일(Tile)로 격상시키는 거대한 아키텍처의 전환을 의미합니다. 최근의 GPU 하드웨어는 비약적으로 발전했습니다. NVIDIA의 Hopper나 Blackwell 아키텍처는 텐서 코어(Tensor Core) 와 같은 강력한 전용 가속 유닛을 탑재하고 있으며, 메모리 계층 구조 또한 더욱 깊고 복잡해졌습니다. 하지만 기존의 프로그래밍 방식으로는 이러한 하드웨어의 잠재력을 100% 끌어내기 위해 개발자가 극도로 복잡한 최적화 과정을 수동으로 수행해야만 했습니다. CUDA Tile은 이러한 문제를 해결하기 위해, 하드웨어의 복잡성을 컴파일러와 중간 표현(IR, Intermediate Representation) 레벨에서 추상화(Abstraction)하고, 개발자는 데이터의 흐름과 레이아웃에 집중할 수 있도록 돕는 포괄적인 생태계입니다. 기존 SIMT 모델의 한계와 CUDA Tile의 필요성 전통적인 CUDA 프로그래밍은 SIMT(Single Instruction, Multiple Threads) 모델을 기반으로 합니다. 이 모델에서 개발자는 수천, 수만 개의 스레드를 개별적으로 제어해야 합니다. 예를 들어, 행렬 곱셈을 수행할 때 개발자는 각 스레드가 처리할 데이터의 인덱스를 계산하고(threadIdx, blockIdx), 글로벌 메모리에서 데이터를 가져와 공유 메모리(Shared Memory)에 적재한 뒤, 레지스터로 이동시켜 연산을 수행하는 파이프라인 전체를 직접 코딩해야 했습니다. 이는 하드웨어를 극한으로 제어할 수 있다는 장점이 있지만, 코드가 복잡해지고 유지보…