Booil Jung

CUDA

이 파트에서는 CUDA의 근본적인 배경을 확립하고, CUDA가 왜 만들어졌으며 그것을 정의하는 핵심 아키텍처 및 실행 원칙이 무엇인지 설명합니다. CPU 스케일링의 한계라는 거시적 산업 동향에서부터 GPU의 미시적 아키텍처 세부 사항까지 다룹니다.

컴퓨팅 기술의 발전은 오랫동안 중앙 처리 장치(CPU)의 성능 향상과 동의어였습니다. 그러나 21세기에 들어서면서 CPU 성능을 향상시키는 전통적인 방법은 물리적 한계에 부딪히기 시작했습니다. 반도체 공정의 미세화는 크기, 발열, 전력 소비와 같은 근본적인 문제에 직면했으며, 이로 인해 클럭 속도를 무한정 높이는 것은 불가능해졌습니다. 이러한 상황에서 업계는 성능 향상을 위한 새로운 돌파구를 모색하기 시작했습니다.

CPU는 본질적으로 소수의 강력한 코어를 사용하여 복잡한 작업을 순차적으로 처리하도록 설계되었습니다. 이는 운영체제 관리나 단일 스레드 애플리케이션 실행과 같은 작업에는 매우 효율적이지만, 대규모 데이터를 동일한 연산으로 처리해야 하는 문제에는 적합하지 않았습니다. 바로 이 지점에서 그래픽 처리 장치(GPU)가 대안으로 부상했습니다. GPU는 원래 3D 그래픽 렌더링을 위해 수많은 데이터를 병렬로 처리하도록 설계된 하드웨어였습니다. 이러한 구조적 특성은 그래픽 작업뿐만 아니라, 대규모 병렬 연산이 필요한 일반적인 컴퓨팅 문제에도 적용될 수 있는 잠재력을 가지고 있었습니다.

이러한 배경에서 GPGPU(General-Purpose computing on Graphics Processing Units), 즉 GPU를 이용한 범용 컴퓨팅이라는 개념이 탄생했습니다. GPGPU의 핵심 아이디어는 애플리케이션의 연산 집약적인 부분을 CPU에서 GPU로 오프로드하여, GPU의 수많은 코어를 활용해 병렬로 처리함으로써 전체 시스템의 성능을 극적으로 향상시키는 것입니다.

GPGPU의 역사는 2000년대 초반, 개발자들이 그래픽 렌더링 파이프라인을 비그래픽 연산에 활용하려는 시도에서 시작되었습니다. 이 흐름에 결정적인 전환점을 마련한 것은 2001년 NVIDIA가 출시한 GeForce3였습니다. 이 GPU는 프로그래밍 가능한 셰이더(Programmable Shader)를 탑재한 최초의 제품으로, 개발자가 고정된 기능 대신 자신만의 코드를 GPU에서 실행할 수 있게 함으로써 고정 함수 파이프라인의 시대를 끝내고 프로그래밍 가능한 그래픽의 시대를 열었습니다.

이러한 하드웨어의 발전은 학계의 연구를 촉진했습니다. 2000년대 초, 스탠포드 대학교에서는 ‘브룩(Brook)’이라는 프로젝트가 진행되었는데, 이는 C와 같은 고수준 언어를 확장하여 GPU에서 범용 계산을 수행하는 프로그래밍 모델을 제시했습니다. 브룩은 CPU가 GPU로 계산 작업을 오프로드하는 핵심 개념을 정립하며 CUDA의 중요한 전신이 되었습니다. NVIDIA는 이러한 학술적 움직임의 잠재력을 일찌감치 파악하고 브룩 프로젝트에 자금을 지원하며 기술의 상업적 가능성을 탐색했습니다. 이 프로젝트를 이끌었던 이안 벅(Ian Buck)과 같은 핵심 인재들은 이후 NVIDIA에 합류하여 CUDA 개발을 주도하게 됩니다.

NVIDIA의 리더십, 특히 젠슨 황(Jensen Huang) CEO는 GPU를 단순히 게임이나 그래픽을 위한 장치가 아닌, 과학 컴퓨팅을 위한 범용 하드웨어로 만들겠다는 명확한 비전을 가지고 있었습니다. 이러한 비전 아래, NVIDIA는 2006년에 CUDA(Compute Unified Device Architecture)를 공식적으로 발표하고 2007년에 첫 버전을 출시했습니다.

CUDA의 핵심 개발 동기는 ‘접근성’이었습니다. 당시 GPGPU를 구현하기 위해서는 Direct3D나 OpenGL과 같은 그래픽 API에 대한 깊고 전문적인 지식이 필요했습니다. 이는 수많은 과학 및 공학 분야의 개발자들에게 높은 진입 장벽으로 작용했습니다. NVIDIA는 이러한 문제를 해결하기 위해, 당시 과학 및 고성능 컴퓨팅 분야에서 가장 널리 사용되던 C 언어를 기반으로 GPU 프로그래밍을 가능하게 하는 것을 목표로 삼았습니다. 개발자들이 완전히 새로운 언어를 배울 필요 없이, 익숙한 C/C++ 구문에 몇 가지 확장 키워드만 추가하여 GPU의 병렬 처리 능력을 활용할 수 있도록 한 것입니다.

이러한 접근 방식은 단순한 기술적 선택을 넘어선 전략적 결정이었습니다. NVIDIA는 CUDA를 통해 GPU를 위한 새로운 시장을 창출하고자 했습니다. C/C++ 개발자 커뮤니티라는 거대한 잠재 고객층에게 가장 낮은 진입 장벽을 제공함으로써, NVIDIA는 도구를 만드는 것을 넘어 하나의 생태계와 그 안의 충성도 높은 사용자 기반을 구축하기 시작했습니다. 이는 CUDA가 단순한 API가 아니라, 개발자들을 NVIDIA 하드웨어에 묶어두는 강력한 ‘해자(moat)’가 되는 첫걸음이었습니다.

CUDA는 단순히 API의 집합이 아니라, 하드웨어와 소프트웨어를 아우르는 포괄적인 병렬 컴퓨팅 플랫폼입니다. 그 핵심 철학은 개발자에게 GPU의 가상 명령어 집합(Virtual Instruction Set)과 병렬 연산 요소에 직접 접근할 수 있는 소프트웨어 계층을 제공하는 것입니다. 이를 통해 개발자는 하드웨어의 복잡한 저수준 제어에 얽매이지 않으면서도 GPU의 성능을 최대한으로 끌어낼 수 있습니다.

이 플랫폼은 여러 구성 요소로 이루어져 있습니다.

결론적으로, CUDA는 개발자 경험을 최우선으로 고려한 전략적 산물입니다. NVIDIA는 초기부터 개방형 표준보다는 사용 편의성에 집중함으로써, CUDA를 기술적 우위뿐만 아니라 강력한 생태계를 갖춘 플랫폼으로 성장시켰습니다. 이는 2012년 AlexNet이 맞춤형 CUDA 커널을 사용하여 딥러닝 혁명을 촉발시킨 사례에서 명확히 드러납니다. 연구자들과 개발자들이 아이디어를 신속하게 구현하고 성능을 극대화할 수 있는 가장 효율적인 경로를 제공함으로써, CUDA는 사실상의 산업 표준으로 자리 잡았고, 경쟁사들이 따라잡기 어려운 깊은 기술적, 생태계적 해자를 구축했습니다. 이러한 초기 전략의 성공이 오늘날 NVIDIA가 AI 및 HPC 시장에서 누리는 지배적인 위치의 근간을 이루고 있습니다.

CUDA의 강력함은 단순히 많은 코어를 병렬로 사용하는 것을 넘어, 이들을 효율적으로 관리하고 제어하는 정교한 실행 모델에 있습니다. 이 모델은 GPU의 물리적 하드웨어 구조를 논리적으로 추상화하여 개발자가 수천, 수만 개의 스레드를 체계적으로 다룰 수 있게 합니다.

GPU의 물리적 아키텍처를 이해하는 것은 CUDA 실행 모델을 파악하는 첫걸음입니다. 현대 GPU는 다수의 스트리밍 멀티프로세서(Streaming Multiprocessors, SM) 로 구성됩니다. 각 SM은 연산을 실제로 수행하는 여러 개의 스트리밍 프로세서(Streaming Processors, SP), 즉 CUDA 코어(CUDA Cores) 를 포함하고 있습니다. 예를 들어, NVIDIA의 Turing 아키텍처 기반 GPU는 72개의 SM을 가질 수 있으며, 각 SM은 64개의 SP를 가집니다.

이러한 구조는 GPU가 대규모 병렬 처리에 적합한 이유를 설명해 줍니다. 수십 개의 SM이 독립적으로 작업을 처리하고, 각 SM 내의 수십 개 CUDA 코어가 동시에 연산을 수행함으로써, 전체 GPU는 수천 개의 스레드를 동시에 실행할 수 있는 능력을 갖추게 됩니다.1

GPU가 수많은 스레드를 효율적으로 제어하는 핵심 원리는 SIMT(Single Instruction, Multiple Thread) 실행 모델입니다.2 SIMT 모델에서 제어 장치는 단일 명령어를 여러 스레드에 동시에 발행합니다. 그러면 각 스레드는 해당 명령어를 각자 할당된 데이터에 대해 병렬로 실행합니다. 이는 대규모 데이터 블록에 동일한 연산을 반복적으로 적용하는 알고리즘(예: 벡터 덧셈, 행렬 곱셈)에서 엄청난 효율성을 발휘합니다.

CUDA에서 SIMT 모델이 실제로 구현되는 단위는 워프(Warp) 입니다. 워프는 32개의 스레드로 구성된 그룹으로, SM에서 스케줄링되는 기본 단위입니다.1 SM의 명령어 유닛은 워프 단위로 명령어를 발행하며, 한 워프에 속한 32개의 스레드는 항상 동일한 명령어를 동시에 실행하는, 이른바 ‘락스텝(lock-step)’ 방식으로 동작합니다.

이 워프라는 개념은 하드웨어 효율성을 위한 중요한 타협점입니다. 모든 개별 스레드에 대해 별도의 명령어 인출 및 해독 장치를 두는 것은 하드웨어적으로 매우 복잡하고 전력 소모가 큽니다. 대신, 32개의 스레드를 하나의 그룹으로 묶어 단일 제어 유닛으로 관리함으로써 하드웨어의 복잡성을 크게 줄이면서도 높은 수준의 병렬성을 유지할 수 있습니다. 하지만 이러한 하드웨어적 제약은 소프트웨어 성능에 직접적인 영향을 미칩니다. 특히, 워프 내 스레드들이 서로 다른 코드 경로를 따르는 제어 흐름 분기(control flow divergence) 가 발생할 때 성능 저하의 주요 원인이 됩니다. 예를 들어, if-else 구문에서 워프 내 일부 스레드는 if 블록을, 나머지는 else 블록을 실행해야 할 경우, 하드웨어는 두 경로를 모두 순차적으로 실행해야 합니다. if 경로가 실행되는 동안 else 경로를 택한 스레드들은 비활성화(idle) 상태가 되고, 그 반대도 마찬가지입니다. 따라서 고성능 CUDA 코드를 작성하기 위해서는 개발자가 항상 ‘워프 단위로 생각’하고, 워프 내 스레드 간의 분기를 최소화하도록 알고리즘을 설계해야 합니다.

CUDA는 개발자가 수많은 스레드를 효과적으로 구성하고 관리할 수 있도록 계층적인 추상화 모델을 제공합니다. 이 계층 구조는 GPU의 물리적 하드웨어 계층(GPU > SM > SP)과 논리적으로 직접 대응되며, 이는 CUDA의 성능과 복잡성을 이해하는 데 매우 중요합니다.1

이러한 계층 구조는 CUDA 프로그래밍의 핵심입니다. 개발자는 처리할 데이터의 구조에 맞춰 스레드, 블록, 그리드의 차원과 크기를 결정합니다. 예를 들어, 2D 이미지 처리를 위해서는 2D 그리드와 2D 블록을 사용하는 것이 자연스럽습니다. 이 구조를 이해하는 것은 GPU의 모든 SM을 유휴 상태 없이 최대한 활용하고, 블록 내 스레드 간의 효율적인 협력을 통해 성능을 극대화하는 데 필수적입니다. 예를 들어 72개의 SM을 가진 GPU의 성능을 모두 끌어내기 위해서는, 최소 72개 이상의 블록으로 구성된 그리드를 실행하여 모든 SM에 작업을 할당해야 합니다.

커널(Kernel) 은 개발자가 작성하지만 GPU 디바이스에서 수많은 스레드에 의해 병렬로 실행되는 C++ 함수입니다. 커널 함수는 __global__ 이라는 CUDA 키워드를 사용하여 선언되며, 이는 해당 함수가 호스트(CPU)에서 호출되어 디바이스(GPU)에서 실행됨을 컴파일러에 알립니다.

호스트 코드에서 커널을 호출하는 것을 커널 실행(Kernel Launch) 이라고 합니다. 커널 실행은 다음과 같은 특별한 <<<...>>> 구문을 사용합니다:

kernel_name<<<grid_dimensions, block_dimensions>>>(argument_list);

여기서 grid_dimensions는 그리드에 포함될 블록의 수와 배치를, block_dimensions는 각 블록에 포함될 스레드의 수와 배치를 지정합니다. 이 커널 실행 구문이 호출되면, CUDA 런타임은 지정된 구성에 따라 스레드의 그리드를 생성하고, 이 스레드들이 GPU에서 커널 코드를 병렬로 실행하도록 작업을 예약합니다. 이처럼 커널은 CPU 중심의 순차적인 프로그램 흐름과 GPU 중심의 대규모 병렬 실행을 연결하는 핵심적인 다리 역할을 합니다.

CUDA의 실행 모델을 이해했다면, 다음으로 정복해야 할 가장 중요한 두 가지 영역은 메모리 시스템과 프로그래밍 구문입니다. 이 파트에서는 개발자가 고성능 코드를 작성하기 위해 반드시 숙달해야 하는 이 두 가지 핵심 요소를 심층적으로 다룹니다.

CUDA 애플리케이션의 성능은 대부분 메모리 접근 패턴에 의해 결정됩니다. GPU는 다양한 속도, 크기, 접근 범위를 가진 복잡한 계층적 메모리 구조를 가지고 있으며, 각 메모리 유형의 특성을 이해하고 데이터를 적재적소에 배치하는 것이 최적화의 핵심입니다.

온칩(On-chip) 메모리는 GPU 칩 내부에 위치하여 접근 지연 시간(latency)이 매우 짧고 대역폭(bandwidth)이 높은, 가장 빠른 메모리 공간입니다.

오프칩(Off-chip) 메모리는 GPU 칩 외부에 위치한 DRAM으로, 용량이 크지만 온칩 메모리에 비해 접근 지연 시간이 길고 대역폭이 상대적으로 낮습니다.

CUDA는 특정 접근 패턴에 최적화된 읽기 전용 캐시 메모리도 제공합니다.

CUDA GPU는 프로그래머가 직접 제어할 수 없는 하드웨어 캐시도 포함하고 있습니다. 각 SM에는 L1 캐시가 있고, 모든 SM이 공유하는 L2 캐시가 존재합니다.3 이 캐시들은 주로 글로벌 메모리와 로컬 메모리 접근의 지연 시간을 줄이기 위해 자동으로 동작합니다. CPU 캐시와의 중요한 차이점 중 하나는, GPU 캐시는 일반적으로 메모리 읽기(load) 동작만 캐싱하고 쓰기(store) 동작은 캐싱하지 않는다는 점입니다. 일부 아키텍처에서는 L1 캐시와 공유 메모리가 동일한 온칩 SRAM을 공유하며, cudaFuncSetCacheConfig와 같은 함수를 통해 두 공간의 크기 비율을 조절할 수 있습니다.3

가장 중요한 점은 CPU(호스트)와 GPU(디바이스)가 물리적으로 분리된 메모리 공간을 가진다는 것입니다. 따라서 GPU에서 연산을 수행하려면, 먼저 호스트 메모리에 있는 데이터를 PCIe 버스를 통해 GPU의 글로벌 메모리로 복사해야 하며, 연산이 끝난 후에는 다시 결과를 호스트 메모리로 가져와야 합니다. 이 데이터 전송 과정은 GPU 내부의 메모리 접근보다 훨씬 느리며, 종종 애플리케이션 전체 성능의 발목을 잡는 심각한 병목 현상을 유발합니다. 따라서 효율적인 CUDA 프로그램은 호스트와 디바이스 간의 데이터 전송 횟수와 양을 최소화하도록 설계되어야 합니다.

아래 표는 CUDA의 복잡한 메모리 계층의 특성을 한눈에 비교하여 개발자가 데이터 배치에 대한 최적의 결정을 내릴 수 있도록 돕습니다.

메모리 유형 위치 범위 접근 권한 지연 시간 대역폭 주요 사용 사례
레지스터 온칩 스레드 단위 R/W 약 1 사이클 매우 높음 빈번하게 접근하는 스레드 전용 변수
로컬 메모리 오프칩 스레드 단위 R/W 높음 낮음 레지스터 스필, 복잡한 지역 배열
공유 메모리 온칩 블록 단위 R/W 낮음 높음 스레드 간 통신, 사용자 관리 캐시
글로벌 메모리 오프칩 그리드 단위 R/W 높음 중간-높음 입출력을 위한 주 데이터 저장소
상수 메모리 오프칩, 캐시됨 그리드 단위 읽기 전용 낮음 (캐시 히트 시) 높음 (브로드캐스트) 모든 스레드에 동일한 값 (예: 계수)
텍스처 메모리 오프칩, 캐시됨 그리드 단위 읽기 전용 낮음 (캐시 히트 시) 중간-높음 2D 공간 지역성을 가진 데이터 (예: 이미지)

효과적인 CUDA 프로그래밍은 본질적으로 메모리 관리 기술입니다. 알고리즘의 계산 자체보다 데이터를 어떻게 배치하고 이동시키느냐가 성능을 좌우하는 경우가 많습니다. 개발자의 사고방식은 “병렬 알고리즘을 작성한다”에서 “계산 유닛에 데이터를 효율적으로 공급하기 위한 데이터 이동 전략을 설계한다”로 전환되어야 합니다. __shared__ 메모리는 이러한 전략의 핵심 도구로, CPU의 자동화된 캐시와 달리 프로그래머가 직접 제어하는 ‘수동 캐시’입니다. 이는 전문가에게 최고의 성능을 달성할 수 있는 강력한 도구를 제공하지만, 동시에 잘못 사용하면(예: 뱅크 충돌) 성능을 저해할 수 있는 복잡성과 책임도 부여합니다.

CUDA C++는 표준 C++를 기반으로 GPU 프로그래밍을 위한 몇 가지 확장 기능과 키워드를 추가한 언어입니다. 이 장에서는 CUDA 개발 환경을 설정하고, 기본적인 프로그래밍 요소들을 사용하는 방법을 다룹니다.

성공적인 CUDA 개발의 첫 단계는 올바른 개발 환경을 구축하는 것입니다. 이 과정은 운영체제와 사용 목적에 따라 약간의 차이가 있지만, 공통적으로 NVIDIA 드라이버, CUDA 툴킷, 그리고 필요에 따라 cuDNN과 같은 라이브러리를 설치해야 합니다.

이처럼 다소 복잡한 설치 과정은 CUDA가 단순한 소프트웨어 라이브러리가 아니라, 하드웨어 드라이버부터 컴파일러, 런타임에 이르기까지 긴밀하게 통합된 하나의 ‘생태계’임을 보여줍니다. 개발자는 이 생태계에 진입하는 첫 단계부터 NVIDIA가 정한 호환성 규칙을 엄격히 따라야 하며, 이는 CUDA의 강력한 성능과 안정성을 보장하는 동시에 특정 벤더에 종속되는 시작점이 되기도 합니다.

CUDA C++는 표준 C++에 몇 가지 중요한 키워드를 추가하여 호스트와 디바이스의 역할, 그리고 메모리 공간을 명시적으로 구분합니다.

CUDA 소스 코드는 일반적으로 .cu 확장자를 가진 파일에 작성됩니다. 이 파일을 컴파일하기 위해 NVIDIA는 NVCC(NVIDIA CUDA Compiler) 라는 전용 컴파일러를 제공합니다.

nvcc의 컴파일 과정은 독특합니다. .cu 파일 내에서 호스트 코드(일반 C++ 코드)와 디바이스 코드(__global__ 함수 등)를 분리합니다. 호스트 코드는 시스템에 설치된 표준 C++ 컴파일러(예: Linux의 GCC, Windows의 MSVC)를 통해 컴파일되고, 디바이스 코드는 NVCC 자체에 의해 PTX(Parallel Thread Execution) 라는 중간 단계의 어셈블리 언어로 컴파일됩니다. 마지막으로 이 둘을 링크하여 최종 실행 파일을 생성합니다.

가장 기본적인 컴파일 명령어는 다음과 같습니다:

nvcc my_code.cu -o my_app

GPU에서 실행되는 수많은 스레드들은 각자 자신이 처리해야 할 데이터가 무엇인지 알아야 합니다. CUDA는 이를 위해 각 스레드에게 자신의 위치 정보를 알려주는 내장 변수(Built-in Variables) 를 제공합니다.1 이 변수들은 커널 함수 내에서 바로 사용할 수 있습니다.

이 변수들을 조합하면 그리드 내에서 각 스레드의 고유한 전역 인덱스(global index) 를 계산할 수 있습니다. 이 인덱싱은 CUDA 프로그래밍의 가장 근본적이고 중요한 개념입니다. 예를 들어, 1차원 배열을 처리하기 위해 1차원 그리드와 1차원 블록을 사용하는 경우, 각 스레드의 전역 인덱스 idx는 다음과 같이 계산됩니다:

int idx = blockIdx.x * blockDim.x + threadIdx.x;

이 공식은 CUDA의 추상적인 병렬 처리 모델을 GPU의 물리적인 실행 계층에 매핑하는 핵심적인 연결고리입니다. “N개의 데이터를 병렬로 처리하라”는 논리적 문제를 “573번 스레드는 573번 데이터를 처리하라”는 구체적인 명령으로 변환하는 역할을 합니다. 이 인덱싱 공식을 이해하고 올바르게 사용하는 것이 CUDA 프로그래머가 되어 첫 번째로 넘어야 할 가장 중요한 관문입니다.

지금까지 다룬 CUDA의 아키텍처, 실행 모델, 메모리 계층, 프로그래밍 환경에 대한 이론적 지식을 바탕으로, 이제 실제 코드를 작성하고 실행하는 과정을 단계별로 살펴봅니다. 이 파트에서는 가장 대표적인 예제인 벡터 덧셈을 통해 CUDA 프로그래밍의 전체 워크플로우를 구체화하고, 이론이 실제 코드에서 어떻게 구현되는지 명확히 보여줍니다.

모든 CUDA 프로그램은 거의 예외 없이 정형화된 절차를 따릅니다. 이 워크플로우를 이해하는 것은 CUDA 프로그래밍의 기본 골격을 세우는 것과 같습니다.

CUDA 프로그래밍은 분리된 두 개의 메모리 공간, 즉 호스트(CPU) 메모리와 디바이스(GPU) 메모리를 다루는 것에서 시작합니다.

  1. 호스트 메모리 할당: 먼저, 입력 데이터와 최종 결과를 저장할 공간을 호스트 메모리에 할당합니다. 이는 표준 C/C++의 malloc 또는 new 연산자를 사용하여 수행됩니다.

    float *h_A, *h_B, *h_C; // h_ prefix for host pointers
h_A = (float*)malloc(size_in_bytes);
h_B = (float*)malloc(size_in_bytes);
h_C = (float*)malloc(size_in_bytes);

  2. 디바이스 메모리 할당: 다음으로, GPU에서 연산을 수행하는 동안 데이터를 저장할 공간을 디바이스 메모리에 할당해야 합니다. 이는 CUDA 런타임 API 함수인 cudaMalloc()을 사용하여 이루어집니다.

    float *d_A, *d_B, *d_C; // d_ prefix for device pointers
cudaMalloc((void**)&d_A, size_in_bytes);
cudaMalloc((void**)&d_B, size_in_bytes);
cudaMalloc((void**)&d_C, size_in_bytes);

  3. 메모리 해제: 프로그램이 종료되기 전, 할당된 모든 메모리는 반드시 해제해야 합니다. 호스트 메모리는 free() 또는 delete로, 디바이스 메모리는 cudaFree() 함수로 해제합니다.

    free(h_A); free(h_B); free(h_C);
cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);

호스트와 디바이스에 각각 메모리 공간이 준비되면, 연산에 필요한 데이터를 두 공간 사이에서 주고받아야 합니다. 이 데이터 전송은 cudaMemcpy() 함수를 통해 이루어지며, 전송 방향을 명시하는 인자를 함께 전달합니다.

  1. 호스트에서 디바이스로 데이터 복사: GPU에서 연산을 시작하기 전에, 호스트 메모리에 초기화된 입력 데이터를 디바이스 메모리로 복사해야 합니다. 이때 cudaMemcpyHostToDevice 플래그를 사용합니다.

    // Initialize h_A and h_B on the host...
cudaMemcpy(d_A, h_A, size_in_bytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size_in_bytes, cudaMemcpyHostToDevice);

  2. 디바이스에서 호스트로 데이터 복사: 커널 실행이 완료된 후, 디바이스 메모리에 저장된 결과 데이터를 호스트 메모리로 다시 가져와야 결과를 확인하거나 후속 처리를 할 수 있습니다. 이때는 cudaMemcpyDeviceToHost 플래그를 사용합니다.

    // After kernel execution...
cudaMemcpy(h_C, d_C, size_in_bytes, cudaMemcpyDeviceToHost);

데이터가 디바이스 메모리에 준비되면, GPU에서 병렬 연산을 수행할 커널을 실행합니다.

  1. 실행 구성(Execution Configuration) 정의: 커널을 실행하기 전에, 얼마나 많은 스레드를 어떤 구조로 실행할지 결정해야 합니다. 이는 그리드(grid)의 차원과 블록(block)의 차원을 정의하는 것을 의미합니다. 일반적으로 처리할 데이터의 총 개수(numElements)와 블록당 스레드 수(threadsPerBlock)를 기반으로 필요한 블록의 수를 계산합니다. 이때 정수 나눗셈으로 인해 일부 데이터가 누락되지 않도록 올림 계산을 하는 것이 일반적입니다.

    int threadsPerBlock = 256;
int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;

  2. 커널 호출: 정의된 실행 구성을 <<<...>>> 구문 안에 넣어 커널 함수를 호출합니다. 이 구문은 커널에 전달할 인자 목록 바로 앞에 위치합니다.

    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);

이제 앞서 설명한 워크플로우를 실제 코드로 구현한 벡터 덧셈 예제를 통해 구체적으로 살펴봅니다. 이 예제는 CUDA의 “Hello, World!”와 같아서, CUDA 프로그래밍의 핵심 개념을 가장 명확하게 보여줍니다.

먼저, 성능 비교를 위한 기준선(baseline)으로 일반적인 C++ CPU 코드를 작성합니다. 이 코드는 for 루프를 사용하여 두 벡터의 각 원소를 순차적으로 더합니다.

void vectorAddCPU(const float* A, const float* B, float* C, int numElements) {
    for (int i = 0; i < numElements; ++i) {
        C[i] = A[i] + B[i];
    }
}

다음은 전체 CUDA 프로그램(.cu 파일)입니다. 이 코드에는 __global__로 선언된 커널 함수와, 앞서 설명한 5단계 워크플로우를 모두 수행하는 main 함수가 포함되어 있습니다.

#include <iostream>
#include <cuda_runtime.h>

// Error checking macro
#define CUDA_CHECK(err) { \
    cudaError_t err_ = (err); \
    if (err_!= cudaSuccess) { \
        std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \
        std::cerr << ": " << cudaGetErrorString(err_) << std::endl; \
        exit(EXIT_FAILURE); \
    } \
}

// CUDA Kernel to add two vectors
__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < numElements) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    // 1. Host and Device Memory Management
    int numElements = 50000;
    size_t size = numElements * sizeof(float);

    // Allocate host memory
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);

    // Initialize host vectors
    for (int i = 0; i < numElements; ++i) {
        h_A[i] = rand() / (float)RAND_MAX;
        h_B[i] = rand() / (float)RAND_MAX;
    }

    // Allocate device memory
    float *d_A, *d_B, *d_C;
    CUDA_CHECK(cudaMalloc((void**)&d_A, size));
    CUDA_CHECK(cudaMalloc((void**)&d_B, size));
    CUDA_CHECK(cudaMalloc((void**)&d_C, size));

    // 2. Data Marshalling (Host to Device)
    CUDA_CHECK(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice));

    // 3. Kernel Launch
    int threadsPerBlock = 256;
    int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);

    // 4. Synchronization
    CUDA_CHECK(cudaGetLastError()); // Check for launch errors
    CUDA_CHECK(cudaDeviceSynchronize());

    // 2. Data Marshalling (Device to Host)
    CUDA_CHECK(cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost));

    // Verify results on host...

    // 1. Free memory
    free(h_A); free(h_B); free(h_C);
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);

    std::cout << "Vector addition completed successfully." << std::endl;

    return 0;
}

커널 코드의 핵심은 각 스레드가 처리할 데이터의 인덱스를 계산하는 부분입니다.

int i = blockIdx.x * blockDim.x + threadIdx.x;

이 한 줄의 코드는 CUDA의 병렬 실행 모델을 데이터에 매핑하는 정수(essence)를 담고 있습니다.

또한, if (i < numElements) 조건문은 매우 중요합니다. 커널 실행 시 생성되는 총 스레드 수(blocksPerGrid * threadsPerBlock)가 처리할 데이터의 총 개수(numElements)와 정확히 일치하지 않을 수 있습니다. 이 조건문은 배열의 크기를 벗어나는 인덱스를 가진 ‘잉여’ 스레드들이 메모리의 잘못된 위치에 접근하는 것을 방지하는 안전장치 역할을 합니다.

이러한 프로그래밍 패턴은 CPU의 for 루프를 GPU의 병렬 그리드로 대체하는 패러다임 전환을 명확하게 보여줍니다. 개발자는 더 이상 루프의 제어 흐름을 직접 작성하지 않습니다. 대신, 루프의 한 번의 반복(iteration)에 해당하는 작업을 커널에 정의하고, CUDA 런타임에 이 작업을 수많은 스레드를 통해 동시에 실행하도록 지시합니다. 이러한 방식을 루프 병렬화(loop parallelism) 라고 하며, 이는 GPU에서의 데이터 병렬 처리의 기본 원리입니다.

위 CUDA 코드를 실행하고 CPU 버전과 성능을 비교하면 중요한 사실을 발견할 수 있습니다. 벡터의 크기가 작을 때는, GPU의 연산 속도가 아무리 빨라도 호스트와 디바이스 간의 데이터를 복사하는 cudaMemcpy의 오버헤드 때문에 오히려 CPU 버전보다 전체 실행 시간이 더 오래 걸릴 수 있습니다.

하지만 벡터의 크기가 수백만 개 이상으로 커지면(예: 16 MiB), GPU의 대규모 병렬 처리 능력이 빛을 발하기 시작합니다. 순수 커널 실행 시간만 비교하면 GPU가 CPU보다 수십 배 더 빠를 수 있습니다. 그러나 공정한 성능 비교를 위해서는 반드시 데이터 전송 시간을 포함한 전체 실행 시간(end-to-end time)을 측정해야 합니다. 예를 들어, 커널 실행이 2ms, 데이터 전송이 62ms가 걸렸다면, 실제 체감 성능 향상은 데이터 전송 시간을 포함한 64ms를 기준으로 평가되어야 합니다.

이 분석은 CUDA 프로그래밍의 핵심 교훈을 다시 한번 강조합니다: 연산은 저렴하지만, 데이터 이동은 비싸다. 따라서 고성능을 달성하기 위한 노력은 종종 알고리즘 자체의 개선보다는 호스트와 디바이스 간, 그리고 디바이스 내 글로벌 메모리와 온칩 메모리 간의 데이터 이동을 최소화하고 최적화하는 데 집중되어야 합니다.

기본적인 CUDA 프로그래밍 방법을 익혔다면, 이제 GPU의 잠재력을 최대한 끌어내는 전문가 수준의 최적화 기법을 탐구할 차례입니다. 이 파트에서는 CUDA 성능을 좌우하는 가장 중요한 두 가지 요소인 메모리 접근과 하드웨어 활용률을 극대화하는 전략을 심층적으로 다룹니다.

CUDA 성능 최적화의 80%는 메모리 최적화라고 해도 과언이 아닙니다. 특히 느린 글로벌 메모리에 대한 접근을 어떻게 효율적으로 만드느냐가 관건입니다.

공유 메모리는 잘못된 메모리 접근 패턴을 바로잡는 가장 강력한 도구입니다. 알고리즘이 본질적으로 병합되지 않는 접근(예: 행렬 전치)을 요구할 때, 공유 메모리를 중간 단계의 스테이징 영역으로 사용하여 접근 패턴을 재구성할 수 있습니다. 이 기법을 타일링(Tiling) 또는 코너 터닝(Corner Turning) 이라고 부릅니다.

  1. 병합된 읽기: 블록 내의 스레드들이 협력하여 글로벌 메모리에서 데이터 ‘타일’을 공유 메모리로 읽어옵니다. 이 읽기 작업은 메모리 병합이 일어나도록 신중하게 설계됩니다.
  2. 동기화: 모든 스레드가 데이터 로딩을 마칠 때까지 __syncthreads()를 호출하여 대기합니다.
  3. 빠른 비병합 접근: 스레드들은 이제 고속의 공유 메모리에서 알고리즘에 필요한 패턴(병합되지 않았더라도)으로 자유롭게 데이터에 접근합니다. 공유 메모리는 온칩에 있어 지연 시간이 매우 짧으므로, 여기서의 비병합 접근은 성능에 큰 영향을 미치지 않습니다.
  4. 병합된 쓰기 (필요 시): 계산 결과를 다시 글로벌 메모리에 쓸 때도 공유 메모리를 버퍼로 사용하여 병합된 쓰기를 수행할 수 있습니다.

이 기법은 행렬 전치(matrix transpose) 예제에서 극적인 효과를 보여줍니다. 순진한 구현에서는 행을 읽는 것은 병합되지만 열을 쓰는 것은 스트라이드 접근이 되어 성능이 저하됩니다. 공유 메모리를 사용하면 데이터를 온칩에서 재정렬한 후 병합된 쓰기를 수행할 수 있어 성능을 크게 향상시킬 수 있습니다.

데이터를 어떻게 구성하는지도 메모리 병합에 직접적인 영향을 미칩니다.

이러한 최적화 기법들은 CUDA 성능 튜닝이 단순한 알고리즘 개선이 아니라, 하드웨어의 숨겨진 규칙에 데이터를 맞추는 과정임을 보여줍니다. 메모리 병합은 소프트웨어 개발자에게 주어진 권장 사항이 아니라, 하드웨어 메모리 컨트롤러가 부과하는 엄격한 규칙입니다. 이를 위반하면 직접적이고 심각한 성능 저하가 발생합니다. 따라서 가장 효율적인 CUDA 개발자는 “어떻게 병렬로 계산할까?”를 묻기 전에 “어떻게 데이터를 배치해야 메모리 접근을 병합할 수 있을까?”를 먼저 고민합니다.

메모리 접근 패턴을 최적화했다면, 다음 단계는 GPU의 계산 유닛을 최대한 쉬지 않고 일하게 만드는 것입니다. 이는 점유율(Occupancy) 과 밀접한 관련이 있습니다.

점유율이 높을수록 항상 성능이 좋은 것은 아니라는 점을 이해하는 것이 중요합니다. 때로는 점유율을 높이려는 시도가 오히려 성능을 저하시킬 수 있습니다.

예를 들어, __launch_bounds__ 지정자나 컴파일러 옵션을 사용하여 스레드당 최대 레지스터 사용량을 강제로 줄이면, 더 많은 블록이 SM에 상주하게 되어 이론적 점유율은 높아질 수 있습니다.3 하지만 스레드가 필요한 만큼의 레지스터를 할당받지 못하면, 컴파일러는 일부 변수를 느린 로컬 메모리로 내보내게 됩니다. 이를 레지스터 스필링(register spilling) 이라고 하며, 이로 인한 성능 저하가 점유율 증가로 얻는 이득보다 더 클 수 있습니다.3

따라서 최적화의 목표는 무조건 점유율을 100%로 만드는 것이 아니라, ‘최적의 균형점’을 찾는 것입니다. 즉, 메모리 지연 시간을 숨길 수 있을 만큼 충분한 점유율을 확보하면서도, 레지스터 스필링이 발생하지 않을 만큼 스레드당 충분한 레지스터를 제공하는 것입니다. 이 균형점은 커널의 특성에 따라 다르며, 종종 여러 실행 구성을 실험적으로 테스트하며 찾아야 합니다.

경험적으로 알려진 몇 가지 모범 사례(heuristics)는 최적의 실행 구성을 찾는 데 좋은 출발점을 제공합니다.2

앞서 언급했듯이, 한 워프 내에서 스레드들이 if-else와 같은 조건문으로 인해 서로 다른 코드 경로를 실행하면, 하드웨어는 이 경로들을 직렬화하여 실행하므로 심각한 성능 저하가 발생합니다. 이를 워프 분기(warp divergence) 라고 합니다.

고성능을 위해서는 워프 분기를 최소화해야 합니다. 예를 들어, 조건문이 스레드 ID에 따라 달라진다면, threadIdx.x와 같은 개별 ID 대신 threadIdx.x / 32 (워프 ID)와 같이 워프 전체에 걸쳐 동일한 값을 갖는 조건을 사용하도록 코드를 재구성하는 것이 좋습니다.

결론적으로, CUDA 성능 최적화는 여러 변수 사이의 균형을 맞추는 복잡한 과정입니다. 블록 크기, 레지스터 사용량, 공유 메모리 사용량은 서로 영향을 주며 점유율을 통해 최종 성능에 기여합니다. 정답은 하나가 아니며, 커널의 특성에 따라 최적의 구성이 달라집니다. 이것이 바로 NVIDIA가 제공하는 프로파일링 도구(Nsight Compute)를 사용하여 애플리케이션의 병목 지점을 ‘평가(Assess)’하고, 병렬화하며, 반복적으로 ‘최적화(Optimize)’하는 APOD(Assess, Parallelize, Optimize, Deploy) 개발 사이클이 중요한 이유입니다.

CUDA 프로그래밍의 가장 강력한 측면 중 하나는 개발자가 직접 저수준 최적화의 수고를 덜 수 있도록 NVIDIA가 제공하는 방대하고 강력한 라이브러리 생태계입니다. 많은 경우, 직접 커널을 작성하는 것보다 고도로 최적화된 라이브러리를 사용하는 것이 훨씬 더 효율적이고 빠른 결과를 가져옵니다.

NVIDIA의 CUDA 라이브러리들은 각 GPU 아키텍처의 특성을 가장 잘 아는 전문가들에 의해 수작업으로 튜닝됩니다. 이 라이브러리들은 새로운 하드웨어 기능(예: 텐서 코어)을 자동으로 활용하며, 개발자가 직접 구현하기 매우 어려운 복잡한 최적화 기법들이 이미 적용되어 있습니다. 따라서 행렬 연산, 딥러닝, 신호 처리와 같은 표준적인 문제에 대해서는 라이브러리를 사용하는 것이 ‘바퀴를 재발명’하는 것을 피하고, 개발 시간을 단축하며, 최고의 성능을 보장하는 가장 현명한 방법입니다.

라이브러리 기반 가속의 실제적인 효과는 NVIDIA가 공개한 컴퓨터 비전 파이프라인 가속 사례에서 명확히 확인할 수 있습니다.9

이 사례는 CUDA 생태계가 어떻게 실제 비즈니스 문제 해결에 기여하는지를 명확히 보여줍니다. CUDA 라이브러리들은 개발자가 저수준 최적화의 어려움에서 벗어나 애플리케이션 로직에 집중하게 함으로써, AI 및 HPC와 같은 고부가가치 분야에서 NVIDIA의 지배력을 공고히 하는 가장 강력한 ‘해자’ 역할을 합니다. 경쟁사가 단순히 GPU 하드웨어를 만드는 것을 넘어, 이처럼 깊고 넓은 소프트웨어 생태계를 구축하고 주요 프레임워크 개발사들을 설득해야만 비로소 경쟁이 가능해집니다. 이는 AMD의 ROCm이 직면한 가장 큰 도전 과제이기도 합니다.

이 마지막 파트에서는 한 걸음 물러나 CUDA를 더 넓은 산업적 맥락에서 평가합니다. CUDA 생태계의 강점과 약점을 분석하고, 주요 경쟁 기술인 OpenCL 및 ROCm과 직접적으로 비교하여 CUDA의 전략적 위치를 조망합니다.

2006년 처음 소개된 이래, CUDA는 수많은 연구 및 상용 애플리케이션에 채택되며 특정 분야에서는 없어서는 안 될 기술로 자리 잡았습니다.

CUDA는 강력한 성능과 생태계를 자랑하지만, 명확한 장점과 함께 본질적인 한계도 가지고 있습니다.

CUDA의 지배적인 위치에 도전하는 주요 대안 기술로는 OpenCL과 ROCm이 있습니다. 각 프레임워크는 서로 다른 철학과 장단점을 가지고 있습니다.

CUDA의 독점적인 특성은 개발자와 기업에 중요한 전략적 시사점을 가집니다.10

아래 표는 개발자나 기술 관리자가 플랫폼을 선택할 때 고려해야 할 핵심적인 차이점들을 요약합니다.

특징 CUDA (NVIDIA) OpenCL (Khronos Group) ROCm (AMD)
라이선스 독점 (Proprietary) 개방형 표준 (Open Standard) 오픈 소스 (Open Source)
하드웨어 지원 NVIDIA GPU 전용 NVIDIA, AMD, Intel GPU, CPU, FPGA 등 AMD GPU 전용 (일부 모델)
이식성 없음 높음 (이론상) 낮음 (HIP가 CUDA 코드 이식 지원)
성능 NVIDIA 하드웨어에서 최고 벤더 구현에 따라 상이, 일반적으로 네이티브보다 낮음 AMD 하드웨어에서 CUDA와 경쟁력 있으나, 편차 존재
라이브러리 생태계 매우 풍부 (cuDNN, cuBLAS 등) 제한적 성장 중 (rocBLAS 등), CUDA에 비해 부족
개발 도구/디버깅 성숙하고 우수 (Nsight) 파편화됨, 벤더별로 상이 개선 중, CUDA보다 미성숙
성숙도/안정성 매우 높음 벤더 구현에 따라 상이 상대적으로 낮음, 과거 안정성 문제 존재
주요 사용 사례 NVIDIA 기반 AI/HPC 이식성이 중요한 이기종 시스템 AMD 기반 AI/HPC

이러한 경쟁 구도는 정체되어 있지 않습니다. NVIDIA의 시장 가치와 AMD의 가치 차이가 스페인의 연간 GDP보다 크다는 사실은, AMD와 그 고객들에게 CUDA의 독점을 깨뜨려야 할 엄청난 재정적 동기를 부여합니다. 이에 따라 AMD는 ROCm에 대한 투자를 늘리고 있으며, CUDA 바이너리를 직접 실행하려는 시도와 같이 전환 장벽을 낮추기 위한 전략적인 움직임을 보이고 있습니다. ROCm은 빠르게 성숙하며 CUDA의 대안으로 부상하고 있습니다.

그러나 ROCm이 극복해야 할 가장 큰 장벽은 기술적인 격차보다 ‘신뢰’의 문제입니다. 과거 AMD의 파편화된 지원과 일관성 없는 로드맵은 개발자 커뮤니티에 깊은 불신을 남겼습니다. 기술만으로는 생태계를 이길 수 없습니다. AMD가 CUDA와 진정으로 경쟁하기 위해서는 기술적 동등성을 넘어, 개발자들에게 장기적이고 일관된 지원을 약속하고 그 신뢰를 구축해야만 합니다.

CUDA는 지난 10여 년간 컴퓨팅의 지형을 근본적으로 바꾸어 놓았습니다. GPU를 단순한 그래픽 처리 장치에서 과학과 인공지능의 발전을 이끄는 핵심 동력으로 변모시킨 CUDA의 영향력은 앞으로도 오랫동안 지속될 것입니다.

이 보고서를 통해 살펴본 효과적인 CUDA 개발의 핵심 원칙은 다음과 같이 요약할 수 있습니다.

  1. 병렬적으로 사고하라: 순차적인 for 루프 대신, 그리드, 블록, 스레드의 계층 구조로 문제를 분해하고 매핑하는 데이터 병렬적 사고방식을 가져야 합니다.
  2. 메모리를 꼼꼼하게 관리하라: 성능은 계산이 아닌 메모리 접근에 의해 결정됩니다. 메모리 병합, 공유 메모리 활용, 데이터 전송 최소화는 선택이 아닌 필수입니다.
  3. 하드웨어를 위해 최적화하라: 워프 분기 최소화, 최적의 점유율 확보 등 GPU 아키텍처의 특성을 이해하고 그에 맞춰 코드를 설계해야 합니다.
  4. 생태계를 적극적으로 활용하라: 직접 모든 것을 구현하려 하지 말고, NVIDIA가 제공하는 고도로 최적화된 라이브러리(cuBLAS, cuDNN 등)를 적극적으로 사용하여 생산성과 성능을 동시에 잡아야 합니다.

CUDA의 미래는 몇 가지 중요한 기술적 흐름 속에서 전개될 것입니다.

결론적으로, CUDA는 강력한 성능과 타의 추종을 불허하는 성숙한 생태계를 바탕으로 가까운 미래에도 AI 및 HPC 분야에서 지배적인 위치를 유지할 것이 확실합니다. 그러나 업계의 개방형 표준에 대한 요구와 경쟁사들의 거센 추격은 CUDA가 끊임없이 혁신하고 발전해야 하는 동력이 될 것입니다. 개발자들은 이러한 거대한 기술적 흐름 속에서 각 플랫폼의 장단점을 명확히 이해하고, 자신의 목표에 가장 적합한 도구를 선택하는 전략적 안목을 갖추어야 할 것입니다.

  1. CUDA 프로그래밍 - CUDA 스레드 계층 - Kudos IT Daily - 티스토리, accessed July 6, 2025, https://kudositdaily.tistory.com/entry/CUDA-%ED%94%84%EB%A1%9C%EA%B7%B8%EB%9E%98%EB%B0%8D-CUDA-%EC%8A%A4%EB%A0%88%EB%93%9C-%EA%B3%84%EC%B8%B5
    1. Preface - CUDA C++ Best Practices Guide 12.9 documentation, accessed July 6, 2025, https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
  3. CUDA Memory Model - 별준 - 티스토리, accessed July 6, 2025, https://junstar92.tistory.com/283
  4. Window에서 CUDA 11.8, cuDNN 8.7.0 설치 - velog, accessed July 6, 2025, https://velog.io/@ksy5098/Window%EC%97%90%EC%84%9C-CUDA-11.8-cuDNN-8.7.0-%EC%84%A4%EC%B9%98
  5. Step-by-Step Guide to Installing CUDA and cuDNN for GPU …, accessed July 6, 2025, https://www.digitalocean.com/community/tutorials/install-cuda-cudnn-for-gpu
  6. Achieved Occupancy - NVIDIA Docs, accessed July 6, 2025, https://docs.nvidia.com/gameworks/content/developertools/desktop/analysis/report/cudaexperiments/kernellevel/achievedoccupancy.htm
    1. Introduction - cuBLAS 12.9 documentation - NVIDIA Docs, accessed July 6, 2025, https://docs.nvidia.com/cuda/cublas/index.html
  8. CUDA 툴킷 다운로드하고, GPU를 자유자재로 활용해보자! - NVIDIA …, accessed July 6, 2025, https://blogs.nvidia.co.kr/blog/cuda-toolkit/
  9. CV-CUDA로 AI 기반 컴퓨터 비전을 위한 처리량 증가 및 비용 절감, accessed July 6, 2025, https://developer.nvidia.com/ko-kr/blog/increasing-throughput-and-reducing-costs-for-computer-vision-with-cv-cuda/
  10. CUDA - 위키백과, 우리 모두의 백과사전, accessed July 6, 2025, https://ko.wikipedia.org/wiki/CUDA