Programming Models: MPI & CUDA

1สามตระกูลหลักของ Parallel Programming Model

2Message-Passing Interface (MPI)

ตัวอย่าง: คำนวณค่า π ด้วยการหาพื้นที่ใต้กราฟ (numerical integration)

#include "mpi.h"
#include <math.h>

int main(argc,argv)
int argc; char *argv[];
{
    int done = 0, n, myid, numprocs, i;
    double PI25DT = 3.141592653589793238462643;
    double mypi, pi, h, sum, x;

    MPI_Init(&argc,&argv);
    MPI_Comm_size(MPI_COMM_WORLD,&numprocs);
    MPI_Comm_rank(MPI_COMM_WORLD,&myid);
    while (!done) {
        if (myid == 0) {
            printf("Enter the number of intervals: (0 quits) ");
            scanf("%d",&n);
        }
        MPI_Bcast(&n, 1, MPI_INT, 0, MPI_COMM_WORLD);
        if (n == 0) break;

        h = 1.0 / (double) n;
        sum = 0.0;
        for (i = myid + 1; i <= n; i += numprocs) {
            x = h * ((double)i - 0.5);
            sum += 4.0 / (1.0 + x*x);
        }
        mypi = h * sum;

        MPI_Reduce(&mypi, &pi, 1, MPI_DOUBLE, MPI_SUM, 0, MPI_COMM_WORLD);
        if (myid == 0)
            printf("pi is approximately %.16f, Error is %.16f\n",
                   pi, fabs(pi - PI25DT));
    }
    MPI_Finalize();
    return 0;
}

3CUDA: เขียนโปรแกรมสำหรับ GPU

เวอร์ชันไร้เดียงสา: หนึ่งเธรดต่อหนึ่งสมาชิกของเมทริกซ์ผลลัพธ์

__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
    float Cvalue = 0;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    for (int e = 0; e < A.width; ++e)
        Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col];
    C.elements[row * C.width + col] = Cvalue;
}

เวอร์ชันปรับปรุง: ใช้ Shared Memory แบบ Tiling

// โหลดชิ้นส่วนย่อย (tile) ของ A และ B มาไว้ใน shared memory ก่อน
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
As[row][col] = GetElement(Asub, row, col);
Bs[row][col] = GetElement(Bsub, row, col);
__syncthreads();               // รอให้ทุกเธรดใน block โหลดเสร็จก่อน

for (int e = 0; e < BLOCK_SIZE; ++e)
    Cvalue += As[row][e] * Bs[e][col];   // อ่านจาก shared memory (เร็วกว่า global มาก)
__syncthreads();               // รอให้ทุกเธรดคำนวณ tile นี้เสร็จก่อนโหลด tile ถัดไป

4MPI และ CUDA ในเครื่องมือปัจจุบัน

5Programming Model คือสัญญาระหว่างความคิดกับเครื่อง

Programming model ไม่ใช่เพียงชุดคำสั่ง แต่กำหนดว่าโปรแกรมเมอร์มองงาน ข้อมูล และการสื่อสารอย่างไร Shared-memory model ทำให้หลาย thread เห็น address space เดียว Message-passing model บังคับให้ส่งข้อมูลอย่างชัดเจน ส่วน accelerator model แยก host กับ device และจัดงานเป็น kernels

6OpenMP: สะพานจากโปรแกรมลำดับไปสู่ Shared Memory

แม้ชื่อบทเน้น MPI และ CUDA แต่ OpenMP เป็นจุดเริ่มที่สำคัญสำหรับ multicore เพราะเปิด parallel region และแบ่ง loop ด้วย directive โดยยังใช้หน่วยความจำร่วม เหมาะกับการทดลอง decomposition ก่อนขยายไป distributed memory

#pragma omp parallel for reduction(+:sum) schedule(static)
for (long i = 0; i < n; ++i) {
    sum += a[i];
}

reduction สร้างผลย่อยต่อ thread แล้วรวมอย่างปลอดภัย ดีกว่าให้ทุก thread atomic update ตัวแปรเดียวทุก iteration ซึ่งสร้าง contention ส่วน schedule(static) เหมาะเมื่องานแต่ละ iteration ใกล้เคียงกัน หากไม่เท่ากันอาจใช้ dynamic หรือ guided แต่ต้องจ่าย scheduling overhead เพิ่ม

คำถามเรื่อง Data Scope

• shared — ทุก thread เห็น object เดียว ต้องระวัง race
• private — แต่ละ thread มีสำเนาของตน
• firstprivate — private แต่เริ่มจากค่าก่อนเข้า region
• reduction — private partial results แล้วรวมด้วย operation ที่กำหนด

7MPI Process, Rank และ Communicator

MPI สร้าง processes ที่มี address space แยกกัน แต่ละ process มี rank ภายใน communicator Rank ไม่ใช่ชื่อเครื่องถาวร แต่เป็นหมายเลขสมาชิกในกลุ่มการสื่อสารเดียวกัน

• MPI_COMM_WORLD คือ communicator เริ่มต้นที่รวมทุก process
• Communicator แยก namespace ของ message และสร้าง subgroup ได้
• Rank ใช้ระบุผู้ส่ง/ผู้รับภายใน communicator นั้น
• Tag แยกชนิดหรือรอบของ message เพื่อจับคู่ให้ถูก

การแยก communicator มีประโยชน์เมื่อ algorithm มีหลายกลุ่ม เช่น process rows กับ columns ใน matrix grid Message ที่มี source/tag เหมือนกันแต่คนละ communicator จะไม่ชนกัน ช่วยให้ library หลายชิ้นใช้ MPI ร่วมกันได้โดยไม่แย่ง namespace

8Point-to-Point Communication: ส่ง รับ และความหมายของคำว่าเสร็จ

MPI_Send กับ MPI_Recv ดูตรงไปตรงมา แต่คำว่า send เสร็จอาจหมายถึง buffer ต้นทางใช้ซ้ำได้ ไม่จำเป็นต้องหมายถึงผู้รับประมวลผลแล้ว MPI implementation อาจ buffer ข้อมูลหรือรอ rendezvous ตามขนาด message

9Deadlock ใน MPI

ถ้าทุก process เรียก blocking send ของ message ใหญ่ไปหาเพื่อนก่อนเรียก receive ทุกคนอาจรอกัน เพราะ implementation ไม่มี buffer พอ วิธีที่ดูทำงานบน input เล็กอาจค้างเมื่อ message โต

// เสี่ยง deadlock: ทั้งสองฝั่งส่งก่อนรับ
MPI_Send(... peer ...);
MPI_Recv(... peer ...);

// ทางเลือก
MPI_Sendrecv(...);
// หรือ post MPI_Irecv ก่อน MPI_Isend แล้ว Waitall

การป้องกันทำได้ด้วยลำดับส่ง/รับที่ไม่เป็นวง ใช้ MPI_Sendrecv หรือ post receives ล่วงหน้า สิ่งสำคัญคืออย่าพึ่งพา eager buffering ซึ่งเป็นรายละเอียด implementation ไม่ใช่ guarantee ของ algorithm

10Collective Communication: ใช้ความหมายระดับกลุ่ม

Collective ไม่ใช่เพียง shorthand ของ loop ส่งทีละคน Library เลือก tree, ring, recursive doubling หรือ hierarchical algorithm ตาม message size และ topology จึงมักเร็วและปลอดภัยกว่าการเขียนเอง

11Communication Cost: Latency กับ Bandwidth

แนวทางลดต้นทุนจึงมีทั้งรวม message เล็ก, ลดจำนวนรอบ, overlap กับ compute และเปลี่ยน decomposition ให้สื่อสารเฉพาะเพื่อนบ้าน แต่การ batch ใหญ่เกินไปเพิ่มเวลารอสะสมข้อมูลและ memory footprint ต้องเลือกตาม latency target

12Domain Decomposition และ Halo Exchange

Simulation grid มักแบ่งพื้นที่เป็น blocks แต่ละ rank คำนวณภายในของตน ส่วนเซลล์ขอบต้องใช้ค่าจากเพื่อนบ้าน จึงเก็บ ghost/halo cells และแลกค่าขอบทุก iteration

1. Post nonblocking receives จากเพื่อนบ้าน
2. Post nonblocking sends ของ boundary
3. คำนวณ interior ที่ไม่ต้องรอข้อมูล
4. Wait ให้ halo มาถึง
5. คำนวณ boundary แล้วเข้าสู่รอบถัดไป

โครงนี้ overlap communication กับ computation ได้ หาก block เล็กเกินไป surface-to-volume ratio สูง ข้อมูลขอบมีสัดส่วนมากและ scaling แย่ นี่เชื่อม granularity กับ geometry ของ decomposition โดยตรง

13MPI Derived Datatype และ Packing

ข้อมูลที่ต้องส่งอาจไม่ต่อเนื่อง เช่น column ของ matrix ใน row-major layout เราสามารถ pack ใส่ buffer ต่อเนื่องเอง หรือสร้าง derived datatype ให้ MPI รู้ stride และ layout

Derived datatype ทำให้โค้ดสื่อความหมายและเปิดโอกาสให้ implementation optimize แต่ไม่ได้รับประกันว่าจะไม่มีการ pack ภายใน ต้องวัดเทียบ manual packing โดยเฉพาะ layout ซับซ้อน การจัดข้อมูลใหม่ให้ communication-friendly อาจดีกว่าพึ่ง datatype แก้ทุกครั้ง

14One-Sided Communication

MPI RMA เปิดให้ process อ่านเขียน window ของอีก process ผ่าน Put/Get/Accumulate โดยผู้รับไม่ต้อง post receive ที่ตรงกันทุก message เหมาะกับ irregular access หรือ shared-data abstraction บางรูปแบบ

แต่ one-sided ไม่ได้แปลว่าไม่ต้อง synchronize ต้องกำหนด epoch และ visibility ว่าเมื่อใดข้อมูลใช้ได้ ความซับซ้อนย้ายจากการจับคู่ send/receive ไปสู่ memory model และ synchronization semantics จึงควรเลือกเพราะ pattern เหมาะ ไม่ใช่เพราะชื่อฟังดูเร็วกว่า

15CUDA Execution Hierarchy

Kernel ถูกเรียกด้วย grid ซึ่งมี thread blocks และแต่ละ block มี threads Scheduler วาง block ลง Streaming Multiprocessor (SM) โดย blocks ต้องเป็นอิสระต่อกันในระดับที่รันลำดับใดก็ได้

• Thread ประมวลผล element หรือกลุ่มงานย่อย
• Warp กลุ่ม threads ที่ execute instruction ร่วมกัน
• Block กลุ่ม threads ที่ใช้ shared memory และ synchronize กันได้
• Grid ชุด blocks ของ kernel invocation

Global synchronization ระหว่าง blocks ปกติต้องจบ kernel แล้วเรียก kernel ถัดไป เพราะ blocks อาจไม่ได้ resident พร้อมกัน การออกแบบที่สมมติว่าทุก block รอ barrier เดียวกันอาจ deadlock หรือใช้ไม่ได้ตาม model

16Indexing: จาก Thread ไปยังข้อมูล

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) c[i] = a[i] + b[i];

สูตรนี้ map thread ID ไป element ของ array หาก n ไม่หาร block size ลงตัวต้องตรวจขอบเขต Block size มีผลต่อ occupancy, scheduling และ memory access แต่ค่าที่หารด้วย warp size ลงตัวเป็นเพียงจุดเริ่ม ไม่ใช่คำตอบสากล

17CUDA Memory Hierarchy

18Coalescing และ Memory Access Pattern

Threads ใน warp ควรอ่าน addresses ที่อยู่ต่อเนื่องหรือรวมเป็น memory transactions จำนวนน้อย หากแต่ละ thread กระโดดห่างหรืออ่านคนละทิศ ระบบต้องสร้าง transactions มากและ bandwidth ที่มีประโยชน์ลดลง

Data layout แบบ Structure of Arrays มักเหมาะกับการอ่าน field เดียวของหลาย objects มากกว่า Array of Structures เพราะค่าเดียวกันอยู่ต่อเนื่อง แต่หากแต่ละ thread ใช้ทุก fields ของ object การตัดสินอาจต่างกัน ต้องวิเคราะห์ access pattern จริง

19Shared Memory Tiling และ Bank Conflict

Matrix multiplication แบบ tiled โหลด A และ B blocks เข้า shared memory แล้ว threads ใช้ซ้ำหลายครั้ง ลด global bytes ต่อ operation หลังโหลดแต่ละ tile ต้อง synchronize threads ใน block ก่อนใช้ข้อมูลร่วม

Shared memory แบ่งเป็น banks หาก threads หลายตัวเข้าคนละ address ที่ตก bank เดียวกัน request อาจถูก serialize เรียกว่า bank conflict Padding matrix tile หนึ่ง column บางครั้งช่วยเปลี่ยน mapping และลด conflict

20Occupancy ไม่ใช่เป้าหมายสุดท้าย

Occupancy คือสัดส่วน active warps ต่อขีดจำกัด SM ช่วยซ่อน latency แต่ occupancy สูงสุดไม่ได้รับประกัน performance สูงสุด Kernel ที่ใช้ register มากอาจลด occupancy แต่คำนวณมีประสิทธิภาพและไม่ spill ส่วนการลด register เพื่อเพิ่ม occupancy อาจทำให้ local-memory traffic สูงขึ้น

เป้าหมายคือมี warps พอซ่อน latency พร้อมใช้ execution units และ bandwidth ได้ดี ต้องอ่าน profiler ทั้ง achieved occupancy, eligible warps, memory throughput, instruction mix และ stall reasons

21CUDA Streams และการซ้อน Transfer กับ Compute

Stream เป็นลำดับงานบน GPU Operations ใน stream เดียวรักษาลำดับ ส่วนคนละ streams อาจ overlap หาก hardware และ dependency อนุญาต เราสามารถแบ่งข้อมูลเป็น chunks แล้วทำ H2D transfer ของ chunk ถัดไประหว่าง kernel ประมวลผล chunk ปัจจุบัน

22Unified Memory: ใช้ง่ายขึ้น แต่การย้ายข้อมูลยังไม่หาย

Unified Memory ให้ pointer ที่ CPU และ GPU ใช้ได้ ลดงานจัดสรรและ copy ด้วยมือ Runtime ย้าย page ตามการเข้าถึง แต่ถ้า access pattern สลับ host/device ถี่ page อาจ thrash และช้ากว่าการวางข้อมูลอย่างชัดเจน

Prefetch และ memory advice ช่วยบอกตำแหน่งใช้งานที่คาดไว้ Unified addressing เป็น abstraction ที่ดี แต่ไม่ควรตีความว่า CPU กับ GPU มีหน่วยความจำราคาเดียวกัน Data movement ยังคงอยู่ เพียงมีระบบจัดการแทนเรา

23Race Condition และการตรวจความถูกต้อง

GPU มี threads จำนวนมาก Race อาจเกิดเฉพาะบาง scheduling จึงผ่านการทดสอบหลายครั้งก่อนผิด Atomics ช่วย update ค่าร่วมอย่างปลอดภัย แต่ถ้าทุก thread atomic ที่ตำแหน่งเดียว performance จะ serialize

• ให้แต่ละ block/thread สร้าง partial result แล้ว reduction เป็นชั้น
• ใช้ synchronization เฉพาะ threads ใน block ตาม scope ที่รับประกัน
• ตรวจ out-of-bounds และ race ด้วยเครื่องมือ sanitizer
• เปรียบเทียบกับ reference implementation บน input เล็ก
• ใช้ tolerance สำหรับ floating point และทดสอบ NaN/Inf

24Hybrid MPI + CUDA

GPU cluster มักใช้ MPI ระหว่าง processes/nodes และ CUDA ภายใน node แต่ต้องวาง rank ให้ใกล้ GPU และ NIC ที่ใช้ รวมถึงหลีกเลี่ยงให้หลาย ranks แย่ง GPU โดยไม่ตั้งใจ

1. แบ่ง global domain ต่อ MPI rank
2. แต่ละ rank เก็บ subdomain บน GPU ของตน
3. คำนวณ interior ด้วย kernel
4. แลก halo ระหว่าง GPUs ผ่าน CUDA-aware MPI หรือ stage ผ่าน host ตามระบบ
5. Overlap halo transfer กับ interior compute

GPUDirect และ topology-aware communication ลดการ copy ผ่าน host แต่ performance ขึ้นกับ PCIe/NVLink/NIC placement การรู้เพียงจำนวน GPUs จึงไม่พอ ต้องรู้เส้นทางของ bytes

25Library, Directive และ Framework: ไม่ต้องเขียน Kernel ทุกครั้ง

การใช้ library ไม่ได้แปลว่าไม่ต้องเข้าใจ architecture เมื่อ performance ไม่ดี เราต้องรู้ว่า tensor layout, batch size, fusion หรือ transfer ทำให้ library ใช้ hardware ไม่เต็มตรงไหน

26กรณีศึกษา: Matrix Multiplication จาก OpenMP สู่ MPI และ CUDA

1. เริ่มจาก blocked serial matrix multiplication เพื่อใช้ cache ให้ดี
2. OpenMP แบ่ง output tiles ระหว่าง CPU threads
3. MPI ใช้ 2D block distribution ระหว่าง nodes และ collective แลก panels
4. CUDA ให้แต่ละ GPU คำนวณ local tiles ด้วย shared-memory tiling
5. Hybrid version วางหนึ่ง rank ต่อ GPU และ overlap communication

แต่ละขั้นไม่ได้แทนขั้นก่อนทั้งหมด Blocking ยังจำเป็นบน GPU และ MPI local computation ยังต้อง optimize ภายใน node Parallelism เป็นชั้นซ้อนกัน และ implementation ที่ดีรักษา locality ในทุกชั้น

27แล็บและการตรวจงาน

28Debugging และ Performance Tooling

Bug ในโปรแกรมขนานมักขึ้นกับลำดับเวลาที่ไม่แน่นอน การใส่ print อาจเปลี่ยน timing จน bug หาย จึงต้องใช้ invariant, input เล็กที่ตรวจด้วยมือ และเครื่องมือเฉพาะ

Timeline ช่วยดู compute, communication, idle และ synchronization ส่วน counter ช่วยอธิบาย occupancy, memory throughput และ stall reasons ควรใส่ application markers ให้ trace เชื่อมกับ phase ของ algorithm ได้

29Portability กับ Performance Portability

Code portability หมายถึงย้ายไปรันอีก platform ได้ ส่วน performance portability หมายถึงยังใช้ทรัพยากรได้ดีบนหลาย architecture โค้ด abstraction สูงอาจ compile ได้ทุกที่แต่ช้า เพราะ memory hierarchy และ execution width ต่างกัน

ควรแยก algorithm จาก backend ใช้ library มาตรฐาน เก็บ architecture-specific kernel เฉพาะ hot path และมี benchmark regression ต่อ platform Framework อย่าง SYCL, Kokkos หรือ directive model ช่วยได้ แต่ยังต้อง tune layout และ granularity

30Task Runtime และ Work Stealing

เมื่อ DAG irregular การกำหนด thread ต่อ task ด้วยมือไม่เหมาะ Task runtime ให้ประกาศ tasks กับ dependencies แล้ว schedule บน worker pool Work stealing ให้ worker ว่างขโมยงานจาก queue อื่น ช่วย balance โดยไม่ใช้ scheduler กลางทุกครั้ง

Task ต้องใหญ่พอและการขโมยอาจเสีย locality จึงควรมี cutoff รวม recursive tasks เล็ก และให้ hint เรื่อง affinity เมื่อข้อมูลมีราคาเคลื่อนย้ายสูง

31Checklist ก่อนส่งงาน MPI/CUDA

• ทุก rank/thread ให้ผลถูกบน input เล็ก
• ไม่มี buffer ถูกแก้ก่อน nonblocking operation complete
• Collectives เรียกตามลำดับเดียวกันใน communicator
• วัด message count, size และ halo volume แล้ว
• GPU indexing ทดสอบขนาดที่หาร block ไม่ลงตัว
• รวม transfer/setup ใน end-to-end benchmark
• ตรวจ race และ out-of-bounds ด้วยเครื่องมือ
• รายงาน placement ของ ranks, threads และ GPUs

คำถามทบทวนการเลือก Programming Model

ก่อนเลือก OpenMP, MPI หรือ CUDA ให้เขียนคำตอบสั้น ๆ ว่า address space เป็นแบบใด งานหนึ่งชิ้นใหญ่แค่ไหน ข้อมูลต้องย้ายผ่านเส้นทางใด ใครต้องรอใคร และผลลัพธ์รวมตรงไหน หากตอบไม่ได้ การเลือก API ยังเร็วเกินไป เพราะ model ที่เหมาะไม่ได้ตัดสินจากจำนวนเครื่องอย่างเดียว แต่ตัดสินจากรูปของ dependency กับ data movement

เมื่อระบบมีหลายชั้น ควรระบุ owner ของ parallelism ด้วย เช่น MPI จัดระหว่าง nodes, OpenMP จัด cores และ CUDA จัด GPU threads เพื่อป้องกัน runtime หลายตัวสร้าง workers ซ้อนกันจน oversubscription และแย่ง bandwidth ครับ

32สรุปและขั้นตอนถัดไป

OpenMP, MPI และ CUDA เปิดให้เราแสดง parallelism คนละระดับ OpenMP แชร์หน่วยความจำในเครื่อง MPI ทำให้การสื่อสารข้าม address space ชัดเจน และ CUDA จัดงานจำนวนมากบน accelerator สิ่งที่เชื่อมทั้งสามคือ decomposition, locality, granularity และการรอ

การเขียนโปรแกรมให้รันหลาย worker เป็นเพียงครึ่งแรก ครึ่งที่ยากกว่าคือพิสูจน์ว่าเร็วขึ้นเพราะอะไรและยังคุ้มเมื่อ scale ต่อหรือไม่ บทถัดไปจึงว่าด้วย performance measurement ตั้งแต่ speedup, efficiency และ cost ไปจนถึง Karp–Flatt, isoefficiency, Roofline และ benchmark ที่เทียบกันได้ครับ