แบ่งผลคูณภายในเป็นไทล์
โหลดไทล์ย่อยของ A และ B ในแต่ละช่วง
แบ่งผลคูณภายในเป็นไทล์ เป็นบทเรียน CUDA Academy ฟรีบน CoddyKit นี่คือบทเรียนที่ 2 จากทั้งหมด 4 บทเรียน คุณสามารถอ่านบทเรียนทั้งหมดด้านล่างฟรี — จากนั้นลองปฏิบัติด้วยตัวคุณเองในเบราว์เซอร์พร้อมตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7 บทเรียนนี้เป็นส่วนหนึ่งของเส้นทางการเรียน CUDA Academy และความก้าวหน้าของคุณจะซิงค์ข้ามเว็บและแอป CoddyKit คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน
แนวคิดการแบ่งเป็นไทล์
การแบ่งเป็นไทล์จะแยกเมทริกซ์ออกเป็น ไทล์ สี่เหลี่ยมขนาดเล็กที่พอดีกับหน่วยความจำบนชิปความเร็วสูง เธรดจะร่วมมือกันโหลดไทล์หนึ่งครั้งแล้วนำกลับมาใช้ซ้ำหลายครั้ง
เหตุผลที่ใช้หน่วยความจำร่วม
ไทล์จะอยู่ในหน่วยความจำ __shared__ ซึ่งเธรดทุกตัวในบล็อกมองเห็นได้ การอ่านจากที่นี่เร็วกว่าการเข้าถึงหน่วยความจำโกลบอลซ้ำแล้วซ้ำอีกมาก ⚡
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];หนึ่งบล็อกต่อหนึ่งไทล์ผลลัพธ์
แต่ละ block รับผิดชอบส่วนของผลลัพธ์ C ขนาด TILE คูณ TILE เธรดในบล็อกจะร่วมมือกันคำนวณส่วนทั้งหมดนี้
ขนาดไทล์ตรงกับบล็อก
คุณกำหนดให้ TILE เท่ากับความกว้างของบล็อก ซึ่งมักเป็น 16 หรือ 32 ด้วยวิธีนี้ แต่ละเธรดจะโหลดองค์ประกอบจากแต่ละ ไทล์ ได้พอดีหนึ่งองค์ประกอบ
#define TILE 16
dim3 threads(TILE, TILE);การจับคู่เธรดกับช่องในไทล์
ภายในไทล์ ช่องของเธรดก็คือ threadIdx ของเธรดนั้น ส่วนแถวและคอลัมน์ในระดับโกลบอลยังคงมาจากดัชนีบล็อกและเธรด
int ty = threadIdx.y, tx = threadIdx.x;
int row = blockIdx.y*TILE + ty;
int col = blockIdx.x*TILE + tx;การโหลดไทล์ของ A
แต่ละเธรดจะคัดลอกองค์ประกอบหนึ่งของไทล์ปัจจุบันจาก A ไปยัง หน่วยความจำร่วม ทั้งบล็อกจะช่วยกันเตรียมบล็อกของ A ขนาด TILE คูณ TILE ให้ครบ
As[ty][tx] = A[row*N + (phase*TILE + tx)];การโหลดไทล์ของ B
ในเวลาเดียวกัน แต่ละเธรดจะโหลดองค์ประกอบหนึ่งของไทล์จาก B ตอนนี้ ไทล์ ทั้งสองอยู่บนชิปและพร้อมสำหรับการอ่านซ้ำอย่างรวดเร็ว
Bs[ty][tx] = B[(phase*TILE + ty)*N + col];ซิงโครไนซ์ก่อนคำนวณ
เรียกใช้ __syncthreads() เพื่อให้เธรดทุกตัวโหลดเสร็จก่อนที่จะมีใครอ่านไทล์ หากข้ามขั้นตอนนี้ ผลลัพธ์จะเป็นข้อมูลขยะ
__syncthreads();คำนวณบนไทล์
ตอนนี้แต่ละเธรดจะวนลูปสั้น ๆ ผ่าน ไทล์ โดยอ่านจากหน่วยความจำร่วมเท่านั้น การอ่านเหล่านี้มีต้นทุนต่ำกว่าการอ่านจากหน่วยความจำโกลบอลอย่างมาก
for (int k = 0; k < TILE; ++k)
sum += As[ty][k] * Bs[k][tx];ประโยชน์จากการใช้ซ้ำ
ค่าที่โหลดแต่ละค่าจะถูกใช้โดย TILE เธรดแทนที่จะใช้โดยเธรดเดียว การใช้ข้อมูลซ้ำ นี้คือเหตุผลทั้งหมดที่ทำให้การคูณเมทริกซ์แบบแบ่งไทล์ทำงานได้รวดเร็ว
ไทล์เดียวไม่เพียงพอ
ไทล์เดียวครอบคลุมผลคูณจุดได้เพียงบางส่วน คุณต้องทำขั้นตอนโหลด-ซิงโครไนซ์-คำนวณซ้ำผ่านหลาย เฟส ซึ่งบทเรียนถัดไปจะอธิบาย
ตรวจสอบอย่างรวดเร็ว
ทบทวนลำดับขั้นตอนเมื่อทำงานกับไทล์ร่วม
สรุป
คุณเตรียมไทล์ของ A และ B ลงใน หน่วยความจำร่วม ซิงโครไนซ์ แล้วคำนวณด้วยการอ่านจากบนชิปที่มีต้นทุนต่ำ การใช้ซ้ำคือข้อได้เปรียบ ขั้นต่อไปคือเฟสต่าง ๆ 🧱
คำถามที่พบบ่อย
บทเรียน “แบ่งผลคูณภายในเป็นไทล์” ฟรีหรือไม่
ใช่ — ข้อความเต็มของ “แบ่งผลคูณภายในเป็นไทล์” ฟรีให้อ่านที่นี่บนเว็บ เพื่อปฏิบัติแบบโต้ตอบ (ตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7) และปลดล็อคส่วนที่เหลือของคอร์ส CUDA Academy ให้อัปเกรดเป็น CoddyKit PRO คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน
คุณจะเรียนรู้อะไรในบทเรียน “แบ่งผลคูณภายในเป็นไทล์”
โหลดไทล์ย่อยของ A และ B ในแต่ละช่วง คุณปฏิบัติ CUDA Academy ด้วยโค้ดที่ใช้งานได้จริงที่คุณเรียกใช้โดยตรงในเบราว์เซอร์ และติวเตอร์ AI ตลอด 24/7 ตอบคำถามของคุณขณะที่คุณไปผ่านบทเรียน
คุณต้องมีประสบการณ์ก่อนที่จะเริ่มเรียน CUDA Academy หรือไม่
ไม่จำเป็นต้องมีประสบการณ์มาก่อน CUDA Academy บน CoddyKit ออกแบบมาสำหรับผู้เริ่มต้นไปจนถึงผู้เรียนขั้นสูง คุณสามารถเริ่มต้นที่นี่หรือเริ่มจากตัวแรกและเรียนด้วยความเร็วของคุณเอง นี่คือบทเรียนที่ 2 จากทั้งหมด 4 บทเรียน
บทเรียน “แบ่งผลคูณภายในเป็นไทล์” ใช้เวลานานแค่ไหน
บทเรียน CoddyKit ส่วนใหญ่ใช้เวลาประมาณ 5–10 นาที แต่ละบทเรียนจึงสั้นและเป็นแบบโต้ตอบ คุณสามารถก้าวหน้าอย่างต่อเนื่องและกลับมาเรียนต่อจากตรงที่เพิ่งหยุดบนเว็บและแอปได้เลย
ฉันเขียนและรันโค้ดในบทเรียน CUDA Academy นี้ได้ไหม
ได้ บทเรียน CUDA Academy ทุกบทมีตัวแก้ไขโค้ดในตัว คุณจึงเขียนและรันโค้ดจริงได้เลยในเบราว์เซอร์ และได้รับข้อเสนอแนะจาก AI ในทันที — ไม่ต้องติดตั้งในเครื่องของคุณ
บทเรียนทั้งหมดในหลักสูตรนี้
- เคอร์เนลคูณเมทริกซ์แบบพื้นฐาน
- แบ่งผลคูณภายในเป็นไทล์
- วนซ้ำตามช่วงของไทล์
- วัดการเพิ่มความเร็ว