บล็อกเธรดเป็นนามธรรมการเขียนโปรแกรมที่แสดงถึงกลุ่มของเธรดที่สามารถดำเนินการแบบอนุกรมหรือแบบขนานได้ เพื่อการแมปกระบวนการและข้อมูล ที่ดีขึ้น เธรดจึงถูกจัดกลุ่มเป็นบล็อกเธรด จำนวนเธรดในบล็อกเธรดเดิมถูกจำกัดโดยสถาปัตยกรรมไว้ที่ 512 เธรดต่อบล็อก แต่ตั้งแต่เดือนมีนาคม 2010 ด้วยความสามารถในการประมวลผล 2.x ขึ้นไป บล็อกอาจมีเธรดได้มากถึง 1024 เธรด เธรดในบล็อกเธรดเดียวกันทำงานบนมัลติโปรเซสเซอร์แบบสตรีมเดียวกัน^{[ 1 ]}เธรดในบล็อกเดียวกันสามารถสื่อสารกันได้ผ่านหน่วยความจำที่ใช้ร่วมกัน การซิงโครไนซ์แบบ กั้น หรือพรีมิที ฟการซิงโครไนซ์อื่นๆ เช่น การดำเนินการอะตอมิก

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

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

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

เนื่องจากแอปพลิเคชันแบบขนานจำนวนมากเกี่ยวข้องกับข้อมูลหลายมิติ จึงสะดวกที่จะจัดระเบียบกลุ่มเธรดเป็นอาร์เรย์เธรดแบบ 1 มิติ 2 มิติ หรือ 3 มิติ กลุ่มเธรดในกริดจะต้องสามารถดำเนินการได้อย่างอิสระ เนื่องจากการสื่อสารหรือความร่วมมือระหว่างกลุ่มเธรดในกริดนั้นเป็นไปไม่ได้ 'เมื่อมีการเรียกใช้เคอร์เนล จำนวนเธรดต่อกลุ่มเธรดและจำนวนกลุ่มเธรดจะถูกระบุ ซึ่งจะกำหนดจำนวนเธรด CUDA ทั้งหมดที่เรียกใช้^{[ 2 ]} ' มิติ x, y และ z สูงสุดของกลุ่มเธรดคือ 1024, 1024 และ 64 และควรจัดสรรให้ x × y × z ≤ 1024 ซึ่งเป็นจำนวนเธรดสูงสุดต่อกลุ่มเธรด^{[ 3 ]}กลุ่มเธรดสามารถจัดระเบียบเป็นกริดหนึ่ง สอง หรือสามมิติได้สูงสุด 2 ³¹ -1, 65,535 และ 65,535 กลุ่มในมิติ x, y และ z ตามลำดับ^{[ 3 ]}ต่างจากจำนวนเธรดสูงสุดต่อบล็อก ไม่มีข้อจำกัดจำนวนบล็อกต่อกริดที่แยกจากขนาดกริดสูงสุด

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

ลองพิจารณาตัวอย่างที่มีอาร์เรย์ขนาด 512 องค์ประกอบ โครงสร้างการจัดระเบียบอย่างหนึ่งคือการใช้กริดที่มีบล็อกเดียวซึ่งมี 512 เธรด สมมติว่ามีอาร์เรย์ C ขนาด 512 องค์ประกอบ ซึ่งได้มาจากการคูณแบบองค์ประกอบของอาร์เรย์ A และ B ซึ่งแต่ละอาร์เรย์มีขนาด 512 องค์ประกอบเท่ากัน แต่ละเธรดมีดัชนี i และทำการคูณองค์ประกอบที่ i ^ของ A และ B จากนั้นเก็บผลลัพธ์ไว้ในองค์ประกอบที่ i ^ของ C โดย i คำนวณจาก blockIdx (ซึ่งในกรณีนี้คือ 0 เนื่องจากมีเพียงบล็อกเดียว), blockDim (512 ในกรณีนี้เนื่องจากบล็อกมี 512 องค์ประกอบ) และ threadIdx ซึ่งแตกต่างกันไปตั้งแต่ 0 ถึง 511 สำหรับแต่ละบล็อก

ดัชนีเกลียว i คำนวณได้จากสูตรต่อไปนี้:

${\displaystyle i=blockIdx.x*blockDim.x+threadIdx.x}$

blockIdx.x คือตัวระบุบล็อกมิติ x

blockDim.x คือมิติ x ของมิติบล็อก

threadIdx.x คือมิติ x ของตัวระบุเธรด

ดังนั้นค่าของ 'i' จะอยู่ในช่วงตั้งแต่ 0 ถึง 511 ซึ่งครอบคลุมทั้งอาร์เรย์

หากเราต้องการพิจารณาการคำนวณสำหรับอาร์เรย์ที่มีขนาดใหญ่กว่า 1024 เราสามารถใช้บล็อกหลายบล็อก โดยแต่ละบล็อกมี 1024 เธรด ลองพิจารณาตัวอย่างที่มีองค์ประกอบในอาร์เรย์ 2048 ตัว ในกรณีนี้ เราจะมีบล็อกเธรด 2 บล็อก โดยแต่ละบล็อกมี 1024 เธรด ดังนั้นค่าของตัวระบุเธรดจะแตกต่างกันไปตั้งแต่ 0 ถึง 1023 ตัวระบุบล็อกจะแตกต่างกันไปตั้งแต่ 0 ถึง 1 และมิติของบล็อกจะเป็น 1024 ดังนั้นบล็อกแรกจะมีค่าดัชนีตั้งแต่ 0 ถึง 1023 และบล็อกสุดท้ายจะมีค่าดัชนีตั้งแต่ 1024 ถึง 2047

ดังนั้นแต่ละเธรดจะคำนวณดัชนีของหน่วยความจำที่ต้องเข้าถึงก่อน จากนั้นจึงดำเนินการคำนวณต่อไป พิจารณาตัวอย่างที่องค์ประกอบจากอาร์เรย์ A และ B ถูกบวกแบบขนานโดยใช้เธรด และผลลัพธ์จะถูกเก็บไว้ในอาร์เรย์ C โค้ดที่เกี่ยวข้องในเธรดแสดงอยู่ด้านล่าง: ^{[ 5 ]}

__global__ void vecAddKernel ( float * A , float * B , float * C , int n ) { int index = blockIdx . x * blockDim . x + threadIdx . x ; if ( index < n ) { C [ index ] = A [ index ] + B [ index ] ; } }

ในทำนองเดียวกัน ในกริดที่มีความซับซ้อนเป็นพิเศษ จำเป็นต้องคำนวณ blockId และ threadId สำหรับแต่ละเธรด โดยขึ้นอยู่กับรูปทรงเรขาคณิตของกริด ลองพิจารณากริด 2 มิติที่มีบล็อก 2 มิติ threadId และ blockId จะถูกคำนวณโดยใช้สูตรต่อไปนี้:

${\displaystyle blockId=blockIdx.x+blockIdx.y*gridDim.x;}$${\displaystyle threadId=blockId*(blockDim.x*blockDim.y)+(threadIdx.y*blockDim.x)+threadIdx.x;}$^{[ 6 ]}

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

แต่ละสถาปัตยกรรมใน GPU (เช่นKeplerหรือFermi ) ประกอบด้วย SM หรือ Streaming Multiprocessors หลายตัว ซึ่งเป็นโปรเซสเซอร์อเนกประสงค์ที่มี อัตรา ความเร็วสัญญาณนาฬิกา ต่ำ และแคชขนาดเล็ก SM สามารถประมวลผลบล็อกเธรดหลายบล็อกพร้อมกันได้ เมื่อบล็อกเธรดหนึ่งเสร็จสิ้นการทำงานแล้ว ก็จะรับบล็อกเธรดถัดไปตามลำดับ โดยทั่วไปแล้ว SM รองรับการทำงานแบบขนานในระดับคำสั่งแต่ไม่รองรับ การคาด การณ์การแตกแขนง^{[ 8 ]}

เพื่อให้บรรลุวัตถุประสงค์นี้ SM ประกอบด้วยสิ่งต่อไปนี้: ^{[ 8 ]}

• หน่วยประมวลผล (หน่วยประมวลผลเลขทศนิยมความแม่นยำเดี่ยว, หน่วยประมวลผลเลขทศนิยมความแม่นยำคู่, หน่วยประมวลผลฟังก์ชันพิเศษ (SFU))
• แคช:

1. แคช L1 (สำหรับลดความล่าช้าในการเข้าถึงหน่วยความจำ)
2. หน่วยความจำร่วม (สำหรับข้อมูลที่ใช้ร่วมกันระหว่างเธรด)
3. แคชคงที่ (สำหรับกระจายการอ่านจากหน่วยความจำแบบอ่านอย่างเดียว )
4. แคชพื้นผิว (สำหรับรวบรวมแบนด์วิดท์จากหน่วยความจำพื้นผิว)

• ตัวกำหนดตารางเวลาสำหรับวาร์ป (ใช้สำหรับออกคำสั่งไปยังวาร์ปตามนโยบายการกำหนดตารางเวลาเฉพาะ)
• ต้องมีรีจิสเตอร์จำนวนมาก (SM อาจกำลังทำงานด้วยเธรดจำนวนมากในเวลาเดียวกัน ดังนั้นจึงจำเป็นต้องมีรีจิสเตอร์หลายพันตัว)

ฮาร์ดแวร์จะจัดสรรบล็อกเธรดให้กับ SM (Single Manager) โดยทั่วไปแล้ว SM หนึ่งตัวสามารถจัดการบล็อกเธรดได้หลายบล็อกพร้อมกัน SM หนึ่งตัวอาจมีบล็อกเธรดได้มากถึง 8 บล็อก โดย SM ที่เกี่ยวข้องจะกำหนด ID ให้กับเธรดแต่ละตัว

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

ตัวกำหนดตารางเวลาวาร์ปของ SM จะตัดสินใจว่าวาร์ปใดจะได้รับความสำคัญเป็นอันดับแรกในระหว่างการออกคำสั่ง^{[ 11 ]}นโยบายการจัดลำดับความสำคัญของวาร์ปบางส่วนได้มีการกล่าวถึงในส่วนต่อไปนี้ด้วย

ในด้านฮาร์ดแวร์ บล็อกเธรดประกอบด้วย 'วาร์ป' (คำนี้มาจากการทอผ้า [ ^{12 ] )} วาร์ปคือชุดของเธรด 32 เธรดภายในบล็อกเธรด ในอดีต เธรดเหล่านี้รับประกันว่าจะทำงาน "พร้อมกัน" (เธรดทั้งหมดภายในวาร์ปทำงานคำสั่งพร้อมกัน) และที่สำคัญคือสามารถเข้าถึงตำแหน่งหน่วยความจำทุกตำแหน่งด้วยเธรดวาร์ปทั้งหมดหรือไม่มีเลย พฤติกรรมนี้อาจนำไปสู่ภาวะเดดล็อกได้ง่าย (เช่น โดยการใช้ if-branch ในลูป) อย่างไรก็ตาม ตั้งแต่สถาปัตยกรรม Voltaเป็นต้นมา การแลกเปลี่ยนข้อมูลภายในวาร์ปผ่านการล็อกที่ละเอียดกว่านั้นเป็นไปได้^{[ 13 ] [ 14 ]}เธรดเหล่านี้จะถูกเลือกตามลำดับโดย SM ^{[ 15 ]}

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

พิจารณา warp ของเธรด 32 เธรดที่กำลังดำเนินการคำสั่ง หากตัวถูกดำเนินการหนึ่งตัวหรือทั้งสองตัวยังไม่พร้อม (เช่น ยังไม่ได้ถูกดึงมาจากหน่วยความจำส่วนกลาง) กระบวนการที่เรียกว่า ' การสลับบริบท ' จะเกิดขึ้น ซึ่งจะถ่ายโอนการควบคุมไปยัง warp อื่น^{[ 16 ]}เมื่อสลับออกจาก warp ใด warp หนึ่ง ข้อมูลทั้งหมดของ warp นั้นจะยังคงอยู่ในไฟล์รีจิสเตอร์เพื่อให้สามารถกลับมาทำงานต่อได้อย่างรวดเร็วเมื่อตัวถูกดำเนินการพร้อม เมื่อคำสั่งไม่มีการพึ่งพาข้อมูลที่ค้างอยู่ กล่าวคือ ตัวถูกดำเนินการทั้งสองตัวพร้อมแล้ว warp ที่เกี่ยวข้องจะถือว่าพร้อมสำหรับการดำเนินการ หากมี warp มากกว่าหนึ่ง warp ที่มีสิทธิ์ในการดำเนินการ SM หลักจะใช้นโยบายการจัดกำหนดการ warp เพื่อตัดสินใจว่า warp ใดจะได้รับคำสั่งที่ดึงมาถัดไป

นโยบายที่แตกต่างกันสำหรับการจัดกำหนดการวาร์ปที่มีสิทธิ์ในการดำเนินการจะกล่าวถึงด้านล่าง: ^{[ 17 ]}

1. Round Robin (RR) - คำสั่งต่างๆ จะถูกดึงมาในลักษณะวนรอบ RR ช่วยให้มั่นใจได้ว่า SM (หน่วยความจำหลัก) จะทำงานอยู่ตลอดเวลา และไม่เสียเวลาไปกับความล่าช้าของหน่วยความจำ
2. ลำดับความสำคัญในการดึงคำสั่ง คือ Least Recently Fetched (LRF) - ในนโยบายนี้ วาร์ปที่ไม่ได้ถูกดึงคำสั่งมานานที่สุดจะได้รับความสำคัญในการดึงคำสั่ง
3. ยุติธรรม (FAIR) ^{[ 17 ]} - ในนโยบายนี้ ตัวกำหนดตารางเวลาจะทำให้แน่ใจว่า warps ทั้งหมดได้รับโอกาสที่ 'ยุติธรรม' ในจำนวนคำสั่งที่ดึงมาสำหรับพวกมัน โดยจะดึงคำสั่งไปยัง warp ที่มีจำนวนคำสั่งที่ดึงมาน้อยที่สุด
4. CAWS แบบบล็อกเธรด^{[ 18 ]} (การจัดกำหนดการวาร์ปที่คำนึงถึงความสำคัญ) - นโยบายการจัดกำหนดการนี้เน้นที่การปรับปรุงเวลาการดำเนินการของบล็อกเธรด โดยจัดสรรทรัพยากรเวลาเพิ่มเติมให้กับวาร์ปที่ใช้เวลานานที่สุดในการดำเนินการ ด้วยการให้ความสำคัญกับวาร์ปที่สำคัญที่สุด นโยบายนี้ช่วยให้บล็อกเธรดเสร็จสิ้นเร็วขึ้น ทำให้ทรัพยากรพร้อมใช้งานได้เร็วขึ้น

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

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