จากศูนย์ถึงหลาย GPU: คลังสอนโอเพนซอร์สของ NVIDIA สอนคุณเขียนโปรแกรม CUDA อย่างเป็นระบบ พร้อมเริ่มต้นใช้งานด้วยการปรับใช้แบบคอนเทนเนอร์เพียงคลิกเดียว

การเขียนโปรแกรม GPU ถูกมองว่าเป็นอุปสรรคทางเทคนิคที่สูงชันมานานแล้ว ผู้เริ่มต้นไม่เพียงแต่ต้องเข้าใจแนวคิดเชิงนามธรรม เช่น โครงสร้างลำดับชั้นของเธรดและโมเดลพื้นที่หน่วยความจำ GPU เท่านั้น แต่ยังต้องต่อสู้กับเครื่องมือคอมไพล์และสภาพแวดล้อมรันไทม์อีกด้วย

คลังโค้ด accelerated-computing-hub ที่ NVIDIA เปิดตัวกำลังพยายามเชื่อมช่องว่างนี้ มันไม่ใช่แค่ชุดตัวอย่างโค้ดที่กระจัดกระจาย แต่เป็นระบบการสอนที่สมบูรณ์ซึ่งผ่านการขัดเกลาอย่างพิถีพิถัน ตั้งแต่ “Hello GPU” ไปจนถึงการคำนวณแบบกระจายบน GPU หลายตัว สิ่งที่น่าสนใจเป็นพิเศษคือ โค้ดของมันสะท้อนถึงแนวคิด “วิศวกรรมซอฟต์แวร์ที่ขับเคลื่อนด้วยการสอน” โดยแต่ละไฟล์ต้นฉบับมุ่งเป้าไปที่เป้าหมายทางปัญญาที่เฉพาะเจาะจง ความซับซ้อนของโค้ดจะเพิ่มขึ้นอย่างเคร่งครัด และใช้ Docker Container เพื่อสร้างสภาพแวดล้อมรันไทม์ที่ไม่ต้องกำหนดค่าใดๆ

• ชุดรวมทรัพยากรการศึกษาที่เกี่ยวข้องกับการเขียนโปรแกรม GPU ทั่วไปที่คัดสรรโดย NVIDIA
• https://github.com/NVIDIA/accelerated-computing-hub
• 4000 คำ อ่าน 20 นาที, พอดแคสต์ 16 นาที

คลังนี้มีโครงร่างบทเรียน ซึ่งรวบรวมบทเรียนการเขียนโปรแกรม GPU แบบโต้ตอบทั้งหมดที่สามารถรันบน NVIDIA Brev และ Google Colab จาก NVIDIA Accelerated Computing Center

ตารางด้านล่างแสดงรายการหัวข้อเนื้อหาบทเรียน ไฟล์กำหนดค่า Docker Compose ที่เกี่ยวข้อง รุ่นอินสแตนซ์ฮาร์ดแวร์ Brev ที่รองรับ และ ผู้ให้บริการ Brev และข้อกำหนดสภาพแวดล้อมระบบ ที่รองรับการปรับใช้และการทำงาน เพื่อให้นักพัฒนาสามารถจับคู่ฮาร์ดแวร์และแพลตฟอร์มสำหรับการเรียนรู้เชิงปฏิบัติได้อย่างรวดเร็ว โดยที่ NVIDIA Brev คือแพลตฟอร์มคลาวด์สำหรับนักพัฒนาของ NVIDIA ที่มุ่งเน้นการพัฒนาและส่งมอบต้นแบบ AI

บทความนี้จะเจาะลึกคลังนี้ โดยเน้นที่รายละเอียดการออกแบบของโค้ดการสอนหลัก และวิธีที่คลังนี้เปิดประตูสู่การเร่งความเร็ว GPU สำหรับนักพัฒนา HPC แบบดั้งเดิมผ่านเส้นทาง “Zero CUDA Code” ที่เรียกว่า Standard Parallelism

unsetunsetสารบัญunsetunset

• เริ่มต้นใช้งานอย่างรวดเร็ว
• หนึ่ง โครงสร้างโดยรวมของคลังและตรรกะการจัดเรียงหลักสูตร
• 1.1 เส้นทางการเรียนรู้แบบขนานสามเส้นทาง
• 1.2 การก้าวหน้าสามขั้นของหลักสูตร CUDA C++
• สอง จากเธรดเดียวไปจนถึงหลายเธรด: การออกแบบโค้ดแบบค่อยเป็นค่อยไปสำหรับการเขียนโปรแกรม Kernel
• 2.1 จุดเริ่มต้น: GPU kernel แบบเธรดเดียว
• 2.2 การแนะนำความขนานของเธรด: รูปแบบ Stride Loop
• 2.3 การสาธิตไวยากรณ์การเปิดตัว Kernel แบบน้อยที่สุด
• สาม หน่วยความจำที่ใช้ร่วมกันและการซิงโครไนซ์เธรด: วงจรปิดที่สมบูรณ์ตั้งแต่หลักการจนถึงการเพิ่มประสิทธิภาพ
• 3.1 ตัวอย่างที่รันได้น้อยที่สุดของหน่วยความจำที่ใช้ร่วมกัน
• 3.2 การปฏิบัติจริง: การเพิ่มประสิทธิภาพฮิสโตแกรมโดยใช้หน่วยความจำที่ใช้ร่วมกัน
• 3.3 อัลกอริทึมการทำงานร่วมกันของ CUB: ยืนบนไหล่ของยักษ์
• สี่ Standard Parallelism: เส้นทางสู่การเร่งความเร็ว GPU โดยไม่มีโค้ด CUDA
• 4.1 การรันอัลกอริทึม C++ มาตรฐานบน GPU โดยตรง
• 4.2 สมการการนำความร้อนสองมิติ: การประยุกต์ใช้ stdpar ในทางปฏิบัติสำหรับ HPC
• ห้า โครงสร้างพื้นฐานแบบคอนเทนเนอร์: การรับประกันทางวิศวกรรมสำหรับสภาพแวดล้อมการสอน
• 5.1 การจัดระเบียบบริการหลายรายการของ Docker Compose
• สรุป

unsetunsetเริ่มต้นใช้งานอย่างรวดเร็วunsetunset

โคลนคลังและเริ่มต้นสภาพแวดล้อม JupyterLab สำหรับบทเรียน CUDA C++ ด้วยคำสั่ง Docker Compose เพียงคำสั่งเดียว:

git clone https://github.com/NVIDIA/accelerated-computing-hub.git
cd accelerated-computing-hub/tutorials/cuda-cpp/brev
docker compose up

หลังจากเริ่มต้น ให้เข้าไปที่ http://127.0.0.1:8888 เพื่อเข้าสู่ Notebook แบบโต้ตอบ นอกจากนี้ยังสามารถรันออนไลน์ผ่าน Google Colab ได้โดยตรง โดยไม่ต้องใช้ GPU ในเครื่อง ดูวิธีการปรับใช้เพิ่มเติม (อินสแตนซ์คลาวด์ NVIDIA Brev, การเลือกรุ่น GPU ฯลฯ) ได้ที่ docs/brev.md

unsetunsetหนึ่ง โครงสร้างโดยรวมของคลังและตรรกะการจัดเรียงหลักสูตรunsetunset

1.1 เส้นทางการเรียนรู้แบบขนานสามเส้นทาง

คลังทั้งหมดจัดระเบียบรอบโมดูลการสอนหลักสามโมดูล ซึ่งสอดคล้องกับกลุ่มเป้าหมายสามกลุ่ม:

ทั้งสามเส้นทางเป็นอิสระต่อกัน แต่ใช้โครงสร้างพื้นฐานเดียวกัน (อิมเมจ Docker, การกำหนดค่า Brev, กรอบงานทดสอบ) การออกแบบนี้ช่วยให้นักพัฒนาที่มีพื้นฐานแตกต่างกันสามารถหาจุดเริ่มต้นที่เหมาะสมกับโครงสร้างความรู้ของตนเองมากที่สุด

1.2 การก้าวหน้าสามขั้นของหลักสูตร CUDA C++

บทเรียน CUDA C++ แบ่งออกเป็นสามขั้นตอนที่ออกแบบมาอย่างพิถีพิถัน:

ขั้นตอนที่หนึ่ง (01.xx): ความรู้เบื้องต้นเกี่ยวกับอัลกอริทึมแบบขนาน – ไม่จำเป็นต้องเขียน Kernel ด้วยตนเอง ใช้สิ่งที่เป็นนามธรรมระดับสูง เช่น thrust/CCCL เพื่อทำความเข้าใจแนวคิดของพื้นที่การดำเนินการและพื้นที่หน่วยความจำ

ขั้นตอนที่สอง (02.xx): Asynchronous และ Stream – เจาะลึกถึงธรรมชาติของการดำเนินการแบบอะซิงโครนัสของ GPU, CUDA Stream และการใช้ Pinned Memory

ขั้นตอนที่สาม (03.xx): Kernel ที่กำหนดเอง – เริ่มต้นจาก Kernel แบบเธรดเดียวพื้นฐานที่สุด ค่อยๆ พัฒนาไปสู่การเพิ่มประสิทธิภาพด้วยหน่วยความจำที่ใช้ร่วมกันและอัลกอริทึมการทำงานร่วมกัน

แนวคิดการสอนแบบ “เรียนรู้การใช้ก่อน เข้าใจหลักการทีหลัง และสุดท้ายเจาะลึกการเพิ่มประสิทธิภาพ” นี้ สะท้อนให้เห็นอย่างชัดเจนในระดับโค้ด

สอง จากเธรดเดียวไปจนถึงหลายเธรด: การออกแบบโค้ดแบบค่อยเป็นค่อยไปสำหรับการเขียนโปรแกรม Kernel

2.1 จุดเริ่มต้น: GPU kernel แบบเธรดเดียว

ส่วน Kernel ของบทเรียนเริ่มต้นจากการใช้งานที่ “ไม่มีประสิทธิภาพ” โดยเจตนา เพื่อให้ผู้เรียนเข้าใจไวยากรณ์พื้นฐานก่อน:

// ที่มา: tutorials/cuda-cpp/notebooks/03.02-Kernels/Sources/simple-kernel.cpp  
#include "ach.h"  

__global__ void single_thread_kernel(ach::temperature_grid_f in, float *out) {  
for (int id = 0; id < in.size(); id++) {  
out[id] = ach::compute(id, in);  
}  
}  

void simulate(ach::temperature_grid_f temp_in, float *temp_out,  
cudaStream_t stream) {  
single_thread_kernel<<<1, 1, 0, stream>>>(temp_in, temp_out);  
}  

สังเกต <<<1, 1>>> ที่นี่ – มันเปิดใช้งานเพียงหนึ่งเธรดภายในหนึ่ง Block จุดประสงค์การสอนของโค้ดนี้ชัดเจนมาก: ให้ผู้เรียนเข้าใจโครงสร้างไวยากรณ์และวิธีการเปิดตัว Kernel ก่อน โดยไม่ต้องกังวลกับปัญหาความขนานในตอนนี้

• ach::temperature_grid_f เป็นประเภทข้อมูลกริดอุณหภูมิที่ห่อหุ้มไว้ในบทเรียน
• ach::compute ห่อหุ้มตรรกะการคำนวณ Stencil เพื่อให้ผู้เรียนสามารถมุ่งเน้นไปที่การปรับเปลี่ยนให้ขนานได้เอง

2.2 การแนะนำความขนานของเธรด: รูปแบบ Stride Loop

จากนั้น โค้ดจะพัฒนาไปสู่เวอร์ชันหลายเธรด ซึ่งแสดงรูปแบบ Stride Loop แบบคลาสสิก:

// ที่มา: tutorials/cuda-cpp/notebooks/03.02-Kernels/Sources/block-256-kernel.cpp  
#include "ach.h"  

const int number_of_threads = 256;  

__global__ void block_kernel(ach::temperature_grid_f in, float *out) {  
int thread_index = threadIdx.x;  

for (int id = thread_index; id < in.size(); id += number_of_threads) {  
out[id] = ach::compute(id, in);  
}  
}  

void simulate(ach::temperature_grid_f temp_in, float *temp_out,  
cudaStream_t stream) {  
block_kernel<<<1, number_of_threads, 0, stream>>>(temp_in, temp_out);  
}  

นี่แสดงรูปแบบ stride loop (การวนซ้ำแบบข้ามขั้น) แบบคลาสสิก: แต่ละเธรดเริ่มต้นจาก threadIdx.x ของตัวเอง และวนซ้ำข้อมูลทั้งหมดด้วยจำนวนเธรดทั้งหมดเป็นขั้นตอน บทเรียนยังมี “เวอร์ชันที่ผิดพลาด” (failed-block-kernel.cpp) ที่มี number_of_threads = 2048 เพื่อให้ผู้เรียนได้สัมผัสกับข้อผิดพลาดรันไทม์ที่เกิดขึ้นเมื่อเกินขีดจำกัดสูงสุดของฮาร์ดแวร์ (1024) วิธีการสอนแบบ “ทำผิดก่อน แล้วค่อยแก้ไข” นี้มีอยู่ทั่วทั้งคลัง

2.3 การสาธิตไวยากรณ์การเปิดตัว Kernel แบบน้อยที่สุด

เพื่ออธิบายไวยากรณ์เฉพาะของ CUDA <<<blocks, threads, shmem, stream>>> แยกต่างหาก บทเรียนจึงมี Kernel “Hello World” ที่เรียบง่ายมาก:

// ที่มา: tutorials/cuda-cpp/notebooks/03.02-Kernels/Sources/kernel-launch.cpp  
#include <cstdio>  

__global__ void kernel(int value) {  
std::printf("value on device = %dn", value);  
}  

int main() {  
int blocks_in_grid = 1;  
int threads_in_block = 1;  
cudaStream_t stream = 0;  
kernel<<<blocks_in_grid, threads_in_block, 0, stream>>>(42);  
cudaStreamSynchronize(stream);  
}  

พารามิเตอร์การเปิดตัวแต่ละตัวถูกกำหนดชื่อตัวแปรเฉพาะ ทำให้เอกสารในตัวเองมีระดับสูงมาก ซึ่งชัดเจนกว่าการแสดง kernel<<<1,1>>>(42) โดยตรงมาก – ชื่อตัวแปรคือคำอธิบายที่ดีที่สุด

สาม หน่วยความจำที่ใช้ร่วมกันและการซิงโครไนซ์เธรด: วงจรปิดที่สมบูรณ์ตั้งแต่หลักการจนถึงการเพิ่มประสิทธิภาพ

3.1 ตัวอย่างที่รันได้น้อยที่สุดของหน่วยความจำที่ใช้ร่วมกัน

หน่วยความจำที่ใช้ร่วมกันเป็นหนึ่งในอาวุธหลักในการเพิ่มประสิทธิภาพ CUDA บทเรียนใช้ตัวอย่างที่เรียบง่ายมากที่มีเพียง 4 เธรดเพื่อแนะนำแนวคิดนี้:

// ที่มา: tutorials/cuda-cpp/notebooks/03.05-Shared-Memory/Sources/simple-shmem.cpp  
#include <cstdio>  

__global__ void kernel()  
{  
__shared__ int shared[4];  
shared[threadIdx.x] = threadIdx.x;  
__syncthreads();  

if (threadIdx.x == 0)  
{  
for (int i = 0; i < 4; i++) {  
std::printf("shared[%d] = %dn", i, shared[i]);  
}  
}  
}  

int main() {  
kernel<<<1, 4>>>();  
cudaDeviceSynchronize();  
return 0;  
}  

โค้ดนี้แสดงประเด็นสำคัญสามประการอย่างแม่นยำ:

• __shared__ ประกาศว่าอาร์เรย์จะถูกวางไว้ในหน่วยความจำที่ใช้ร่วมกันบนชิป
• __syncthreads() รับประกันว่าเธรดทั้งหมดเขียนเสร็จแล้วจึงจะอ่าน
• มีเพียงเธรดที่ 0 เท่านั้นที่ดำเนินการพิมพ์ เพื่อหลีกเลี่ยงความยุ่งเหยิงของเอาต์พุต

ผู้เรียนสามารถคอมไพล์และรันโค้ดนี้ได้ทันที สังเกตเอาต์พุต shared[0]=0, shared[1]=1, ... และสร้างความเข้าใจโดยสัญชาตญาณเกี่ยวกับการมองเห็นของหน่วยความจำที่ใช้ร่วมกัน

3.2 การปฏิบัติจริง: การเพิ่มประสิทธิภาพฮิสโตแกรมโดยใช้หน่วยความจำที่ใช้ร่วมกัน

หลังจากสร้างพื้นฐานแล้ว บทเรียนจะเข้าสู่การปฏิบัติจริงทันที – สอนวิธีเพิ่มประสิทธิภาพเคอร์เนลการคำนวณฮิสโตแกรมโดยใช้หน่วยความจำที่ใช้ร่วมกัน:

// ที่มา: tutorials/cuda-cpp/notebooks/03.05-Shared-Memory/Solutions/shmem.cpp  
__global__ void histogram_kernel(cuda::std::span<float> temperatures,  
cuda::std::span<int> histogram) {  
__shared__ int block_histogram[10];  

if (threadIdx.x < 10) {  
block_histogram[threadIdx.x] = 0;  
}  
__syncthreads();  

int cell = blockIdx.x * blockDim.x + threadIdx.x;  
int bin = static_cast<int>(temperatures[cell] / 10);  

cuda::atomic_ref<int, cuda::thread_scope_block> block_ref(  
block_histogram[bin]);  
block_ref.fetch_add(1);  
__syncthreads();  

if (threadIdx.x < 10) {  
cuda::atomic_ref<int, cuda::thread_scope_device> ref(  
histogram[threadIdx.x]);  
ref.fetch_add(block_histogram[threadIdx.x]);  
}  
}

โค้ดนี้มีคุณค่าสูงมากในระดับการสอน:

• กลไกการลดขนาดสองระดับ: ขั้นแรก ทำการสะสมในหน่วยความจำที่ใช้ร่วมกันระดับ block (โดยใช้การดำเนินการอะตอมมิกของ thread_scope_block) จากนั้น รวบรวมผลลัพธ์บางส่วนของแต่ละ block ไปยังหน่วยความจำส่วนกลาง (โดยใช้การดำเนินการอะตอมมิกของ thread_scope_device)
• สไตล์ CUDA C++ สมัยใหม่: ใช้ cuda::std::span แทนพอยน์เตอร์เปล่าแบบดั้งเดิม ใช้ cuda::atomic_ref แทนฟังก์ชัน atomicAdd แบบเก่า แสดงให้เห็นถึงอินเทอร์เฟซที่ทันสมัยและปลอดภัยต่อชนิดข้อมูลที่ไลบรารี libcu++ จัดเตรียมไว้
• ความสำคัญด้านประสิทธิภาพ: การดำเนินการอะตอมมิกบนหน่วยความจำที่ใช้ร่วมกันภายใน block นั้นเร็วกว่าการดำเนินการโดยตรงบนหน่วยความจำส่วนกลางถึงหนึ่งลำดับความสำคัญขึ้นไป สาเหตุหลักคือหลีกเลี่ยงการแย่งชิงการเชื่อมต่อระหว่าง Streaming Multiprocessor (SM) โดยสิ้นเชิง

3.3 อัลกอริทึมการทำงานร่วมกันของ CUB: ยืนบนไหล่ของยักษ์

จุดสุดท้ายของบทเรียน แนะนำ primitives ระดับ block ในไลบรารี CUB แสดงให้ผู้เรียนเห็นถึงการใช้งานการลดขนาดแบบร่วมมือระดับอุตสาหกรรม:

// ที่มา: tutorials/cuda-cpp/notebooks/03.06-Cooperative-Algorithms/Sources/block-sum.cpp  
#include <cub/block/block_reduce.cuh>  

constexpr int block_threads = 128;  

__global__ void block_sum()  
{  
using block_reduce_t = cub::BlockReduce<int, block_threads>;  
using storage_t = block_reduce_t::TempStorage;  

__shared__ storage_t storage;  

int block_sum = block_reduce_t(storage).Sum(threadIdx.x);  

if (threadIdx.x == 0)  
{  
printf("block sum = %dn", block_sum);  
}  
}  

int main() {  
block_sum<<<1, block_threads>>>();  
cudaDeviceSynchronize();  
return 0;  
}

BlockReduce ที่ไลบรารี CUB จัดเตรียมไว้นั้นใช้คำสั่ง warp-level shuffle และแผนผังการลดขนาดหลายรอบในระดับ底层 ซึ่งมีประสิทธิภาพสูงกว่าเวอร์ชันการลดขนาดหน่วยความจำที่ใช้ร่วมกันที่เขียนด้วยมือมาก ตรรกะการจัดเรียงของบทเรียนทั้งหมดชัดเจนมาก: ขั้นแรกให้คุณเขียนโค้ดหน่วยความจำที่ใช้ร่วมกันด้วยตัวเองเพื่อเข้าใจหลักการพื้นฐาน จากนั้นจึงบอกคุณว่า “ในงานวิศวกรรมจริง ควรใช้ CUB” – การเปลี่ยนผ่านจากหลักการไปสู่แนวทางปฏิบัติที่ดีที่สุดนี้ถือเป็นระดับตำราเรียน

สี่ Standard Parallelism: เส้นทางสู่การเร่งความเร็ว GPU โดยไม่มีโค้ด CUDA

4.1 การรันอัลกอริทึม C++ มาตรฐานบน GPU โดยตรง

บทเรียน stdpar แสดงให้นักพัฒนาเห็นเส้นทางการเขียนโปรแกรม GPU ที่แตกต่างไปจากเดิมอย่างสิ้นเชิง – การเร่งความเร็ว GPU โดยใช้อัลกอริทึมแบบขนานมาตรฐาน ISO C++ เท่านั้น:

// ที่มา: tutorials/stdpar/notebooks/cpp/lab1_daxpy/solutions/exercise5.cpp  
#include <algorithm>  
#include <execution>  
#include <ranges>  
#include <vector>  

void initialize(std::vector<double> &x, std::vector<double> &y) {  
std::for_each_n(std::execution::par, std::views::iota(0).begin(), x.size(),  
[x = x.data()](int i) { x[i] = (double)i; });  
std::fill_n(std::execution::par, y.data(), y.size(), 2.);  
}  

void daxpy(double a, std::vector<double> const &x, std::vector<double> &y) {  
std::for_each_n(std::execution::par,  
std::views::iota(0).begin(), x.size(),  
[a, x = x.data(), y = y.data()](int i) {  
y[i] += a * x[i];  
});  
}

ในโค้ดนี้ ไม่เห็นคีย์เวิร์ดเฉพาะของ CUDA เลย – ไม่มี __global__, ไม่มี <<<>>>, ไม่มี cudaMalloc นักพัฒนาเพียงแค่ต้องส่งนโยบายการดำเนินการ std::execution::par และใช้ตัวเลือกการคอมไพล์ nvc++ -stdpar=gpu อัลกอริทึมมาตรฐานเหล่านี้ก็จะทำงานบน GPU โดยอัตโนมัติ ที่ดียิ่งกว่านั้น บทเรียนยังเปรียบเทียบผลลัพธ์การทำงานของคอมไพเลอร์สี่ตัว (g++, clang++, nvc++ CPU backend, nvc++ GPU backend) เพื่อให้ผู้เรียนเห็นความแตกต่างของประสิทธิภาพของโค้ดชุดเดียวกันบนแบ็กเอนด์การดำเนินการที่แตกต่างกันอย่างเป็นรูปธรรม

4.2 สมการการนำความร้อนสองมิติ: การประยุกต์ใช้ stdpar ในทางปฏิบัติสำหรับ HPC

บทเรียนสาธิตเพิ่มเติมถึงวิธีการแก้ปัญหาการคำนวณทางวิทยาศาสตร์จริงด้วยอัลกอริทึมมาตรฐาน – การคำนวณ Stencil แบบขนานสำหรับสมการการนำความร้อนสองมิติ:

// ที่มา: tutorials/stdpar/notebooks/cpp/lab2_heat/exercise2.cpp  
double apply_stencil(double* u_new, double* u_old, grid g, parameters p) {
auto xs = std::views::iota(g.x_begin, g.x_end);
auto ys = std::views::iota(g.y_begin, g.y_end);
auto ids = std::views::cartesian_product(xs, ys);
return std::transform_reduce(
std::execution::par, ids.begin(), ids.end(),
        0., std::plus{}, [u_new, u_old, p](auto idx) {
auto [x, y] = idx;
return stencil(u_new, u_old, x, y, p);
});
}

ด้วยการใช้ std::views::cartesian_product สร้างพื้นที่ดัชนีสองมิติ จากนั้นใช้ std::transform_reduce เพื่อดำเนินการ Stencil แบบขนานพร้อมกับทำการลดขนาดผลรวมพลังงาน – นี่คือการแสดงออกที่กระชับของรูปแบบ map-reduce ใน C++ มาตรฐาน เมื่อรวมกับการสื่อสาร GPU หลายตัวผ่าน MPI โค้ดนี้แสดงให้เห็นถึงความสามารถที่สมบูรณ์ของอัลกอริทึมแบบขนานมาตรฐานในการประยุกต์ใช้ HPC ระดับอุตสาหกรรม

ห้า โครงสร้างพื้นฐานแบบคอนเทนเนอร์: การรับประกันทางวิศวกรรมสำหรับสภาพแวดล้อมการสอน

5.1 การจัดระเบียบบริการหลายรายการของ Docker Compose

บทเรียนนี้ใช้การกำหนดค่า Docker Compose ที่ออกแบบมาอย่างพิถีพิถัน เพื่อให้ “ปรับใช้ได้ด้วยคลิกเดียว”:

# ที่มา: tutorials/cuda-cpp/brev/docker-compose.yml  
services:
base:
<<: [*gpu-config, *common-service]
image: *image
entrypoint: ["/accelerated-computing-hub/brev/entrypoint.bash", "base"]
jupyter:
<<: [*gpu-config, *common-service, *persistent-service]
image: *image
entrypoint: ["/accelerated-computing-hub/brev/entrypoint.bash", "jupyter"]
ports:
      - "127.0.0.1:8888:8888"
nsys:
image: nvcr.io/nvidia/devtools/nsight-streamer-nsys:2026.1.1
ports:
      - "127.0.0.1:8080:8080"

เพียงรัน docker compose up เพียงครั้งเดียว ก็สามารถเริ่มบริการสามรายการพร้อมกันได้: JupyterLab, โปรไฟล์ประสิทธิภาพ Nsight Systems และโปรไฟล์ Nsight Compute ผู้เรียนหลังจากเขียน Kernel ใน Notebook เสร็จแล้ว สามารถดูไทม์ไลน์ GPU และการวิเคราะห์ระดับคำสั่งในเบราว์เซอร์ได้โดยตรง – วงจรป้อนกลับทันทีแบบ “เขียนโค้ด-วิเคราะห์” นี้เป็นสิ่งที่วิธีการสอนแบบดั้งเดิมทำได้ยาก

สรุป

accelerated-computing-hub ของ NVIDIA ไม่ใช่แค่คลังโค้ดเท่านั้น มันเป็นตัวแทนของ กระบวนทัศน์วิศวกรรมซอฟต์แวร์ที่ขับเคลื่อนด้วยการสอน:

1. การเพิ่มภาระทางปัญญาแบบค่อยเป็นค่อยไป: แต่ละไฟล์โค้ดแนะนำแนวคิดใหม่เพียงแนวคิดเดียว การเปลี่ยนจาก <<<1,1>>> ไปเป็น cub::BlockReduce ถูกแบ่งออกเป็นขั้นตอนย่อยๆ มากกว่าสิบขั้นตอนที่ย่อยง่าย
2. การออกแบบทางเข้าหลายทาง: สามเส้นทาง (CUDA C++/stdpar/Python) ช่วยให้นักพัฒนาที่มีพื้นฐานแตกต่างกันสามารถหาจุดเริ่มต้นที่เหมาะสมได้
3. สภาพแวดล้อมคือโค้ด: Docker Compose และอิมเมจที่สร้างไว้ล่วงหน้าช่วยขจัด “การกำหนดค่าสภาพแวดล้อม” ซึ่งเป็นอุปสรรคการเรียนรู้ที่ใหญ่ที่สุด
4. ให้ความสำคัญกับ C++ สมัยใหม่: ใช้อินเทอร์เฟซสมัยใหม่อย่างเต็มที่ เช่น cuda::std::span, cuda::atomic_ref, std::mdspan แทน CUDA API แบบ C ดั้งเดิม

สำหรับนักพัฒนาที่ต้องการเรียนรู้การเขียนโปรแกรม GPU อย่างเป็นระบบ คลังนี้มีเส้นทางการเรียนรู้แบบโอเพนซอร์สที่มีโครงสร้างมากที่สุดและทำซ้ำได้มากที่สุดในปัจจุบัน และสำหรับนักออกแบบการสอน มันแสดงให้เห็นถึงวิธีการถ่ายทอดความรู้ที่ซับซ้อนอย่างมีประสิทธิภาพผ่านการจัดระเบียบโค้ดแบบค่อยเป็นค่อยไป

เอกสารอ้างอิง

[1] บทเรียน CUDA Tile: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/cuda-tile/README.md

[2] บทเรียน CUDA C++: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/cuda-cpp/README.md

[3] บทเรียน Standard Parallelism: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/stdpar/README.md

[4] บทเรียน Accelerated Python: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/accelerated-python/README.md

[5] บทเรียน nvmath-python: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/nvmath-python/README.md

[6] บทเรียน CUDA Python – CuPy, cuDF, CCCL และเคอร์เนล (เวอร์ชัน 8 ชั่วโมง): https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/accelerated-python/notebooks/syllabi/cuda_python__cupy_cudf_cccl_kernels__8_hours.ipynb

[7] บทเรียน CUDA Python – cuda.core และ CCCL (เวอร์ชัน 2 ชั่วโมง): https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/accelerated-python/notebooks/syllabi/cuda_python__cuda_core_cccl__2_hours.ipynb

[8] บทเรียน PyHPC – NumPy, CuPy และ mpi4py (เวอร์ชัน 4 ชั่วโมง): https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/accelerated-python/notebooks/syllabi/pyhpc__numpy_cupy_mpi4py__4_hours.ipynb

บทเรียนการปรับใช้ GPU: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/gpu-deployment/gpu-deployment-from-scratch.md

อ่านเพิ่มเติม

• ☕ สนับสนุนค่ากาแฟทีมงาน

  หากคุณชอบบทความนี้ สามารถสนับสนุนเราได้ผ่าน PromptPay

  
  SCAN TO PAY WITH ANY BANK