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

สร้างฮิสโตแกรม

ใช้การดำเนินการอะตอมิกร่วมกับการแยกข้อมูลในหน่วยความจำร่วม

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

บางส่วนของบทเรียนนี้ยังไม่ได้รับการแปล และแสดงเป็นภาษาอังกฤษ

What a Histogram Counts

A histogram counts how many inputs fall into each bin. Many threads will want to increment the same bin, so this is an atomics problem. 📊

The Naive Approach

Each thread reads one element, finds its bin, and increments that bin. Without protection, popular bins lose counts to races.

The Global Atomic Version

The simplest correct fix is one atomicAdd per element straight into global memory. It works, but hot bins serialize threads.

atomicAdd(&hist[bin], 1);

The Contention Problem

When data clusters into a few bins, thousands of threads pile onto the same address. That contention can make global atomics painfully slow.

Privatization to the Rescue

Privatization gives each block its own private histogram in fast shared memory. Threads collide only inside their block, not across the whole grid.

Declare the Shared Histogram

Each block declares a shared array sized to the number of bins. It lives on-chip, so atomics there are far cheaper than global ones.

__shared__ int local[NBINS];

Step 1: Clear the Bins

Threads cooperatively zero the shared histogram, then call __syncthreads so no one counts before clearing is done.

local[tid] = 0;
__syncthreads();

Step 2: Count Locally

Now each thread atomically bumps its bin in shared memory. Same atomicAdd, but on the fast on-chip copy instead of global.

atomicAdd(&local[bin], 1);

Step 3: Merge to Global

After a barrier, threads add each shared bin into the global histogram with one atomicAdd per bin. Far fewer global atomics than before.

atomicAdd(&hist[i], local[i]);

Why This Is Faster

Shared-memory atomics are quick, and the costly global atomics now fire once per bin per block instead of once per element.

Watch the Bin Count

The private histogram must fit in shared memory. With too many bins, split them into passes or fall back to global atomics.

Quick Check

One question on the histogram strategy.

Recap: Building a Histogram

You built a histogram with global atomics, then sped it up using shared-memory privatization: clear, count locally, merge. ✅

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

บทเรียน “สร้างฮิสโตแกรม” ฟรีหรือไม่

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

คุณจะเรียนรู้อะไรในบทเรียน “สร้างฮิสโตแกรม”

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

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

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

บทเรียน “สร้างฮิสโตแกรม” ใช้เวลานานแค่ไหน

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

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

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

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

  1. เงื่อนไขแข่งขันบน GPU
  2. atomicAdd และฟังก์ชันที่เกี่ยวข้อง
  3. สร้างฮิสโตแกรม
  4. การดำเนินการอะตอมิกแบบกำหนดเองด้วย atomicCAS
← กลับไปที่ CUDA Academy