Tiny GPU: GPU tối giản được triển khai bằng Verilog

 (github.com/adam-maj)
2 điểm bởi GN⁺ 2024-04-27 | 1 bình luận | Chia sẻ qua WhatsApp
  • tiny-gpu là một triển khai GPU tối giản dựa trên Verilog để học từ nền tảng cách GPU hoạt động ở cấp phần cứng; nó tập trung vào các nguyên lý chung của GPGPU và bộ tăng tốc ML hơn là phần cứng chuyên dụng cho đồ họa
  • Phần triển khai gồm chưa đến 15 tệp Verilog có tài liệu, tài liệu kiến trúc và ISA, kernel cộng/nhân ma trận, cùng hỗ trợ mô phỏng kernel và truy vết thực thi
  • GPU chạy một kernel duy nhất tại một thời điểm; sau khi nạp bộ nhớ chương trình và bộ nhớ dữ liệu, đặt thread_count, rồi đưa tín hiệu start lên cao để bắt đầu kernel
  • Để đơn giản hóa, mỗi core xử lý một block tại một thời điểm; mỗi thread có ALU, LSU, PC và register file riêng, nhưng giả định rằng tất cả thread hội tụ về cùng PC sau mỗi lệnh
  • Hầu hết các tính năng như cache nhiều tầng, shared memory, memory coalescing, pipelining, warp scheduling, branch divergence và barrier của GPU hiện đại đều bị lược bỏ, ưu tiên cấu trúc phục vụ học tập

Vấn đề tiny-gpu muốn giải quyết

  • Có nhiều tài liệu để học về CPU, từ kiến trúc đến tín hiệu điều khiển, nhưng các chi tiết kỹ thuật cấp thấp của GPU hiện đại phần lớn vẫn là độc quyền do thị trường cạnh tranh
  • Có nhiều tài liệu lập trình GPU, nhưng hầu như không có tài liệu để học cách GPU hoạt động ở cấp phần cứng
  • Các triển khai GPU mã nguồn mở như Miaow và VeriGPU hướng tới tính hoàn chỉnh và khả năng vận hành nên có cấu trúc phức tạp
  • tiny-gpu loại bỏ nhiều độ phức tạp của card đồ họa cấp sản xuất và tập trung vào các thành phần cốt lõi phổ biến trong bộ tăng tốc phần cứng hiện đại
    • Các thành phần quan trọng của kiến trúc GPU
    • Cách mô hình lập trình SIMD được triển khai trong phần cứng
    • Cách GPU xử lý băng thông bộ nhớ hạn chế

Kiến trúc tổng thể

  • tiny-gpu được thiết kế để chỉ chạy một kernel tại một thời điểm
  • Quy trình chạy kernel như sau
    • Nạp mã kernel vào bộ nhớ chương trình toàn cục
    • Nạp dữ liệu cần thiết vào bộ nhớ dữ liệu
    • Chỉ định số thread sẽ chạy trong thanh ghi điều khiển thiết bị
    • Đặt tín hiệu start lên high để chạy kernel
  • GPU gồm các đơn vị sau
    • Thanh ghi điều khiển thiết bị
    • Dispatcher
    • Số lượng compute core có thể thay đổi
    • Memory controller cho bộ nhớ dữ liệu và bộ nhớ chương trình
    • Cache

Thực thi kernel và phân bổ thread

  • Thanh ghi điều khiển thiết bị có vai trò lưu metadata thực thi kernel; trong tiny-gpu, nó chỉ lưu thread_count, tức tổng số thread sẽ chạy
  • Dispatcher phân phối thread cho nhiều compute core khi kernel bắt đầu
    • Gom các nhóm thread có thể chạy song song thành block
    • Gửi block đến core khả dụng để xử lý
    • Thông báo hoàn tất thực thi kernel khi tất cả block đã xử lý xong
  • Core đã đơn giản hóa xử lý một block tại một thời điểm
  • Mỗi thread có ALU, LSU, PC và register file chuyên dụng
  • Việc quản lý thực thi lệnh của thread trên các tài nguyên này là một trong những bài toán khó của GPU

Cấu trúc bộ nhớ và controller

  • GPU được tạo để giao tiếp với bộ nhớ toàn cục bên ngoài; để đơn giản hóa, bộ nhớ dữ liệu và bộ nhớ chương trình được tách riêng
  • Đặc tả bộ nhớ dữ liệu
    • Địa chỉ hóa 8 bit
    • Tổng cộng 256 hàng
    • Dữ liệu 8 bit
    • Mỗi hàng lưu giá trị nhỏ hơn 256
  • Đặc tả bộ nhớ chương trình
    • Địa chỉ hóa 8 bit
    • Tổng cộng 256 hàng
    • Dữ liệu 16 bit
    • Theo ISA, mỗi lệnh là 16 bit
  • Memory controller theo dõi các yêu cầu bộ nhớ đến từ core, giới hạn yêu cầu theo băng thông bộ nhớ ngoài thực tế và chuyển phản hồi đến đúng tài nguyên
  • Mỗi memory controller có số kênh cố định tùy theo băng thông bộ nhớ toàn cục
  • Cache là tính năng đang được phát triển; nó lưu dữ liệu lấy từ bộ nhớ ngoài vào SRAM trong thiết bị để các yêu cầu sau có thể lấy nhanh hơn và băng thông bộ nhớ có thể được dùng cho dữ liệu mới

Cấu trúc bên trong core

  • Mỗi core có một scheduler duy nhất để quản lý thực thi thread
  • Scheduler của tiny-gpu thực thi lệnh của một block đến hết rồi lấy block mới, đồng thời chạy lệnh của tất cả thread theo thứ tự đồng bộ
  • Với scheduler nâng cao hơn, có thể tăng mức sử dụng tài nguyên bằng pipelining và warp scheduling
  • Ràng buộc chính của scheduler phát sinh từ độ trễ khi load và store dữ liệu từ bộ nhớ toàn cục
    • Hầu hết lệnh có thể chạy đồng bộ
    • Các thao tác load-store như LDR và STR là bất đồng bộ, nên việc thực thi lệnh phải được tổ chức quanh thời gian chờ dài
  • Fetcher lấy bất đồng bộ lệnh tại program counter hiện tại từ bộ nhớ chương trình
  • Decoder giải mã lệnh đã lấy thành tín hiệu điều khiển để thực thi thread
  • Register file của mỗi thread lưu dữ liệu đang được tính toán và cho phép mẫu SIMD
    • Các thanh ghi chỉ đọc chứa %blockIdx, %blockDim, %threadIdx
    • Kernel có thể chạy với dữ liệu khác nhau tùy theo ID thread cục bộ
  • ALU của mỗi thread xử lý các lệnh số học ADD, SUB, MUL, DIV
  • CMP xuất ra kết quả cho biết hiệu của hai thanh ghi là âm, 0 hay dương, rồi lưu kết quả vào thanh ghi NZP của đơn vị PC
  • LSU của mỗi thread truy cập bộ nhớ dữ liệu toàn cục, xử lý LDR, STR và độ trễ bộ nhớ bất đồng bộ
  • PC của mỗi thread quyết định lệnh tiếp theo sẽ thực thi
    • Mặc định tăng 1 sau mỗi lệnh
    • BRnzp rẽ nhánh đến một hàng bộ nhớ chương trình cụ thể nếu điều kiện trong thanh ghi NZP do CMP trước đó thiết lập là đúng
    • Vòng lặp và câu lệnh điều kiện được triển khai theo cách này
  • Để đơn giản hóa, tiny-gpu giả định tất cả thread hội tụ về cùng PC sau mỗi lệnh
  • Trong GPU thực tế, từng thread có thể rẽ nhánh đến các PC khác nhau; khi đó nhóm thread đang được xử lý cùng nhau sẽ tách thành nhiều luồng thực thi, gọi là branch divergence

ISA

  • tiny-gpu triển khai ISA 11 lệnh để chạy các kernel đơn giản dạng proof-of-concept như cộng ma trận và nhân ma trận
  • Các lệnh được hỗ trợ
    • BRnzp: nhảy đến hàng khác trong bộ nhớ chương trình nếu điều kiện NZP khớp
    • CMP: so sánh giá trị của hai thanh ghi và lưu kết quả vào thanh ghi NZP
    • ADD, SUB, MUL, DIV: các phép toán số học cơ bản cho tensor math
    • LDR: load dữ liệu từ bộ nhớ toàn cục
    • STR: store dữ liệu vào bộ nhớ toàn cục
    • CONST: load giá trị hằng vào thanh ghi
    • RET: tín hiệu kết thúc thực thi thread hiện tại
  • Mỗi thanh ghi được chỉ định bằng 4 bit, tổng cộng có 16 thanh ghi
    • 13 thanh ghi từ R0 đến R12 là thanh ghi tự do có thể đọc/ghi
    • 3 thanh ghi cuối là thanh ghi đặc biệt chỉ đọc, cung cấp %blockIdx, %blockDim, %threadIdx cần thiết cho SIMD

Luồng thực thi

  • Khi thực thi lệnh, mỗi core tuân theo luồng điều khiển gồm các bước sau
    • FETCH: lấy lệnh tiếp theo tại PC hiện tại
    • DECODE: giải mã lệnh thành tín hiệu điều khiển
    • REQUEST: nếu cần LDR hoặc STR, gửi yêu cầu dữ liệu đến bộ nhớ toàn cục
    • WAIT: nếu cần, chờ phản hồi từ bộ nhớ toàn cục
    • EXECUTE: thực hiện tính toán trên dữ liệu
    • UPDATE: cập nhật register file và thanh ghi NZP
  • Luồng điều khiển này được xây dựng nhằm đơn giản và dễ hiểu
  • Trong triển khai thực tế, có thể gộp một số bước để tối ưu thời gian xử lý, hoặc điều phối nhiều lệnh trên tài nguyên core bằng pipelining
  • Mỗi thread tính toán trên dữ liệu trong register file chuyên dụng của mình theo cùng một đường thực thi
  • Nó giống sơ đồ CPU, nhưng khác ở chỗ %blockIdx, %blockDim, %threadIdx nằm trong các thanh ghi chỉ đọc, cho phép chức năng SIMD

Kernel ví dụ

  • Các kernel cộng ma trận và nhân ma trận được viết để làm proof-of-concept cho ISA
  • Các tệp kiểm thử trong repository có thể mô phỏng hoàn chỉnh các kernel này trên GPU, đồng thời tạo trạng thái bộ nhớ dữ liệu và toàn bộ truy vết thực thi

  • Cộng ma trận

    • matadd.asm cộng hai ma trận 1 x 8
    • Phép cộng theo từng phần tử của 8 phần tử được thực hiện trên các thread riêng biệt
    • Sử dụng các thanh ghi %blockIdx, %blockDim, %threadIdx để minh họa lập trình SIMD
    • Bao gồm quản lý bộ nhớ bất đồng bộ bằng các lệnh LDR và STR

  • Nhân ma trận

    • matmul.asm nhân hai ma trận 2 x 2
    • Tính tích vô hướng của hàng và cột tương ứng theo từng phần tử
    • Dùng CMP và BRnzp để minh họa rẽ nhánh bên trong thread
    • Vì tất cả nhánh đều hội tụ lại, nó hoạt động trong triển khai tiny-gpu hiện tại

Mô phỏng

  • Để chạy mô phỏng kernel, cần iverilog và cocotb
  • Các bước chuẩn bị
    • Cài trình biên dịch Verilog và cocotb bằng brew install icarus-verilog và pip3 install cocotb
    • Tải phiên bản mới nhất của sv2v, giải nén và thêm binary vào $PATH
    • Chạy mkdir build tại root của repository
  • Chạy mô phỏng kernel bằng make test_matadd và make test_matmul
  • Kết quả thực thi được xuất ra các tệp log trong test/logs
    • Trạng thái bộ nhớ dữ liệu ban đầu
    • Toàn bộ truy vết thực thi của kernel
    • Trạng thái bộ nhớ dữ liệu cuối cùng
  • Ở phần đầu mỗi tệp log có ma trận đầu vào, và ở bộ nhớ dữ liệu cuối cùng ở phần cuối có ma trận kết quả
  • Truy vết thực thi bao gồm trạng thái thực thi của tất cả thread trên mọi core ở từng chu kỳ
    • Lệnh hiện tại
    • PC
    • Giá trị thanh ghi
    • Thông tin trạng thái

Các tính năng GPU nâng cao cố ý lược bỏ

  • tiny-gpu lược bỏ hầu hết các yếu tố cải thiện hiệu năng và tính năng của GPU hiện đại để đơn giản hóa

  • Cache nhiều tầng và shared memory

    • GPU hiện đại dùng nhiều tầng cache để giảm truy cập bộ nhớ toàn cục
    • tiny-gpu chỉ triển khai một tầng cache duy nhất để lưu dữ liệu gần đây giữa tài nguyên yêu cầu và memory controller
    • Cache nhiều tầng giảm thời gian load bằng cách cache dữ liệu thường dùng gần nơi sử dụng hơn
    • GPU cũng có thể dùng shared memory để các thread trong cùng block trao đổi kết quả chung

  • Memory coalescing

    • Nhiều thread đang chạy song song thường truy cập các địa chỉ liên tiếp, chẳng hạn các phần tử liền kề của ma trận
    • Memory coalescing phân tích các yêu cầu bộ nhớ đang xếp hàng và gộp các yêu cầu liền kề thành một transaction
    • Mục đích là giảm thời gian dành cho địa chỉ hóa và xử lý các yêu cầu cùng nhau

  • Pipelining

    • Core của tiny-gpu chỉ bắt đầu lệnh tiếp theo sau khi thực thi xong một lệnh của một nhóm thread
    • GPU hiện đại stream việc thực thi nhiều lệnh tuần tự, trong khi vẫn đảm bảo lệnh có phụ thuộc được chạy theo thứ tự
    • Điều này tăng mức sử dụng tài nguyên để tài nguyên core không bị nhàn rỗi trong các tình huống như chờ yêu cầu bộ nhớ bất đồng bộ

  • Warp scheduling

    • Chia block thành các warp, tức các batch thread có thể chạy cùng nhau
    • Khi một warp đang chờ, core thực thi lệnh của warp khác, qua đó xử lý nhiều warp đồng thời trên một core
    • Tương tự pipelining nhưng xử lý lệnh của các thread khác nhau

  • Branch divergence

    • tiny-gpu giả định tất cả thread trong một batch đều ở cùng PC sau mỗi lệnh
    • Trên thực tế, từng thread có thể rẽ nhánh đến dòng khác tùy theo dữ liệu
    • Các thread có PC khác nhau sẽ được tách thành các luồng thực thi riêng, và cũng cần quản lý thời điểm chúng hội tụ lại

  • Đồng bộ hóa và barrier

    • GPU hiện đại có thể đặt barrier để buộc nhóm thread trong cùng block chờ cho đến khi tất cả đều tới một điểm nhất định
    • Điều này hữu ích khi các thread cần trao đổi dữ liệu dùng chung, nhằm đảm bảo việc xử lý dữ liệu đã hoàn tất

Công việc tiếp theo

  • Các hạng mục cải tiến trong tương lai như sau
    • Thêm cache lệnh đơn giản
    • Xây dựng adapter để có thể dùng GPU trong Tiny Tapeout 7
    • Thêm branch divergence cơ bản
    • Thêm memory coalescing cơ bản
    • Thêm pipelining cơ bản
    • Tối ưu luồng điều khiển và cách dùng thanh ghi để cải thiện cycle time
    • Viết kernel đồ họa cơ bản hoặc thêm phần cứng đồ họa đơn giản để minh họa chức năng đồ họa
  • Người dùng muốn cải thiện repository có thể đóng góp bằng PR

Bài viết liên quan

1 bình luận

 
GN⁺ 2024-04-27
Ý kiến trên Hacker News
  • Thị trường GPU vốn cạnh tranh khốc liệt nên phần lớn chi tiết kỹ thuật cấp thấp của các kiến trúc hiện đại vẫn không được công khai
    Ngoại lệ là Intel đã công bố khá nhiều tài liệu kỹ thuật về GPU: https://kiwitree.net/~lina/intel-gfx-docs/prm/
    Sổ tay i810/815 cũng có thể tìm thấy trên mạng, và nếu bỏ qua khoảng trống kỳ lạ ở giai đoạn trước 965 khi thiếu 855/910/915/945, thì việc tài liệu hóa nhìn chung khá đều đặn
    • AMD cũng công bố khá nhiều tài liệu: https://www.amd.com/en/developer/browse-by-resource-type/documentation.html
      Bao gồm cả tài liệu kiến trúc tập lệnh cho các sản phẩm hiện tại và trước đây, nhưng có vẻ giống tài liệu nhắm đến người triển khai hơn là phần giải thích cấp cao cho những người đam mê quan tâm
    • Driver Linux của Intel cũng có chất lượng tốt và đã nằm trong mainline
      Ước gì mọi công ty đều làm theo cách này
    • Tài liệu từ năm 2018 nhưng vẫn có liên quan ở mức nào đó: The Thirty Million Line Problem - Casey Muratori
  • Dự án thật sự rất hay, và thật tốt khi thấy những dự án phần cứng như thế này được phát triển công khai
    Tuy nhiên tôi thấy nó gần với bộ đồng xử lý SIMD hơn
    Tôi nghĩ để gọi là GPU thì tối thiểu phải có một dạng xuất hiển thị nào đó
    Tôi biết thuật ngữ này đã trở nên khá lỏng lẻo khi gần đây Nvidia và các hãng khác cũng bán các biến thể kiến trúc đồ họa chỉ dành cho máy chủ dưới tên GPU, nhưng trong thiết kế GPU, phần đồ họa đến nay vẫn chiếm một phần đáng kể độ phức tạp
    • Nếu nó xử lý đồ họa thì tôi nghĩ vẫn có thể xem là GPU dù không có đầu ra
      GPU không xuất hình vẫn hữu ích
      Ở chỗ tôi làm có khoảng 75 workstation gắn Quadro tầm trung, nhưng card chỉ có mini-DisplayPort còn công ty chỉ mua cáp HDMI, nên tất cả đều đang cắm vào đồ họa tích hợp
      Dù vậy các card đó vẫn tăng tốc phần mềm và xử lý đồ họa, chỉ là không xuất ra màn hình mà thôi
  • Hay. Tôi rất ủng hộ công việc về GPU lõi mở
    Cũng có ví dụ khác: https://github.com/jbush001/NyuziProcessor
    • Sẽ rất tuyệt nếu có một triển khai CUDA tối thiểu cho một trong các bộ xử lý lõi mở như thế này
      Cần sản lượng ở mức nào để TSMC hay một foundry khác có thể sản xuất các bộ xử lý kiểu này một cách kinh tế?
  • Dự án thật sự xuất sắc
    Tôi muốn thử FPGA, nhưng nói thật là ngay cả việc biết bắt đầu từ đâu cũng khó, và cả lĩnh vực này tạo cảm giác khá đáng ngại
    Mục tiêu cuối cùng là làm một card tăng tốc cho LLM; dù đó là mục tiêu hoàn toàn tự đặt ra, tôi nghĩ nó sẽ có nhiều phần trùng với dự án này, và có lẽ chỉ khác ở phần offload bộ nhớ để tải các mô hình lớn hơn
    • Cần thay đổi khung suy nghĩ trong đầu
      Nhập môn FPGA phải được chia nhỏ thành nhiều kỹ năng con, và cũng cần điều chỉnh kỳ vọng
      Ta không kỳ vọng một kỹ sư phần mềm ngay từ đầu đã xây dựng cả máy tính từ nguyên lý, viết kiến trúc tập lệnh, hiểu mã máy, chuyển nó thành assembly, rồi phát triển cả ngôn ngữ lập trình để tạo ứng dụng bằng mã Python
      Cách đúng là bắt đầu từ phía trên rồi đi xuống dưới của stack
      Nếu tập trung vào việc trừu tượng hóa độ phức tạp và xây dựng hệ thống bằng các IP có sẵn, thiết kế FPGA khá dễ
      Thường thì tôi sẽ khuyên dùng thứ như MATLAB, vì có thể tạo ứng dụng ban đầu bằng HDL Coder trên DevKit có reference design
      Nếu không, bạn sẽ gánh một khối lượng học khổng lồ: kiến trúc tính toán số, Verilog, timing, transceiver/I/O, lập kế hoạch chân, Quartus/Vivado, mô phỏng/xác minh, hệ thống nhúng, v.v.
      Tóm lại, hãy bắt đầu từ thiết kế cấp hệ thống, học cách lấy các IP kiểu cắm là chạy và kết nối chúng ở tầng cao nhất, rồi đưa mô-đun đó vào một thiết kế tham chiếu dựng sẵn
      Sau đó có thể dần bóc từng lớp để lộ độ phức tạp bên dưới
    • Tôi cũng ở hoàn cảnh tương tự, và kế hoạch của tôi là thế này
      1. Đọc Digital Design and Computer Architecture của Harris, Harris. (2022). Elsevier: https://doi.org/10.1016/c2019-0-00213-0
      2. Theo khóa RVFpga của tác giả để tạo một CPU RISC-V thực sự trên FPGA: https://www.youtube.com/watch?v=ePv3xD3ZmnY
    • Tôi khuyên đi theo lộ trình này
      1. Clone kho lưu trữ giáo dục https://github.com/yuri-panchul/basics-graphics-music. Đây là tập hợp các bài thực hành đơn giản cho người học Verilog từ con số 0, do Yuri Panchul, người từng làm tại Imagination để phát triển GPU, viết
      2. Kiếm một trong hàng chục board FPGA được hỗ trợ cùng các phụ kiện như nút bấm, LED
      3. Cài Yosys và các công cụ liên quan
      4. Bắt đầu từ lab01 DeMorgan và làm càng nhiều bài thực hành trong kho càng tốt
        Có thể vừa đọc Harris&Harris vừa làm thực hành
        Khi hoàn thành các bài thực hành và cuốn sách, đó là lúc bắt đầu dự án của riêng mình
        Nhân tiện, HackerMojo cũng có buổi gặp hằng tuần, và dù không ở Valley bạn vẫn có thể tham gia qua Zoom
    • Tôi không biết bạn đang ở giai đoạn nào, nhưng các tài liệu này đã giúp tôi hiểu rõ hơn về logic số và kiến trúc CPU/GPU
      1. https://learn.saylor.org/course/CS301
      2. https://www.coursera.org/learn/comparch
  1. https://hdlbits.01xz.net/wiki/Main_Page
  • Nếu muốn tăng tốc LLM thì trước hết cần hiểu kiến trúc
    Bắt đầu từ đó là được
    Phần cứng thực ra vừa là phần dễ, vừa là phần khó nếu xét về mặt chế tạo
  • Có lý do gì để trộn toán tử gán non-blocking và gán blocking trong khối always tuần tự ở đây không?
    • Cái đó trông giống biến cục bộ
    • Nếu không quá ám ảnh với việc kết quả mô phỏng và tổng hợp phải khớp nhau thì làm vậy cũng được
  • Lâu rồi tôi từng làm một thứ tương tự bằng VHDL
    Khi đó có một trang tên là opencores, tập hợp nhiều dự án HDL mã nguồn mở
    Không biết ngày nay có trình mô phỏng HDL phân tán quy mô lớn cấp HPC nào dùng ổn không
    Việc tận dụng GPU hiện đại cho mô phỏng ở mức RTL nghe có vẻ hợp lý
  • ALU mà triển khai nguyên lệnh DIV ở mức phần cứng sao?
    Ở những thứ như CUDA core hiện đại, việc có phép chia như một lệnh thật sự có phổ biến không, hay thường được mô phỏng bằng phần mềm?
    Mạch chia phần cứng thực sự chiếm rất nhiều diện tích, nên tôi không ngờ nó lại nằm trong ALU của GPU
    Trong Verilog thì viết một dòng DIV: begin alu_out_reg <= rs / rt; end quá dễ, nhưng một dòng đó ngốn rất nhiều silicon
    Nếu chỉ mô phỏng Verilog thì có thể không thấy được điều đó
    • Đây chỉ là một dự án để ai đó học Verilog
      Dự án dừng ở mức mô phỏng; để biến nó thành phần cứng thật thì cần nhiều việc hơn rất nhiều
  • Lại là một “GPU” không có chức năng đồ họa
    Cá nhân tôi nghĩ những thứ như thế nên được gọi bằng tên khác
    • Câu hỏi đầu tiên ngay từ đầu là vì sao CPU và GPU lại tách ra
      Khoảng cách giữa hai bên đang thu hẹp và cả hai đều bổ sung tính năng của nhau, nhưng vẫn có khác biệt đáng kể
      Theo tôi, điều này liên quan đến định luật Amdahl [0]
      Theo nghĩa đó, CPU có thể được gọi là bộ xử lý tối ưu cho độ trễ, còn GPU là bộ xử lý tối ưu cho thông lượng
      Cụ thể hơn, [1] cũng có thể gọi CPU là bộ xử lý cho phụ thuộc dữ liệu dài và sâu, còn GPU là bộ xử lý cho phụ thuộc dữ liệu rộng và phẳng
      [0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
      [1]: https://en.wikipedia.org/wiki/Data_dependency
    • Có thể gọi là TPU, tức bộ xử lý tensor
      Tensor chỉ là mảng n chiều
      Có thể đặt phần mềm hoặc firmware lên đó để khiến nó hoạt động giống GPU
    • Tôi từng nghĩ đến việc bắt đầu một dự án tạo “bộ chuyển đổi hiển thị”, nhưng còn chưa bắt đầu đã bị kẹt vì không nắm được giao thức giao tiếp giữa driver GOP của UEFI và bộ chuyển đổi hiển thị
      Tôi đã cố ghép các mảnh từ mã nguồn EDK2, nhưng không rõ phần nào là đặc thù của QEMU
    • Gọi là MPU, tức bộ xử lý ma trận, là được
    • Tôi nghĩ thuật ngữ đang dần ổn định là AIA, tức bộ tăng tốc AI
  • Việc tiny-gpu giả định mọi luồng “hội tụ” về cùng một bộ đếm chương trình sau mỗi lệnh là một đơn giản hóa quá ngây thơ
    Trong GPU thực tế, từng luồng có thể rẽ nhánh sang PC khác, và xuất hiện phân kỳ nhánh, nơi một nhóm luồng ban đầu được xử lý cùng nhau tách ra thành các luồng thực thi riêng
    Lẽ ra nên thử lập trình GPU trước khi làm GPU bằng silicon
    Hơn nữa, gọi là SIMD cũng có vẻ không hẳn đúng
    Đây chính là người trước đây từng nối các mạch của người khác lại để làm LED nhấp nháy rồi nói là mình đã tạo CPU
    • Điều đầu tiên chẳng phải tương đương với việc gọi __syncthreads() sau mỗi lần thực thi sao?