Tiny GPU: Verilog로 구현한 최소한의 GPU

 (github.com/adam-maj)
2P by GN⁺ | ★ favorite | 댓글 1개
  • tiny-gpu는 GPU가 하드웨어 수준에서 어떻게 동작하는지 밑바닥부터 학습하기 위한 Verilog 기반 최소 GPU 구현체이며, 그래픽 전용 하드웨어보다 GPGPU와 ML 가속기의 공통 원리에 초점을 둠
  • 구현은 문서화된 Verilog 15개 미만 파일, 아키텍처와 ISA 문서, 행렬 덧셈·곱셈 커널, 커널 시뮬레이션과 실행 추적 지원으로 구성됨
  • GPU는 한 번에 단일 커널을 실행하며, 프로그램 메모리와 데이터 메모리를 로드하고 thread_count를 설정한 뒤 start 신호를 올려 커널을 시작함
  • 단순화를 위해 각 코어는 한 번에 하나의 block을 처리하고, 각 스레드마다 ALU·LSU·PC·레지스터 파일을 갖지만 모든 스레드가 매 명령 뒤 같은 PC로 수렴한다고 가정함
  • 현대 GPU의 다층 캐시, 공유 메모리, 메모리 병합, 파이프라이닝, 워프 스케줄링, 분기 발산, 배리어 같은 기능은 대부분 제외되어 학습용 구조를 우선함

tiny-gpu가 풀려는 문제

  • CPU는 아키텍처부터 제어 신호까지 학습할 수 있는 자료가 많지만, 현대 GPU의 저수준 기술 세부사항은 경쟁적인 시장 때문에 대부분 독점적으로 남아 있음
  • GPU 프로그래밍 자료는 많지만, GPU가 하드웨어 수준에서 어떻게 동작하는지 배울 수 있는 자료는 거의 없음
  • 오픈소스 GPU 구현체인 Miaow와 VeriGPU는 기능 완성과 동작을 목표로 해 구조가 복잡함
  • tiny-gpu는 생산급 그래픽 카드의 많은 복잡성을 제거하고, 현대 하드웨어 가속기에 공통적인 핵심 요소에 집중함
    • GPU 아키텍처의 중요한 구성 요소
    • SIMD 프로그래밍 모델이 하드웨어에서 구현되는 방식
    • 제한된 메모리 대역폭을 GPU가 다루는 방식

전체 아키텍처

  • tiny-gpu는 한 번에 하나의 커널만 실행하도록 설계됨
  • 커널 실행 절차는 다음과 같음
    • 전역 프로그램 메모리에 커널 코드 로드
    • 데이터 메모리에 필요한 데이터 로드
    • 장치 제어 레지스터에 실행할 스레드 수 지정
    • start 신호를 high로 설정해 커널 실행
  • GPU는 다음 유닛으로 구성됨
    • 장치 제어 레지스터
    • 디스패처
    • 가변 개수의 컴퓨트 코어
    • 데이터 메모리와 프로그램 메모리용 메모리 컨트롤러
    • 캐시

커널 실행과 스레드 배분

  • 장치 제어 레지스터는 커널 실행 메타데이터를 저장하는 역할을 하며, tiny-gpu에서는 실행할 전체 스레드 수인 thread_count만 저장함
  • 디스패처는 커널이 시작되면 스레드를 여러 컴퓨트 코어에 분배함
    • 병렬 실행 가능한 스레드 묶음을 block으로 구성함
    • 사용 가능한 코어에 block을 보내 처리함
    • 모든 block 처리가 끝나면 커널 실행 완료를 알림
  • 단순화된 코어는 한 번에 하나의 block을 처리함
  • 각 스레드는 전용 ALU, LSU, PC, 레지스터 파일을 가짐
  • 스레드 명령 실행을 이 자원 위에서 관리하는 일이 GPU의 어려운 문제 중 하나임

메모리 구조와 컨트롤러

  • GPU는 외부 전역 메모리와 인터페이스하도록 만들어졌으며, 단순화를 위해 데이터 메모리와 프로그램 메모리를 분리함
  • 데이터 메모리 사양
    • 8비트 주소 지정
    • 총 256행
    • 8비트 데이터
    • 각 행은 256 미만 값을 저장
  • 프로그램 메모리 사양
    • 8비트 주소 지정
    • 총 256행
    • 16비트 데이터
    • ISA에 따라 각 명령은 16비트
  • 메모리 컨트롤러는 코어에서 오는 메모리 요청을 추적하고, 실제 외부 메모리 대역폭에 맞춰 요청을 제한하며, 응답을 올바른 자원으로 전달함
  • 각 메모리 컨트롤러는 전역 메모리 대역폭에 따라 고정된 수의 채널을 가짐
  • 캐시는 작업 중인 기능이며, 외부 메모리에서 가져온 데이터를 장치 내 SRAM에 저장해 이후 요청에서 더 빠르게 가져오고 메모리 대역폭을 새 데이터에 사용할 수 있게 함

코어 내부 구성

  • 각 코어에는 단일 스케줄러가 있어 스레드 실행을 관리함
  • tiny-gpu 스케줄러는 하나의 block 명령을 끝까지 실행한 뒤 새 block을 가져오며, 모든 스레드의 명령을 동기화된 순서로 실행함
  • 더 발전된 스케줄러에서는 파이프라이닝과 워프 스케줄링으로 자원 활용률을 높일 수 있음
  • 스케줄러의 주요 제약은 전역 메모리에서 데이터를 로드하고 저장할 때 발생하는 지연 시간임
    • 대부분의 명령은 동기적으로 실행 가능함
    • LDR와 STR 같은 load-store 작업은 비동기라 긴 대기 시간을 중심으로 명령 실행을 구성해야 함
  • Fetcher는 현재 프로그램 카운터의 명령을 프로그램 메모리에서 비동기로 가져옴
  • Decoder는 가져온 명령을 스레드 실행용 제어 신호로 디코드함
  • 각 스레드의 레지스터 파일은 계산 중인 데이터를 보관하며 SIMD 패턴을 가능하게 함
    • 읽기 전용 레지스터에는 %blockIdx, %blockDim, %threadIdx가 들어 있음
    • 커널은 로컬 스레드 ID에 따라 다른 데이터로 실행될 수 있음
  • 각 스레드의 ALU는 ADD, SUB, MUL, DIV 산술 명령을 처리함
  • CMP는 두 레지스터 차이의 결과가 음수·0·양수인지 출력하고, 결과를 PC 유닛의 NZP 레지스터에 저장함
  • 각 스레드의 LSU는 전역 데이터 메모리에 접근하며, LDR와 STR 및 비동기 메모리 대기 시간을 처리함
  • 각 스레드의 PC는 다음 실행 명령을 결정함
    • 기본적으로 명령마다 1씩 증가함
    • BRnzp는 이전 CMP가 설정한 NZP 레지스터 조건이 맞으면 특정 프로그램 메모리 행으로 분기함
    • 루프와 조건문은 이 방식으로 구현됨
  • tiny-gpu는 단순화를 위해 모든 스레드가 매 명령 뒤 같은 PC로 수렴한다고 가정함
  • 실제 GPU에서는 개별 스레드가 서로 다른 PC로 분기할 수 있고, 이때 함께 처리되던 스레드 그룹이 여러 실행 흐름으로 갈라지는 분기 발산이 발생함

ISA

  • tiny-gpu는 행렬 덧셈과 행렬 곱셈 같은 개념 증명용 단순 커널을 실행하기 위해 11개 명령 ISA를 구현함
  • 지원 명령
    • BRnzp: NZP 조건이 맞으면 다른 프로그램 메모리 행으로 점프
    • CMP: 두 레지스터 값을 비교하고 결과를 NZP 레지스터에 저장
    • ADD, SUB, MUL, DIV: 텐서 수학을 위한 기본 산술 연산
    • LDR: 전역 메모리에서 데이터 로드
    • STR: 전역 메모리에 데이터 저장
    • CONST: 상수 값을 레지스터에 로드
    • RET: 현재 스레드 실행 종료 신호
  • 각 레지스터는 4비트로 지정되어 총 16개 레지스터를 가짐
    • R0부터 R12까지 13개는 읽기·쓰기 가능한 자유 레지스터
    • 마지막 3개는 SIMD에 필요한 %blockIdx, %blockDim, %threadIdx를 제공하는 읽기 전용 특수 레지스터

실행 흐름

  • 각 코어는 명령을 실행할 때 다음 단계의 제어 흐름을 따름
    • FETCH: 현재 PC의 다음 명령 가져오기
    • DECODE: 명령을 제어 신호로 디코드
    • REQUEST: LDR 또는 STR가 필요한 경우 전역 메모리에 데이터 요청
    • WAIT: 필요한 경우 전역 메모리 응답 대기
    • EXECUTE: 데이터에 대한 계산 실행
    • UPDATE: 레지스터 파일과 NZP 레지스터 갱신
  • 이 제어 흐름은 단순성과 이해 가능성을 위해 구성됨
  • 실제 구현에서는 일부 단계를 압축해 처리 시간을 최적화하거나, 파이프라이닝으로 코어 자원에서 여러 명령 실행을 조율할 수 있음
  • 각 스레드는 전용 레지스터 파일의 데이터를 대상으로 같은 실행 경로를 따라 계산함
  • CPU 다이어그램과 비슷하지만, %blockIdx, %blockDim, %threadIdx가 읽기 전용 레지스터에 있어 SIMD 기능을 가능하게 하는 점이 다름

예제 커널

  • ISA의 개념 증명을 위해 행렬 덧셈과 행렬 곱셈 커널이 작성됨
  • 저장소의 테스트 파일은 이 커널들을 GPU에서 완전히 시뮬레이션하고, 데이터 메모리 상태와 전체 실행 추적을 생성할 수 있음

  • 행렬 덧셈

    • matadd.asm는 1 x 8 행렬 두 개를 더함
    • 8개 원소별 덧셈을 각각 별도 스레드에서 수행함
    • %blockIdx, %blockDim, %threadIdx 레지스터를 사용해 SIMD 프로그래밍을 보여줌
    • LDR와 STR 명령을 사용해 비동기 메모리 관리를 포함함

  • 행렬 곱셈

    • matmul.asm는 2 x 2 행렬 두 개를 곱함
    • 관련 행과 열의 내적을 원소별로 계산함
    • CMP와 BRnzp를 사용해 스레드 내부 분기를 보여줌
    • 모든 분기가 다시 수렴하기 때문에 현재 tiny-gpu 구현에서 동작함

시뮬레이션

  • 커널 시뮬레이션을 실행하려면 iverilog와 cocotb가 필요함
  • 준비 절차
    • brew install icarus-verilog와 pip3 install cocotb로 Verilog 컴파일러와 cocotb 설치
    • sv2v의 최신 버전을 내려받아 압축을 풀고 바이너리를 $PATH에 추가
    • 저장소 루트에서 mkdir build 실행
  • 커널 시뮬레이션은 make test_matadd와 make test_matmul로 실행함
  • 실행 결과는 test/logs의 로그 파일로 출력됨
    • 초기 데이터 메모리 상태
    • 커널의 전체 실행 추적
    • 최종 데이터 메모리 상태
  • 각 로그 파일의 시작 부분에는 입력 행렬이 보이고, 끝부분의 최종 데이터 메모리에는 결과 행렬이 보임
  • 실행 추적은 각 사이클마다 모든 코어의 모든 스레드 실행 상태를 포함함
    • 현재 명령
    • PC
    • 레지스터 값
    • 상태 정보

의도적으로 빠진 고급 GPU 기능

  • tiny-gpu는 단순화를 위해 현대 GPU의 성능·기능 개선 요소를 대부분 제외함

  • 다층 캐시와 공유 메모리

    • 현대 GPU는 전역 메모리 접근을 줄이기 위해 여러 캐시 계층을 사용함
    • tiny-gpu는 요청 자원과 메모리 컨트롤러 사이에 최근 데이터를 저장하는 단일 캐시 계층만 구현함
    • 여러 계층 캐시는 자주 쓰는 데이터를 사용 위치에 더 가깝게 캐시해 로드 시간을 줄임
    • GPU는 같은 block의 스레드가 공유 결과를 주고받을 수 있도록 공유 메모리를 사용하기도 함

  • 메모리 병합

    • 병렬 실행 중인 여러 스레드는 행렬의 인접 원소처럼 연속된 주소에 접근하는 경우가 많음
    • 메모리 병합은 큐에 쌓인 메모리 요청을 분석해 인접 요청을 하나의 트랜잭션으로 합침
    • 주소 지정에 쓰는 시간을 줄이고 요청을 함께 처리하는 데 목적이 있음

  • 파이프라이닝

    • tiny-gpu의 코어는 한 스레드 묶음의 한 명령 실행이 끝나야 다음 명령을 시작함
    • 현대 GPU는 의존성이 있는 명령은 순차 실행을 보장하면서도 여러 순차 명령 실행을 스트리밍함
    • 비동기 메모리 요청 대기 같은 상황에서 코어 자원이 놀지 않도록 자원 활용률을 높임

  • 워프 스케줄링

    • block을 함께 실행 가능한 스레드 배치인 warp로 나눔
    • 한 warp가 대기 중일 때 다른 warp의 명령을 실행해 단일 코어에서 여러 warp를 동시에 처리함
    • 파이프라이닝과 비슷하지만 서로 다른 스레드의 명령을 다룸

  • 분기 발산

    • tiny-gpu는 단일 배치의 모든 스레드가 각 명령 후 같은 PC에 있다고 가정함
    • 실제로는 데이터에 따라 개별 스레드가 다른 줄로 분기할 수 있음
    • 다른 PC를 가진 스레드는 별도 실행 흐름으로 나뉘며, 다시 수렴하는 시점도 관리해야 함

  • 동기화와 배리어

    • 현대 GPU는 같은 block의 스레드 그룹이 특정 지점에 모두 도달할 때까지 기다리도록 배리어를 설정할 수 있음
    • 스레드들이 공유 데이터를 교환해야 하는 경우, 데이터 처리가 완료됐는지 보장하는 데 유용함

다음 작업

  • 향후 개선 항목은 다음과 같음
    • 간단한 명령 캐시 추가
    • Tiny Tapeout 7에서 GPU를 사용할 수 있는 어댑터 구축
    • 기본 분기 발산 추가
    • 기본 메모리 병합 추가
    • 기본 파이프라이닝 추가
    • 사이클 시간을 개선하기 위한 제어 흐름과 레지스터 사용 최적화
    • 그래픽 기능을 보여주기 위한 기본 그래픽 커널 작성 또는 단순 그래픽 하드웨어 추가
  • 저장소 개선을 원하는 사용자는 PR로 기여할 수 있음

함께 보면 좋은 글 β

GN⁺   [-]
Hacker News 의견들
  • GPU 시장이 워낙 경쟁적이라 현대 아키텍처의 저수준 기술 세부사항은 대부분 비공개로 남아 있음
    예외적으로 Intel은 GPU 기술 문서를 많이 공개함: https://kiwitree.net/~lina/intel-gfx-docs/prm/
    i810/815 매뉴얼도 온라인에서 찾을 수 있고, 855/910/915/945가 빠진 965 이전 구간의 이상한 공백을 빼면 문서화가 꽤 꾸준했던 편임
  • 정말 멋진 프로젝트고, 이런 하드웨어 프로젝트가 공개적으로 진행되는 걸 보는 게 좋음
    다만 이건 SIMD 보조 프로세서에 가깝다고 봄
    GPU라고 부르려면 최소한 어떤 형태의 디스플레이 출력은 있어야 한다고 생각함
    최근 Nvidia 등이 서버 전용 그래픽 아키텍처 변형도 GPU로 팔면서 용어가 꽤 느슨해진 건 알지만, GPU 설계에서 그래픽 부분은 지금도 복잡도의 상당 부분을 차지함
    • 그래픽을 처리한다면 출력이 없어도 GPU로 볼 수 있다고 생각함
      출력하지 않는 GPU도 여전히 쓸모가 있음
      내 직장에는 중급 Quadro가 들어간 워크스테이션이 75대쯤 있는데, 카드에는 mini-DisplayPort만 있고 회사는 HDMI 케이블만 사줘서 모두 내장 그래픽에 연결돼 있음
      그래도 그 카드들은 소프트웨어를 가속하고 그래픽을 처리하며, 단지 화면 출력만 하지 않을 뿐임
  • 좋음. 오픈 코어 GPU 작업은 적극 응원함
    다른 예시도 있음: https://github.com/jbush001/NyuziProcessor
    • 이런 오픈 코어 프로세서 중 하나에 대해 최소한의 CUDA 구현이 있으면 좋겠음
      TSMC나 다른 파운드리에서 이런 프로세서를 경제적으로 생산하려면 어느 정도 물량이 필요할까?
  • 정말 훌륭한 프로젝트임
    FPGA를 해보고 싶은데, 솔직히 어디서 시작해야 할지 감을 잡기도 어렵고 분야 전체가 꽤 위압적으로 느껴짐
    최종 목표는 LLM용 가속 카드를 만드는 것인데, 완전히 임의로 정한 목표이긴 해도 이 프로젝트와 겹치는 부분이 많을 것 같고, 아마 더 큰 모델을 로드하기 위한 메모리 오프로딩 부분 정도만 다를 듯함
    • 머릿속에서 프레임을 바꿔야 함
      FPGA 입문은 여러 하위 기술로 쪼개야 하고, 기대치도 조정해야 함
      소프트웨어 엔지니어에게 처음부터 원리에서 컴퓨터 전체를 만들고, 명령어 집합 구조를 쓰고, 기계어를 이해하고, 이를 어셈블리로 바꾸고, Python 코드로 애플리케이션을 만들기 위해 프로그래밍 언어까지 개발하라고 기대하지는 않음
      위에서 시작해 스택 아래로 내려가는 게 맞음
      복잡도를 추상화하고 미리 만들어진 IP로 시스템을 만드는 데 집중하면 FPGA 설계는 꽤 쉬움
      보통 MATLAB 같은 걸 추천하는데, Reference design이 있는 DevKit에서 HDL Coder로 초기 애플리케이션을 만들 수 있기 때문임
      그렇지 않으면 디지털 컴퓨팅 아키텍처, Verilog, 타이밍, 트랜시버/입출력, 핀 계획, Quartus/Vivado, 시뮬레이션/검증, 임베디드 시스템 등을 배우는 막대한 부담이 생김
      요약하면 시스템 수준 설계부터 시작해 플러그 앤드 플레이 IP를 가져와 최상위에서 연결하는 법을 배우고, 그 모듈을 미리 만들어진 기준 설계에 넣어보면 됨
      이후 점차 층을 벗기면서 아래의 복잡도를 드러내면 됨
    • 나도 같은 처지이고, 계획은 이렇다
      1. Harris, Harris의 Digital Design and Computer Architecture. (2022). Elsevier를 읽기: https://doi.org/10.1016/c2019-0-00213-0
      2. 저자의 RVFpga 과정을 따라가며 FPGA 위에 실제 RISC-V CPU를 만들기: https://www.youtube.com/watch?v=ePv3xD3ZmnY
    • 이런 경로를 추천함
      1. 교육용 저장소 https://github.com/yuri-panchul/basics-graphics-music를 클론하기. 처음부터 Verilog를 배우는 사람을 위한 간단한 실습 모음이고, GPU 개발을 위해 Imagination에서 일했던 Yuri Panchul이 작성했음
      2. 지원되는 수십 종의 FPGA 보드 중 하나와 키, LED 같은 액세서리를 구하기
      3. Yosys와 관련 도구를 설치하기
      4. lab01 DeMorgan부터 시작해 저장소의 실습을 최대한 많이 해보기
        Harris&Harris를 읽으면서 실습을 병행할 수 있음
        실습과 책을 끝내면 자기 프로젝트를 시작할 시점이 됨
        참고로 HackerMojo에서 주간 모임도 하고, Valley에 있지 않아도 Zoom으로 참여할 수 있음
    • 어느 단계에 있는지는 모르겠지만, 디지털 논리와 CPU/GPU 아키텍처를 더 잘 이해하는 데 이 자료들이 도움이 됐음
      1. https://learn.saylor.org/course/CS301
      2. https://www.coursera.org/learn/comparch
      3. https://hdlbits.01xz.net/wiki/Main_Page
    • LLM을 가속하고 싶다면 먼저 아키텍처를 알아야 함
      거기서 시작하면 됨
      하드웨어는 실제로 쉬운 부분이기도 하고, 제조 측면에서는 어려운 부분이기도 함
  • 여기 순차 always 블록에서 non-blocking 대입과 blocking 대입 연산자를 섞어 쓰는 이유가 있나?
    • 저건 지역 변수처럼 보임
    • 시뮬레이션과 합성 결과의 일치에 너무 집착하지 않는다면 그렇게 해도 됨
  • 오래전에 VHDL로 비슷한 걸 해본 적이 있음
    여러 오픈소스 HDL 프로젝트가 모여 있던 opencores라는 사이트가 있었음
    요즘 HPC 수준의 대규모 분산 HDL 시뮬레이터가 괜찮은 게 있는지 궁금함
    RTL 수준 시뮬레이션에 현대 GPU를 활용하는 건 타당해 보임
  • ALU가 DIV 명령어를 하드웨어 수준에서 그대로 구현한다고?
    현대 CUDA 코어 같은 데서 실제 명령어로 나눗셈을 갖는 게 보통인가, 아니면 보통은 소프트웨어로 에뮬레이션하나?
    실제 하드웨어 나눗셈 회로는 공간을 엄청 차지하니까 GPU ALU에 들어갈 거라고는 예상하지 못했음
    Verilog에서는 DIV: begin alu_out_reg <= rs / rt; end 한 줄 쓰기가 너무 쉽지만, 그 한 줄이 실리콘을 많이 잡아먹음
    단순히 Verilog만 시뮬레이션한다면 그 사실을 못 볼 수도 있음
    • 이건 그냥 누군가 Verilog를 배우는 프로젝트임
      프로젝트가 시뮬레이션에서 멈추고, 실제 하드웨어로 만들려면 훨씬 더 많은 작업이 필요함
  • 또 그래픽 기능이 없는 “GPU”임
    개인적으로 이런 건 다른 이름으로 불러야 한다고 봄
    • 첫 질문은 애초에 왜 CPU와 GPU가 나뉘었는가임
      둘 사이의 간격이 좁혀지고 있고 양쪽 모두 서로의 기능을 추가하고 있지만, 여전히 상당한 차이가 있음
      내 생각에는 Amdahl의 법칙과 관련이 있음 [0]
      그런 의미에서 CPU는 지연 시간 최적화 프로세서, GPU는 처리량 최적화 프로세서라고 부를 수 있음
      더 구체적으로는 [1] CPU를 길고 깊은 데이터 의존성 프로세서, GPU를 넓고 평평한 데이터 의존성 프로세서라고 부를 수도 있음
      [0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
      [1]: https://en.wikipedia.org/wiki/Data_dependency
    • TPU, 즉 텐서 처리 장치라고 부를 수 있음
      텐서는 그냥 n차원 배열임
      그 위에 소프트웨어나 펌웨어를 올려 GPU처럼 동작하게 만들 수 있음
    • ‘디스플레이 어댑터’를 만드는 프로젝트를 시작해볼까 생각해왔지만, 시작도 하기 전에 UEFI의 GOP 드라이버와 디스플레이 어댑터 사이의 통신 프로토콜을 파악하지 못해 막혔음
      EDK2 소스에서 조각을 맞춰보려 했지만, 어느 정도가 QEMU에 특화된 건지 불분명함
    • MPU, 즉 행렬 처리 장치라고 부르면 됨
    • 자리 잡아가는 용어는 AIA, 즉 AI 가속기라고 봄
  • tiny-gpu가 모든 스레드가 각 명령어 뒤에 같은 프로그램 카운터로 “수렴”한다고 가정한다는 건 너무 순진한 단순화임
    실제 GPU에서는 개별 스레드가 다른 PC로 분기할 수 있고, 처음엔 함께 처리되던 스레드 그룹이 별도 실행으로 갈라지는 분기 발산이 생김
    실리콘으로 GPU를 만들기 전에 GPU 프로그래밍부터 해보는 게 좋았을 듯함
    게다가 SIMD라고 하기에도 뭔가 맞지 않음
    이 사람은 예전에 남의 회로들을 이어 붙여 LED를 깜빡이게 해놓고 CPU를 만들었다고 했던 그 사람임
    • 첫 번째는 모든 실행마다 __syncthreads()를 호출하는 것과 같은 셈 아닌가?
답변달기