0Pricing
CUDA Academy · บทเรียน

เคอร์เนลคูณเมทริกซ์แบบพื้นฐาน

วิธีพื้นฐานที่ใช้ดัชนีสองมิติและข้อจำกัด

เคอร์เนลคูณเมทริกซ์แบบพื้นฐาน เป็นบทเรียน CUDA Academy ฟรีบน CoddyKit นี่คือบทเรียนที่ 1 จากทั้งหมด 4 บทเรียน คุณสามารถอ่านบทเรียนทั้งหมดด้านล่างฟรี — จากนั้นลองปฏิบัติด้วยตัวคุณเองในเบราว์เซอร์พร้อมตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7 บทเรียนนี้เป็นส่วนหนึ่งของเส้นทางการเรียน CUDA Academy และความก้าวหน้าของคุณจะซิงค์ข้ามเว็บและแอป CoddyKit คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน

การคูณเมทริกซ์ในสไตล์ GPU

การคูณเมทริกซ์เป็นหัวใจของกราฟิกและปัญญาประดิษฐ์ วันนี้คุณจะสร้างเวอร์ชัน GPU แบบพื้นฐานก่อน แล้วเรียนรู้ว่าเหตุใดจึงยังดึงประสิทธิภาพออกมาได้ไม่เต็มที่

คณิตศาสตร์ในบรรทัดเดียว

เซลล์ผลลัพธ์แต่ละเซลล์ C[row][col] คือผลคูณจุด: นำแถวเต็มของ A คูณกับคอลัมน์เต็มของ B แล้วหาผลรวมของผลลัพธ์ 🧮

C[row][col] = sum over k of A[row][k] * B[k][col]

หนึ่งเธรดต่อหนึ่งผลลัพธ์

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

ตารางเธรดแบบสองมิติ

เนื่องจาก C เป็นกริด 2 มิติ คุณจึงเปิดใช้งานเธรดใน สองมิติ โดยดัชนี x จะจับคู่กับคอลัมน์ และดัชนี y จะจับคู่กับแถว

dim3 threads(16, 16);
dim3 blocks((N+15)/16, (N+15)/16);

การค้นหาเซลล์ของเธรดนี้

ภายในเคอร์เนล แต่ละเธรดจะคำนวณ แถวและคอลัมน์ ของตนเองจากดัชนีบล็อกและเธรด เช่นเดียวกับการทำดัชนีแบบ 1 มิติ แต่ทำบนทั้งสองแกน

int row = blockIdx.y*blockDim.y + threadIdx.y;
int col = blockIdx.x*blockDim.x + threadIdx.x;

การตรวจสอบขอบเขต

กริดจะปัดค่าขึ้น ดังนั้นเธรดบางส่วนจึงอยู่นอกเมทริกซ์ ให้ป้องกันด้วย if (row < N && col < N) ก่อนเข้าถึงหน่วยความจำ

if (row < N && col < N) {
  // safe to compute
}

ลูปด้านใน

แต่ละเธรดจะทำงานผ่าน ลูปของ k โดยสะสมผลคูณไว้ในผลรวมเฉพาะที่ ตัวแปรเฉพาะที่นี้จะอยู่ในรีจิสเตอร์ความเร็วสูง

float sum = 0.0f;
for (int k = 0; k < N; ++k)
  sum += A[row*N+k] * B[k*N+col];

การเขียนผลลัพธ์

เมื่อลูปทำงานเสร็จ เธรดจะเก็บผลรวมสะสมลงใน C เพียงครั้งเดียว หนึ่งเธรดต่อหนึ่งการเขียนที่ชัดเจน

C[row*N + col] = sum;

การทำให้แบนแบบเรียงตามแถว

เมทริกซ์เป็นอาร์เรย์ 1 มิติแบบแบน ดังนั้นคุณจึงใช้ดัชนีเป็น row*N + col การจัดเลย์เอาต์ให้ถูกต้องถือเป็นครึ่งหนึ่งของงานในการคูณเมทริกซ์

เหตุใดจึงทำงานได้ แต่ช้า

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

ต้นทุนที่ซ่อนอยู่

เธรดที่อยู่ติดกันอ่านแถวของ A และคอลัมน์ของ B ซ้ำแล้วซ้ำอีก การรับส่งข้อมูลในหน่วยความจำ ที่สูญเปล่านี้คือสิ่งที่การแบ่งเป็นไทล์จะเข้ามาแก้ไขต่อไป

ตรวจสอบอย่างรวดเร็ว

ลองคิดดูว่าเคอร์เนลแบบพื้นฐานจับคู่งานกับเธรดอย่างไร

สรุป

คุณจับคู่หนึ่งเธรดกับหนึ่งเซลล์ผลลัพธ์ วนลูปผ่าน k โดยอ่านจาก หน่วยความจำโกลบอล และเห็นการอ่านซ้ำซ้อนแล้ว ขั้นต่อไปคือการลดการรับส่งข้อมูลนี้ด้วยการแบ่งเป็นไทล์ 🚀

คำถามที่พบบ่อย

บทเรียน “เคอร์เนลคูณเมทริกซ์แบบพื้นฐาน” ฟรีหรือไม่

ใช่ — ข้อความเต็มของ “เคอร์เนลคูณเมทริกซ์แบบพื้นฐาน” ฟรีให้อ่านที่นี่บนเว็บ เพื่อปฏิบัติแบบโต้ตอบ (ตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7) และปลดล็อคส่วนที่เหลือของคอร์ส CUDA Academy ให้อัปเกรดเป็น CoddyKit PRO คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน

คุณจะเรียนรู้อะไรในบทเรียน “เคอร์เนลคูณเมทริกซ์แบบพื้นฐาน”

วิธีพื้นฐานที่ใช้ดัชนีสองมิติและข้อจำกัด คุณปฏิบัติ CUDA Academy ด้วยโค้ดที่ใช้งานได้จริงที่คุณเรียกใช้โดยตรงในเบราว์เซอร์ และติวเตอร์ AI ตลอด 24/7 ตอบคำถามของคุณขณะที่คุณไปผ่านบทเรียน

คุณต้องมีประสบการณ์ก่อนที่จะเริ่มเรียน CUDA Academy หรือไม่

ไม่จำเป็นต้องมีประสบการณ์มาก่อน CUDA Academy บน CoddyKit ออกแบบมาสำหรับผู้เริ่มต้นไปจนถึงผู้เรียนขั้นสูง คุณสามารถเริ่มต้นที่นี่หรือเริ่มจากตัวแรกและเรียนด้วยความเร็วของคุณเอง นี่คือบทเรียนที่ 1 จากทั้งหมด 4 บทเรียน

บทเรียน “เคอร์เนลคูณเมทริกซ์แบบพื้นฐาน” ใช้เวลานานแค่ไหน

บทเรียน CoddyKit ส่วนใหญ่ใช้เวลาประมาณ 5–10 นาที แต่ละบทเรียนจึงสั้นและเป็นแบบโต้ตอบ คุณสามารถก้าวหน้าอย่างต่อเนื่องและกลับมาเรียนต่อจากตรงที่เพิ่งหยุดบนเว็บและแอปได้เลย

ฉันเขียนและรันโค้ดในบทเรียน CUDA Academy นี้ได้ไหม

ได้ บทเรียน CUDA Academy ทุกบทมีตัวแก้ไขโค้ดในตัว คุณจึงเขียนและรันโค้ดจริงได้เลยในเบราว์เซอร์ และได้รับข้อเสนอแนะจาก AI ในทันที — ไม่ต้องติดตั้งในเครื่องของคุณ

บทเรียนทั้งหมดในหลักสูตรนี้

  1. เคอร์เนลคูณเมทริกซ์แบบพื้นฐาน
  2. แบ่งผลคูณภายในเป็นไทล์
  3. วนซ้ำตามช่วงของไทล์
  4. วัดการเพิ่มความเร็ว
← กลับไปที่ CUDA Academy