GPU আর্কিটেকচার — হাজারো সরল কোরের জগৎ

আগে এটা বুঝি

গত লেসনে আমরা SIMD দেখেছি — একটা CPU core-এর ভেতরে একটা AVX register 8টা float একসাথে যোগ করে। একই instruction, একসাথে বহু data — কিন্তু সেই “বহু” সংখ্যাটা ছিল ৪ থেকে ১৬।

এখন প্রশ্নটা করি: যদি “বহু” মানে ৮ না হয়ে ১০,০০০ হয়?

একটা 4096 × 4096 matrix multiply-এ প্রায় ৬৮৭ কোটি multiply-add লাগে। একটা 1920 × 1080 ছবির প্রতিটা pixel-এ যদি একটা filter চালাতে হয়, সেটা ২০ লক্ষ স্বাধীন হিসাব। প্রতিটা হিসাব একে অপরের থেকে সম্পূর্ণ স্বাধীন — ঠিক map pattern-এর মতো, যা আমরা পরের লেসনে নাম দেব।

এই স্কেলে CPU-র approach ভেঙে পড়ে। একটা আধুনিক CPU-তে ৪ থেকে ৬৪টা core, প্রতিটা জটিল: out-of-order execution (আগের লেসন), branch prediction, বড় cache, deep pipeline — প্রতিটা transistor একটামাত্র instruction stream-কে যতটা সম্ভব দ্রুত চালানোর জন্য উৎসর্গ করা। ২০ লক্ষ independent pixel-এর জন্য ৬৪টা জটিল কোর h ব্যবহার করলে প্রতিটাকে ~৩১,০০০টা pixel সিরিয়ালি করতে হয়।

GPU সম্পূর্ণ ভিন্ন বাজি ধরে: হাজার হাজার সরল, in-order কোর। কোনো branch prediction নেই (বা খুবই সাধারণ), কোনো out-of-order execution নেই, ছোট cache। প্রতিটা কোর এককভাবে CPU কোরের চেয়ে অনেক ধীর আর বোকা — কিন্তু একটা NVIDIA A100 GPU-তে ৬,৯১২টা এমন কোর আছে। একই silicon budget, সম্পূর্ণ ভিন্ন সিদ্ধান্ত: latency optimize করার বদলে throughput optimize করা।

একই প্রশ্নের দুইটা উত্তর

“সেকেন্ডে কত কাজ করতে পারি?” — CPU উত্তর দেয়: “কয়েকটা কাজ, কিন্তু প্রতিটা যত দ্রুত সম্ভব, আর জটিল নির্ভরশীলতা সামলাতে পারি (branch, speculation, cache)।” GPU উত্তর দেয়: “একটামাত্র কাজ ধীরে করব, কিন্তু একসাথে হাজার হাজার কাজ করব, যদি সেগুলো একে অপরের থেকে স্বাধীন হয়।”

GPU CPU-র প্রতিস্থাপন নয় — এটা একটা different point on the same design space, যে design space-এর ভেতরে SIMD একটা ছোট পদক্ষেপ ছিল, GPU সেই একই পদক্ষেপের চরম রূপ।

এই লেসনে আমরা দেখব GPU হার্ডওয়্যার আসলে কেমন গঠিত, কোন শব্দগুলো (SM, warp, SIMT) কী বোঝায়, আর কোথায় এই মডেলটা ভেঙে পড়ে (warp divergence)। পরের লেসনে আমরা এই হার্ডওয়্যারকে সফটওয়্যার দিয়ে কীভাবে address করা হয় (CUDA) সেটা দেখব।

মূল ধারণা

Streaming Multiprocessor (SM)

GPU-র মূল building block একটা core একা নয় — একটা Streaming Multiprocessor (SM)। একটা SM হলো সরল কোরের একটা cluster যারা control logic আর দ্রুত on-chip memory ভাগাভাগি করে।

একটা SM-এর ভেতরে থাকে:

বহু CUDA core (NVIDIA-র শব্দ সরল ALU-র জন্য — integer/float add, multiply করতে পারে, কিন্তু একা কোনো instruction fetch/decode করতে পারে না)
একটা বা কয়েকটা warp scheduler (কোন warp-টা এই cycle-এ চলবে সিদ্ধান্ত নেয়)
Register file — হাজার হাজার thread-এর জন্য register রাখার জায়গা
Shared memory / L1 cache (একই SRAM, programmable ভাগে ভাগ করা যায় — এই লেসনের পরের অংশে বিস্তারিত)
বিশেষায়িত unit — Tensor core (matrix multiply), RT core (ray tracing), Special Function Unit (sin/cos/sqrt)

একটা NVIDIA A100 GPU-তে ১০৮টা SM আছে, প্রতিটায় ৬৪টা FP32 CUDA core — মোট ৬,৯১২টা। তুলনায় একটা ৬৪-core সার্ভার CPU-তে ৬৪টা সম্পূর্ণ, স্বাধীন, out-of-order core। সংখ্যাটা ~১০৮× বেশি, কিন্তু প্রতিটা কোর ব্যাপকভাবে সরল।

CPU die:                          GPU die:
┌─────┬─────┬─────┬─────┐        ┌──┬──┬──┬──┬──┬──┬──┬──┐
│ কোর │ কোর │ কোর │ কোর │        ├──┼──┼──┼──┼──┼──┼──┼──┤
│বড়  │বড়  │বড়  │বড়  │  L3    │  SM (৬৪+ সরল কোর প্রতিটায়)  │
│OoO  │OoO  │OoO  │OoO  │ cache  ├──┼──┼──┼──┼──┼──┼──┼──┤
│bpred│bpred│bpred│bpred│        │  ... আরো ১০০+ SM ...        │
└─────┴─────┴─────┴─────┘        └──┴──┴──┴──┴──┴──┴──┴──┘
৪–৬৪টা জটিল কোর                   হাজার হাজার সরল কোর
transistor খরচ: control/         transistor খরচ: ALU সংখ্যা
speculation/cache-এ               বাড়ানোয়, control সরল রাখা

CPU বনাম GPU — একই silicon বাজেট, ভিন্ন বন্টন।

Warp — SIMD-এর GPU নাম

এখানেই lesson 10-এর সাথে সরাসরি যোগসূত্র। একটা SM তার কোরগুলোকে এককভাবে চালায় না — সে থ্রেডগুলোকে warp-এ ভাগ করে, সাধারণত ৩২টা thread-এর একটা দল। একটা warp-এর সবগুলো thread ঠিক একই instruction, একই cycle-এ, lockstep-এ চালায়।

এটা লক্ষ্য করুন: এটা হুবহু SIMD-এর ধারণা — “একটা instruction, বহু data lane” — শুধু ভিন্ন শব্দভান্ডারে। AVX-এ lane গুলো একটা 256-বিট register-এর ভেতরে bit-field; GPU-তে lane গুলোকে “thread” বলা হয়, প্রতিটার নিজস্ব register আছে বলে মনে হয়। কিন্তু hardware level-এ warp scheduler একটামাত্র instruction fetch/decode করে, আর সেই একই ALU operation ৩২টা lane-এ broadcast করে — precisely SIMD execution।

৩২ সংখ্যাটা arbitrary নয় — এটা NVIDIA hardware-এর একটা fixed physical property, যেমন AVX-512-এর register 512 বিট fixed। (AMD GPU-তে সমতুল্য “wavefront” ৬৪ thread-এর, ধারণাটা একই।)

SIMT — কেন GPU একে “SIMD” বলে না

NVIDIA ইচ্ছাকৃতভাবে এই মডেলটাকে SIMT (Single Instruction, Multiple Threads) বলে, SIMD নয় — আর এই পার্থক্যটা genuinely গুরুত্বপূর্ণ, প্রায়ই ভুল বোঝা হয়।

SIMD (CPU-তে): প্রোগ্রামার সরাসরি একটা wide register নিয়ে কাজ করে। vaddps ymm0, ymm1, ymm2 — এটা স্পষ্টভাবে “৮টা float একসাথে যোগ করো” বলছে। প্রোগ্রামারকে জানতে হয় সে vector data নিয়ে কাজ করছে।

SIMT (GPU-তে): প্রোগ্রামার ৩২টা independent thread লেখে, প্রতিটার নিজস্ব register, নিজস্ব program counter আছে বলে মনে হয়, নিজস্ব if decision নিতে পারে বলে মনে হয়। কোড দেখতে সাধারণ scalar কোডের মতো:

__global__ void add_one(int *data) {
    int i = threadIdx.x;
    data[i] = data[i] + 1;   // প্রতিটা thread "স্বাধীনভাবে" লেখে বলে মনে হয়
}

কিন্তু হার্ডওয়্যার নিচে ঠিক SIMD-র মতোই কাজ করে — একটা warp-এর ৩২টা thread একই cycle-এ একই instruction চালায়, একটা shared instruction fetch/decode ব্যবহার করে। SIMT হলো একটা abstraction layer: “independent thread” নামক illusion, যার নিচে “lockstep SIMD execution” নামক বাস্তবতা।

SIMT vs SIMD — নির্ভুল পার্থক্য

	SIMD	SIMT
প্রোগ্রামিং মডেল	Explicit vector — programmer জানে সে lane নিয়ে কাজ করছে	Implicit — programmer একটা scalar thread লেখে
Register	একটা shared wide register	প্রতিটা thread-এর নিজস্ব register (hardware-এ বরাদ্দ)
Divergence	Compile-time-এ ধরা পড়ে, programmer manually predicate লেখে	Runtime-এ ঘটে, hardware automatically handle করে (masking)
নাম যা বোঝায়	Data-এর উপর জোর	Thread-এর illusion-এর উপর জোর

মূল কথা: SIMT = “SIMD execution, কিন্তু thread-এর ছদ্মবেশে পরিবেশন করা, যাতে প্রোগ্রামারকে explicit vectorization নিয়ে ভাবতে না হয়।” এটা একটা genuinely আলাদা programming model হলেও, performance-এর দিক থেকে সেই একই SIMD সীমাবদ্ধতাগুলো (lockstep, divergence cost) নিচে বহাল থাকে।

ভেতরে কী ঘটছে

Warp Divergence — যেখানে illusion ভেঙে পড়ে

SIMT-এর illusion বলে প্রতিটা thread স্বাধীন, তার নিজের if নিতে পারে। কিন্তু হার্ডওয়্যার একটাই instruction stream চালাতে পারে প্রতি cycle-এ পুরো warp-এর জন্য। তাহলে কী হয় যখন একটা warp-এর কিছু thread if নেয়, বাকিরা else?

__global__ void branchy(int *data) {
    int i = threadIdx.x;
    if (data[i] % 2 == 0) {
        data[i] = expensive_even(data[i]);   // পথ A
    } else {
        data[i] = expensive_odd(data[i]);    // পথ B
    }
}

ধরুন একটা warp-এর ৩২টা thread-এর মধ্যে ১৬টার data জোড় সংখ্যা, ১৬টার বিজোড়। হার্ডওয়্যার দুটো পথই সিরিয়ালি চালায়:

প্রথম pass: পথ A চালানো হয়, কিন্তু শুধু সেই ১৬টা lane-এর জন্য যাদের condition true — বাকি ১৬টা lane mask out করা হয় (তাদের ALU idle থাকে, কিন্তু cycle খরচ হয়)
দ্বিতীয় pass: পথ B চালানো হয়, এবার উল্টো ১৬টা lane active, আগেরটা mask out
দুই পথ শেষে ফলাফল merge হয়, warp আবার একসাথে চলা শুরু করে

Warp-এর ৩২টা lane, condition অনুযায়ী রঙ করা (■ = true, □ = false)

Divergence ছাড়া (সব lane একমত):
  cycle 1: ■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■   ← এক pass, সব active

Divergence সহ (অর্ধেক-অর্ধেক):
  cycle 1: ■■■■■■■■■■■■■■■□□□□□□□□□□□□□□□□□   ← path A, ১৬টা lane idle
  cycle 2: □□□□□□□□□□□□□□□□■■■■■■■■■■■■■■■■   ← path B, বাকি ১৬টা idle
                                                    মোট = ২× কাজ, সময়ও ২×

Warp divergence — দুটো পথই সিরিয়ালি, প্রতিবার অর্ধেক lane অলস।

সবচেয়ে খারাপ ক্ষেত্রে — যদি warp-এর প্রতিটা thread ভিন্ন পথ নেয় (৩২-ওয়ে switch statement, প্রতি thread আলাদা case) — হার্ডওয়্যারকে সবগুলো পথ সিরিয়ালি চালাতে হয়, প্রতিটায় মাত্র ১/৩২ lane active। থিওরিটিক্যাল worst case: ৩২ গুণ ধীর, কোনো divergence না থাকলে যা লাগত তার তুলনায়।

Branch prediction-এর সাথে গুলিয়ে ফেলবেন না

computer-architecture/branch-prediction লেসনে আমরা দেখেছি CPU branch predict করে, ভুল হলে speculative কাজ বাতিল করে আর সঠিক পথে ফিরে যায় — misprediction penalty pipeline-এর গভীরতার সমানুপাতিক (~১৫–২০ cycle)।

GPU warp divergence সম্পূর্ণ ভিন্ন mechanism, ভিন্ন cost model:

	Branch misprediction (CPU)	Warp divergence (GPU)
কারণ	ভুল অনুমান	একই warp-এ ভিন্ন thread ভিন্ন সিদ্ধান্ত নেয়
প্রতিকার	ভালো predictor, branch hint	কোড পুনর্গঠন — একই warp-এ একমুখী condition
খরচ	pipeline flush, একবার (~১৫–২০ cycle)	উভয় পথ চালানো, protruding — condition-এর জটিলতার সমানুপাতিক
Predict করে এড়ানো যায়?	হ্যাঁ, ভালো predictor দিয়ে	না — divergence deterministic, prediction অপ্রাসঙ্গিক

CPU-তে ভালো prediction দিয়ে branch cost প্রায় শূন্যে নামানো যায়। GPU-তে divergence prediction দিয়ে এড়ানো যায় না — একমাত্র উপায় কোড এমনভাবে লেখা যাতে একই warp-এর thread-রা একই পথে যায়।

বাস্তব প্রতিকার: ডেটা এমনভাবে সাজান/sort করুন যাতে একই warp-এর thread-দের condition একই হয় (data-driven divergence কমানো), অথবা branch-heavy কাজ পুরোপুরি GPU-তে না দিয়ে সেই অংশটা CPU-তে রাখুন। Ray tracing-এ এই সমস্যা কুখ্যাত — প্রতিটা ray ভিন্ন বস্তুতে আঘাত করে, ভিন্ন material shader চালায় — যে কারণে RT core-এর মতো বিশেষায়িত hardware দরকার হয়েছে।

GPU-র memory hierarchy

CPU-র মতোই GPU-র memory hierarchy আছে, কিন্তু একটা মৌলিক দার্শনিক পার্থক্য সহ: CPU cache automatic, GPU shared memory programmer-managed।

একটা GPU thread-এর memory access, দ্রুততম থেকে ধীরতম

Register~১ cycle — প্রতিটা thread-এর নিজস্ব, hardware বরাদ্দ করে
Shared memory~২০–৩০ cycle — SM-এর ভেতরে on-chip SRAM, একটা block-এর সব thread ভাগ করে
L1 / texture cache~৩০ cycle — shared memory-র সাথে একই SRAM, ভিন্নভাবে কনফিগার করা
L2 cache~২০০ cycle — পুরো GPU জুড়ে ভাগাভাগি
Global memory (GDDR/HBM)~৪০০–৮০০ cycle — GPU-র নিজস্ব DRAM, CPU-র DRAM-এর সমতুল্য

Global memory — CPU-র DRAM-এর সরাসরি সমতুল্য। GPU-র নিজস্ব DRAM chip (GDDR6X বা HBM2e), CPU-র system RAM থেকে সম্পূর্ণ আলাদা, PCIe বা NVLink দিয়ে সংযুক্ত। বড়, কিন্তু ধীর — একটা global memory access ৪০০–৮০০ cycle লাগতে পারে, CPU-র DRAM access (~২০০ cycle, computer-architecture/memory-hierarchy লেসন দ্রষ্টব্য)-এর চেয়েও বেশি ধীর, আপেক্ষিকভাবে (GPU clock সাধারণত CPU-র চেয়ে কম, ~১.৫ GHz বনাম ~৪ GHz)।

Shared memory — প্রতিটা SM-এর নিজস্ব ছোট (সাধারণত ৪৮–১৬৪ KB), অত্যন্ত দ্রুত (~২০–৩০ cycle) on-chip SRAM। একই SM-এ চলা একটা thread block-এর সব thread এটা ভাগাভাগি করে। ধারণাগতভাবে CPU-র L1 cache-এর মতো (ছোট, দ্রুত, কাছে) — কিন্তু মৌলিক পার্থক্য: এটা automatic নয়।

CPU cache বনাম GPU shared memory — automatic বনাম manual

CPU-র L1/L2/L3 cache transparent — প্রোগ্রামার কিছু লেখে না, hardware নিজে থেকে সিদ্ধান্ত নেয় কোন cache line রাখবে, কখন evict করবে (computer-architecture/cache-organization দ্রষ্টব্য)। প্রোগ্রামার শুধু access pattern ভালো রাখার চেষ্টা করে (locality), সরাসরি নিয়ন্ত্রণ করে না।

GPU shared memory সম্পূর্ণ programmer-managed — আপনি explicitly __shared__ keyword দিয়ে ঘোষণা করেন কী data সেখানে যাবে, নিজে সেখানে copy করেন global memory থেকে, নিজে __syncthreads() দিয়ে নিশ্চিত করেন সব thread copy শেষ করেছে তারপর ব্যবহার করে (পরের লেসনে বিস্তারিত)।

এই পার্থক্যের কারণ: GPU-তে হাজার হাজার concurrent thread-এর জন্য hardware-চালিত cache replacement policy ব্যবহারিকভাবে খুবই জটিল ও ব্যয়বহুল হতো। তার বদলে নিয়ন্ত্রণটা প্রোগ্রামারের হাতে দেওয়া হয়েছে — বেশি কাজ, কিন্তু predictable performance।

Latency hiding — GPU কেন cache-এর উপর কম নির্ভর করে

CPU গভীরভাবে speculative execution আর branch prediction-এর উপর নির্ভর করে memory latency লুকাতে (আগের লেসনগুলো)। GPU সম্পূর্ণ ভিন্ন কৌশল ব্যবহার করে: massive thread-level parallelism দিয়ে latency hide করা, speculation দিয়ে নয়।

একটা SM-এ একসাথে হাজার হাজার thread “resident” থাকতে পারে (তাদের সবার register বরাদ্দ করা আছে)। যখন একটা warp global memory থেকে data-র জন্য অপেক্ষা করে (~৫০০ cycle), warp scheduler সেটাকে আলাদা রেখে অন্য একটা প্রস্তুত warp চালানো শুরু করে — বিনা খরচে context switch, কারণ প্রতিটা warp-এর register আলাদাভাবে hardware-এ বরাদ্দ, cache করা নয়। যথেষ্ট warp resident থাকলে GPU প্রায় কখনো idle বসে থাকে না, যদিও প্রতিটা individual memory access ধীর।

এটাকেই বলে occupancy — একটা SM কতগুলো warp একসাথে resident রাখতে পারছে, তার শতাংশ। বেশি occupancy মানে বেশি latency hiding সম্ভাবনা। কিন্তু register বা shared memory বেশি ব্যবহার করলে occupancy কমে যায় (কারণ একটা SM-এর মোট register/shared memory budget fixed) — এটাই CUDA performance tuning-এর একটা কেন্দ্রীয় trade-off, পরের লেসনগুলোয় ফিরে আসবে।

উদাহরণ

একটা concrete উদাহরণ — ইমেজ ব্রাইটনেস বাড়ানো

ধরুন একটা 1920 × 1080 ছবির প্রতিটা pixel-এর brightness ৫০ বাড়াতে হবে। মোট ২০,৭৩,৬০০ টা স্বাধীন হিসাব, প্রতিটা pixel[i] = min(255, pixel[i] + 50)।

CPU-তে (single core, scalar): ২০ লক্ষ iteration, প্রতিটা কয়েক nanosecond। মোট ~৫–১০ ms।

CPU-তে SIMD (AVX2, ৮-wide): ২০ লক্ষ / ৮ ≈ ২,৫৯,০০০ iteration। ~১–২ ms।

GPU-তে: ২০,৭৩,৬০০ pixel-কে ৬৪-thread block-এ ভাগ করলে ~৩২,৪০০টা block লাগে। একটা মাঝারি GPU-তে (~৫০টা SM, প্রতিটায় ৪টা warp scheduler) এই সব block একসাথে বা প্রায় একসাথে চলে, কারণ প্রতিটা pixel-এর কাজ ছোট আর সম্পূর্ণ independent — কোনো warp divergence নেই (সব thread একই কাজ করছে, min operation predication-এই সস্তায় হয়)। ফলাফল: ~০.১ ms বা তার কম — মূলত data GPU-তে transfer করার সময়ই dominant খরচ হয়ে দাঁড়ায় (PCIe-র মাধ্যমে, যা পরের লেসনে cudaMemcpy আলোচনায় গুরুত্বপূর্ণ হবে)।

কেন GPU এত ভালো এখানে: এই কাজটা ঠিক GPU যা ভালোবাসে — বিশাল সংখ্যক, ছোট, independent, branch-free, uniform কাজ। এটাই map pattern (পরের-পরের লেসনে formalize করব), আর SIMD-এর lesson-এর সাথে সরাসরি যোগসূত্র: এটা সেই একই “একই instruction, বহু data” ধারণা, শুধু ৮-wide থেকে হাজার-wide-এ স্কেল করা।

নিজে চালিয়ে দেখুন

EXPERIMENT

Warp Divergence Simulator — divergence-এর cost সরাসরি দেখুন

Python 3 (কোনো GPU দরকার নেই)· ২০ মিনিট

আসল GPU হার্ডওয়্যার নেই এমন যে কারো জন্য, আমরা warp scheduler-এর আচরণ simulate করতে পারি: একটা warp-এর ৩২টা “lane”-এর প্রতিটার জন্য একটা predicate (branch condition) দেওয়া থাকবে, আমরা গুনব কয়টা সিরিয়াল pass লাগবে সবাইকে সব পথ চালাতে, আর কত শতাংশ lane-cycle সত্যিকারের কাজে ব্যয় হলো (বাকিটা mask-out-এ নষ্ট)।

import random

WARP_SIZE = 32

def simulate_warp(predicates, n_paths=2):
    """
    predicates: প্রতিটা lane কোন path নেবে (0..n_paths-1)-এর একটা তালিকা।
    রিটার্ন করে: (মোট pass সংখ্যা, utilization %)
    """
    passes = 0
    total_lane_cycles = 0
    active_lane_cycles = 0

    unique_paths = set(predicates)
    for path in unique_paths:
        passes += 1
        active = sum(1 for p in predicates if p == path)
        total_lane_cycles += WARP_SIZE      # পুরো warp সবসময় "চলে", মাস্ক থাকুক বা না থাকুক
        active_lane_cycles += active        # শুধু matching lane-ই আসল কাজ করে

    utilization = active_lane_cycles / total_lane_cycles * 100
    return passes, utilization


def scenario(name, predicates):
    passes, util = simulate_warp(predicates)
    print(f"{name:35s}  pass={passes}   utilization={util:5.1f}%")


# দৃশ্য ১ — কোনো divergence নেই, সব lane একমত
scenario("সব lane একই path (uniform)", [0] * WARP_SIZE)

# দৃশ্য ২ — অর্ধেক-অর্ধেক
scenario("৫০/৫০ split (if/else)", [0]*16 + [1]*16)

# দৃশ্য ৩ — সম্পূর্ণ random binary condition (data % 2)
random.seed(42)
data = [random.randint(0, 1000) for _ in range(WARP_SIZE)]
scenario("data % 2 (গড়ে ৫০/৫০, বাস্তব)", [d % 2 for d in data])

# দৃশ্য ৪ — worst case, প্রতিটা lane আলাদা path (৩২-ওয়ে switch)
scenario("worst case — সবাই আলাদা path", list(range(WARP_SIZE)))

# দৃশ্য ৫ — একটা outlier, বাকি সব uniform (বাস্তবে খুব সাধারণ:
# boundary check, "যদি এটাই শেষ element হয়")
scenario("৩১ uniform + ১টা boundary-case", [0]*31 + [1])

প্রত্যাশিত output:

সব lane একই path (uniform)          pass=1   utilization=100.0%
৫০/৫০ split (if/else)                pass=2   utilization= 50.0%
data % 2 (গড়ে ৫০/৫০, বাস্তব)         pass=2   utilization= 50.0%
worst case — সবাই আলাদা path         pass=32  utilization=  3.1%
৩১ uniform + ১টা boundary-case       pass=2   utilization= 51.6%

লক্ষ্য করার মতো: শেষ দৃশ্যে মাত্র একটা outlier lane পুরো warp-কে দুই pass-এ বাধ্য করছে — utilization প্রায় অর্ধেকে নেমে যায়, শুধু একটা lane-এর জন্য। এটাই বাস্তব কোডে সবচেয়ে সাধারণ divergence pattern: array-এর শেষ প্রান্তে boundary check, বা কোনো বিরল error condition — একটামাত্র thread ভিন্ন পথ নিলেও পুরো warp-কে সেই খরচ বহন করতে হয়।

নিজে বাড়ান: n_paths বাড়িয়ে দেখুন worst-case utilization কীভাবে ১/n_paths-এর দিকে যায়। আর একটা nested-if কেস simulate করুন (dvergence-এর ভেতরে আরেকটা divergence) — pass সংখ্যা কি গুণিতক হারে বাড়ে?

যাদের সত্যিকারের GPU আছে

CUDA toolkit ইনস্টল থাকলে নিচের kernel-টা compile করে nvprof বা nsys দিয়ে profile করুন — warp_execution_efficiency metric সরাসরি এই utilization সংখ্যাটাই দেখায়:

nsys profile --stats=true ./divergent_kernel

তুলনা করুন uniform branch (if (blockIdx.x == 0), পুরো block-এর জন্য একই) বনাম per-thread branch (if (threadIdx.x % 2 == 0)) — প্রথমটায় কোনো warp-level divergence নেই (পুরো warp একই সিদ্ধান্ত নেয়), দ্বিতীয়টায় প্রতিটা warp-এ ঠিক ৫০/৫০ split।

এটা কী প্রমাণ করে

একই কাজ, শুধু condition-এর distribution পাল্টে, warp-এর সিরিয়াল pass সংখ্যা আর কার্যকর utilization কীভাবে বদলায় — এটাই divergence-এর প্রকৃত ব্যয়।

নিজে বানান

BUILD IT

Occupancy Estimator

Python · ●●○○○

একটা SM-এর সম্পদ সীমা মডেল করুন — max thread, max register, max shared memory
একটা kernel-এর per-thread register ব্যবহার ও per-block shared memory ইনপুট নিন
হিসাব করুন কতগুলো block একসাথে একটা SM-এ resident থাকতে পারবে
সেখান থেকে occupancy শতাংশ বের করুন এবং bottleneck (register নাকি shared memory নাকি thread-slot) চিহ্নিত করুন

বাস্তব CUDA occupancy calculator (NVIDIA নিজেই একটা spreadsheet সরবরাহ করে) যা করে, তার একটা সরলীকৃত সংস্করণ বানাই। ধারণাটা: একটা SM-এর তিনটা স্বাধীন সীমা আছে — যেটা প্রথমে শেষ হয়ে যায়, সেটাই bottleneck।

# একটা টিপিক্যাল আধুনিক GPU SM-এর সীমা (Ampere-এর কাছাকাছি সংখ্যা)
SM_MAX_THREADS = 2048
SM_MAX_BLOCKS = 32
SM_MAX_REGISTERS = 65536
SM_MAX_SHARED_MEM_BYTES = 164 * 1024
WARP_SIZE = 32


def occupancy(threads_per_block, regs_per_thread, shared_mem_per_block):
    # সীমা ১ — thread slot দিয়ে কতগুলো block আঁটবে
    blocks_by_threads = SM_MAX_THREADS // threads_per_block

    # সীমা ২ — register budget দিয়ে
    regs_per_block = regs_per_thread * threads_per_block
    blocks_by_regs = SM_MAX_REGISTERS // regs_per_block if regs_per_block > 0 else SM_MAX_BLOCKS

    # সীমা ৩ — shared memory budget দিয়ে
    blocks_by_shared = (SM_MAX_SHARED_MEM_BYTES // shared_mem_per_block
                         if shared_mem_per_block > 0 else SM_MAX_BLOCKS)

    # সীমা ৪ — hardware-এর নিজস্ব max block সংখ্যা
    active_blocks = min(blocks_by_threads, blocks_by_regs, blocks_by_shared, SM_MAX_BLOCKS)
    active_warps = active_blocks * (threads_per_block // WARP_SIZE)
    max_warps = SM_MAX_THREADS // WARP_SIZE

    occ = active_warps / max_warps * 100

    limits = {
        "thread slot": blocks_by_threads,
        "register": blocks_by_regs,
        "shared memory": blocks_by_shared,
    }
    bottleneck = min(limits, key=limits.get)

    print(f"block/SM = {active_blocks}, warp/SM = {active_warps}/{max_warps}, "
          f"occupancy = {occ:.1f}%   (bottleneck: {bottleneck})")
    return occ


print("কেস ১ — হালকা kernel (৩২ রেজিস্টার, শেয়ার্ড মেমরি নেই):")
occupancy(threads_per_block=256, regs_per_thread=32, shared_mem_per_block=0)

print("\nকেস ২ — রেজিস্টার-ক্ষুধার্ত kernel (১২৮ রেজিস্টার/thread):")
occupancy(threads_per_block=256, regs_per_thread=128, shared_mem_per_block=0)

print("\nকেস ৩ — shared-memory-ভারী kernel (৪৮KB/block):")
occupancy(threads_per_block=256, regs_per_thread=32, shared_mem_per_block=48*1024)

প্রত্যাশিত output:

কেস ১ — হালকা kernel (৩২ রেজিস্টার, শেয়ার্ড মেমরি নেই):
block/SM = 8, warp/SM = 64/64, occupancy = 100.0%   (bottleneck: thread slot)

কেস ২ — রেজিস্টার-ক্ষুধার্ত kernel (১২৮ রেজিস্টার/thread):
block/SM = 2, warp/SM = 16/64, occupancy = 25.0%   (bottleneck: register)

কেস ৩ — shared-memory-ভারী kernel (৪৮KB/block):
block/SM = 3, warp/SM = 24/64, occupancy = 37.5%   (bottleneck: shared memory)

মূল শিক্ষা: বেশি register ব্যবহার করা একটা kernel-কে দ্রুত করতে পারে (কম memory access, বেশি data register-এ রাখা যায়) — কিন্তু সেই একই register ব্যবহারই occupancy কমিয়ে দিতে পারে, ফলে latency hide করার মতো যথেষ্ট warp resident থাকে না। এটা একটা সত্যিকারের trade-off, “যত বেশি তত ভালো” নয় — ঠিক যেমন lesson ৩-৯-এ cache/register allocation নিয়ে দেখা trade-off গুলো।

নিজে বাড়ান: threads_per_block ৩২-এর গুণিতক না হলে কী হয়, সেটা মডেলে যোগ করুন (partial warp = নষ্ট lane, ঠিক occupancy calculator-এর মতো)।

বাস্তব সিস্টেমে

GPU আর্কিটেকচার যেখানে সিদ্ধান্ত নেয়

Deep learning training/inference। একটা matrix multiply (GEMM) — যা neural network-এর প্রতিটা layer-এর মূল হিসাব — পুরোপুরি uniform, branch-free, আর massively parallel। এজন্যই NVIDIA-র Tensor Core (একটা বিশেষ SM unit যা 4×4 matrix multiply একটা instruction-এ করে) modern AI hardware-এর কেন্দ্রে। H100 GPU-র Tensor Core FP16 matmul-এ ~১০০০ TFLOPS দিতে পারে, সাধারণ CUDA core FP32-এর তুলনায় ~৩০× বেশি।

Cryptocurrency mining (ঐতিহাসিক)। Bitcoin-এর SHA-256 hashing প্রতিটা nonce-এর জন্য সম্পূর্ণ স্বাধীন, branch-free, uniform কাজ — GPU mining-এর স্বর্ণযুগ (২০১০–১৩) তৈরি করেছিল। পরে ASIC (একেবারে সেই একটা hash function-এর জন্য বানানো chip) GPU-কেও হারিয়ে দেয়, কারণ ASIC-এ general-purpose flexibility-র কোনো transistor খরচ নেই — শুধু SHA-256।

Ray tracing। এখানে GPU-র সবচেয়ে বড় শত্রু — warp divergence। প্রতিটা ray ভিন্ন বস্তুতে আঘাত করে, ভিন্ন material shader চালায়, ভিন্ন গভীরতায় bounce করে। একই warp-এর ৩২টা ray প্রায়ই সম্পূর্ণ ভিন্ন কোড পথ নেয়। এই সমস্যা সমাধানে NVIDIA RT core যোগ করেছে — bounding-volume-hierarchy traversal-এর জন্য বিশেষায়িত fixed-function hardware, যাতে সেই divergence-ভারী অংশটা general SIMT core-এ না গিয়ে আলাদা, optimized সার্কিটে চলে।

Video encoding। H.264/HEVC encoding-এর motion estimation ধাপ — প্রতিটা 16×16 pixel block-এর জন্য স্বাধীন search — একটা classic map pattern, GPU/dedicated encoder ASIC (NVENC)-এর জন্য উপযুক্ত।

বৈজ্ঞানিক সিমুলেশন। Molecular dynamics-এ প্রতিটা কণার উপর বলের হিসাব বাকি সব কণা থেকে (বা কাছের কণাগুলো থেকে) স্বাধীনভাবে করা যায় — N-body simulation classic GPU workload, uniform force calculation, ন্যূনতম branching।

ডেটাবেস analytics (GPU-accelerated)। RAPIDS/cuDF-এর মতো framework GPU-তে column-wise scan, filter, aggregate চালায় — WHERE age > 30-এর মতো filter প্রতিটা row-এ uniform, predicate হিসেবে সহজে map হয়।

যেখানে GPU ব্যবহার হয় না। একটা web server-এর request routing, একটা compiler-এর syntax parsing, একটা OS scheduler — এগুলো branch-heavy, data-dependent, sequential নির্ভরতাযুক্ত কাজ। এখানে CPU-র জটিল branch predictor আর out-of-order execution-ই সঠিক হাতিয়ার। GPU CPU-কে প্রতিস্থাপন করেনি — heterogeneous computing-এ দুটো একসাথে কাজ করে, CPU serial/branchy অংশ সামলায়, GPU massive uniform parallel অংশ।

যে ভুলগুলো সবাই করে

“GPU-তে হাজার হাজার 'core' মানে CPU-র মতোই হাজার হাজার স্বাধীন কোর।”

CUDA core CPU core-এর মতো নয়। একটা CPU core সম্পূর্ণ স্বাধীনভাবে নিজস্ব instruction fetch, decode, branch prediction করতে পারে। একটা CUDA core একটা সরল ALU — সে নিজে কোনো instruction fetch/decode করে না; সেটা করে warp scheduler, একবারে পুরো warp-এর (৩২ core) জন্য।

তাই বাস্তবে একটা GPU-র “স্বাধীন instruction stream” সংখ্যা তার CUDA core সংখ্যা নয় — এটা তার warp scheduler সংখ্যা (প্রতি SM-এ সাধারণত ৪টা)। একটা A100-এ ১০৮ SM × ৪ scheduler = ৪৩২টা সত্যিকারের স্বাধীন instruction stream, ৬,৯১২টা “core” নয়।

সঠিক মানসিক মডেল: GPU = কয়েকশ স্বাধীন SIMD unit, প্রতিটা ৩২-wide, CPU-র AVX-এর মতোই — শুধু সংখ্যায় অনেক বেশি।

“বেশি thread launch করলেই সবসময় বেশি speed পাওয়া যাবে।”

Occupancy-র উপরের build exercise ঠিক এই ভুলটা দেখায়। GPU-র প্রতিটা SM-এর register file আর shared memory fixed, সীমিত। বেশি thread launch করলে যদি প্রতিটা thread বেশি register চায়, তাহলে একই সময়ে কম thread resident থাকতে পারবে — occupancy কমে যায়, latency hide করার ক্ষমতা কমে যায়, performance কমতেও পারে।

আরেকটা সীমা: ৩২-এর গুণিতক না এমন block size (যেমন ৫০) launch করলে প্রতিটা warp-এর একটা অংশ সবসময় অব্যবহৃত থাকে (৫০ thread মানে ২ warp = ৬৪ lane বরাদ্দ, ১৪টা lane সবসময় idle)। সবসময় ৩২-এর গুণিতক block size বেছে নেওয়া উচিত।

নিয়ম: থ্রেড সংখ্যা GPU-র সম্পদ বাজেটের সাপেক্ষে বেছে নিতে হয়, “যত বেশি তত ভালো” নয়।

“Warp divergence শুধু 'একটু ধীর' করে দেয়, বড় সমস্যা না।”

Experiment section-এর simulation দেখিয়েছে worst-case utilization ৩.১%-এ নামতে পারে — মানে ৩২ গুণ ধীর, শুধু একটা branch-এর কারণে। এটা কোনো ছোট প্রভাব নয়; এটা এমন কিছু যা একটা algorithm-কে GPU-র জন্য সম্পূর্ণভাবে অনুপযুক্ত করে দিতে পারে।

আর সবচেয়ে বিপজ্জনক রূপ সেটা নয় যেটা চোখে পড়ে (৫০/৫০ split) — সেটা যেটা লুকিয়ে থাকে: একটা বিরল boundary check, একটা error-handling path, একটা if (threadIdx.x == 0) guard (যা প্রায়ই দেখা যায় যখন শুধু একটা thread কোনো shared variable initialize করে) — এগুলো প্রতিটা warp-কে একটা বাড়তি pass করতে বাধ্য করে, ছোট দেখতে হলেও।

বুঝেছেন কি না দেখুন

SM, warp, আর SIMT — তিনটা শব্দের সংজ্ঞা এক বাক্যে দিন, আর বলুন কোনটা hardware-এর একটা physical অংশ, কোনটা একটা execution group, আর কোনটা একটা programming/execution model-এর নাম।

স্মরণ

SM (Streaming Multiprocessor) — hardware-এর physical অংশ: সরল কোরের একটা cluster যারা control logic ও on-chip memory ভাগ করে (register file, shared memory, warp scheduler)।

Warp — একটা execution group: ৩২টা thread-এর একটা দল যারা একই instruction একই cycle-এ lockstep-এ চালায়। এটা hardware-এর একটা fixed গঠন (৩২ সংখ্যাটা GPU vendor-নির্দিষ্ট, পরিবর্তনযোগ্য নয়) কিন্তু সফটওয়্যারে সরাসরি দেখা যায় না — CUDA প্রোগ্রামার warp manually তৈরি করে না, compiler/runtime automatically thread-দের warp-এ ভাগ করে।

SIMT — একটা execution model-এর নাম: “independent thread”-এর illusion দিয়ে SIMD lockstep execution বাস্তবায়ন করার পদ্ধতি। এটা কোনো physical জিনিস নয়, বরং hardware আর programming model-এর মধ্যে সম্পর্কটা বর্ণনা করার একটা concept।

একটা warp-এ ৩২টা thread আছে। একটা kernel-এ একটা switch statement আছে ৪টা case-এর সাথে, আর warp-এর thread-গুলো এভাবে ভাগ হয়েছে: case ০-এ ২০টা thread, case ১-এ ৮টা, case ২-এ ৩টা, case ৩-এ ১টা। এই warp-টা সম্পূর্ণ switch চালাতে কয়টা সিরিয়াল pass লাগবে, আর গড় utilization কত?

প্রয়োগ

হার্ডওয়্যারকে প্রতিটা unique path আলাদাভাবে চালাতে হয়, সেই path-এ কতগুলো thread আছে তা বিবেচ্য না — এমনকি একটা thread হলেও।

Pass সংখ্যা = unique case সংখ্যা = ৪।

প্রতিটা pass-এ পুরো warp-এর ৩২টা lane “চলে” (হয় active নয়তো masked), কিন্তু শুধু matching thread-গুলোই আসল কাজ করে:

Pass	Active lane	Idle (masked) lane
১ (case ০)	২০	১২
২ (case ১)	৮	২৪
৩ (case ২)	৩	২৯
৪ (case ৩)	১	৩১

মোট lane-cycle খরচ হলো ৪ pass × ৩২ lane = ১২৮ lane-cycle। আসল কাজে ব্যয় হলো ২০+৮+৩+১ = ৩২ lane-cycle (যা হওয়ারই কথা, কারণ মোট thread ৩২টাই)।

$\text{Utilization} = \frac{32}{128} = 25\%$

তুলনা করুন: যদি এই ৩২টা thread ৪টা আলাদা warp-এ ভাগ করা যেত (প্রতিটা warp-এ শুধু একই case-এর thread), তাহলে প্রতিটা warp-ই ১টা pass-এ, ১০০% utilization-এ শেষ হতো। এটাই মূল শিক্ষা: divergence-এর cost data-র warp-এর ভেতরে বিতরণের উপর নির্ভর করে, মোট কাজের উপর নয় — আর এজন্যই ডেটা sort/group করে warp-uniform করাটা একটা বাস্তব optimization কৌশল।

CPU speculative execution/branch prediction দিয়ে memory latency আর branch cost লুকায় (আগের লেসনগুলো)। GPU সেটা করে না বললেই চলে — তাহলে GPU memory latency (৪০০–৮০০ cycle!) কীভাবে সামলায়?

যুক্তি

GPU speculation দিয়ে নয়, massive parallelism দিয়ে latency হাইড করে — এটাকে বলে latency hiding।

একটা SM-এ একসাথে হাজার হাজার thread (বহু warp) “resident” থাকতে পারে, প্রতিটার নিজস্ব register hardware-এ বরাদ্দ (cache করা নয়, তাই context switch বিনামূল্যে)। যখন একটা warp memory-র জন্য অপেক্ষা করে, warp scheduler সেই cycle-এই অন্য একটা প্রস্তুত warp চালানো শুরু করে — কোনো stall নেই, কোনো speculation-এর দরকার নেই, কারণ ভুল অনুমান করার প্রশ্নই নেই।

এটা কার্যকর হওয়ার শর্ত: যথেষ্ট independent warp থাকতে হবে memory latency ঢাকার জন্য। যদি ৫০০ cycle memory latency লুকাতে হয় আর প্রতিটা warp গড়ে ২৫ cycle কাজ করে memory-র জন্য অপেক্ষা করার আগে, তাহলে ৫০০/২৫ = ২০টা independent warp দরকার সেই latency সম্পূর্ণ হাইড করতে — এটাই occupancy-র গুরুত্বের মূল কারণ।

দার্শনিক তুলনা: CPU বলে “আমি অনুমান করব ভবিষ্যতে কী ঘটবে, আর ভুল হলে ফিরে আসব” (speculation)। GPU বলে “আমি অপেক্ষা করব না — অন্য কোনো স্বাধীন কাজ তো আছেই, সেটা করি এই ফাঁকে” (parallelism)। দুটোই একই সমস্যার (memory latency) সমাধান, সম্পূর্ণ ভিন্ন দর্শনে — ঠিক যেমন CPU/GPU-র সামগ্রিক design philosophy ভিন্ন।

আপনাকে একটা algorithm দেওয়া হয়েছে যা একটা ১ কোটি-নোডের গ্রাফে প্রতিটা নোডের জন্য একটা ভিন্ন, জটিল, data-dependent traversal করে (কিছু নোডে গভীর recursion, কিছুতে অগভীর, condition-গুলো সম্পূর্ণ unpredictable)। এটা GPU-তে চালানো কি ভালো সিদ্ধান্ত? যুক্তি দিন।

ডিজাইন

সম্ভবত না — অন্তত সরলভাবে নয়। এই সমস্যাটা GPU-র শক্তির বিপরীত দিকে পড়ে, প্রতিটা মাপকাঠিতে:

১. Uniformity নেই। প্রতিটা নোডের কাজ ভিন্ন (গভীরতা, condition) — একই warp-এর ৩২টা thread প্রায় নিশ্চিতভাবে ভিন্ন পথ নেবে। Warp divergence খুব বেশি হবে, সম্ভবত worst-case-এর কাছাকাছি (Question ২-এর মতো), utilization হয়তো ১০%-এর নিচে নামবে।

২. Memory access pattern অনির্দিষ্ট। Graph traversal মানে pointer-chasing ধরনের access (mathematics/asymptotic-notation লেসনের linked-list উদাহরণ মনে করুন) — GPU-র global memory latency (৪০০–৮০০ cycle) CPU-র DRAM-এর চেয়েও বেশি, আর latency hiding-এর জন্য দরকার predictable, coalesced access — random pointer chasing ঠিক তার বিপরীত।

৩. Recursion depth ভিন্ন ভিন্ন। GPU thread-এর stack ছোট আর সীমিত; গভীর, unpredictable recursion GPU kernel-এ সমস্যাযুক্ত (stack overflow বা বিশাল stack বরাদ্দ প্রয়োজন, যা register/occupancy budget খেয়ে ফেলে)।

বিকল্প চিন্তা:

যদি গ্রাফটা কোনোভাবে restructure করা যায় (level-order BFS-এর মতো layer-ভিত্তিক processing, যেখানে একই layer-এর সব নোড একই ধরনের কাজ করে) — তাহলে প্রতিটা layer আলাদা kernel launch হিসেবে GPU-তে uniform ভাবে চলতে পারে। এটাই বাস্তবে GPU BFS implementation-গুলো করে।
নাহলে CPU-ই সঠিক হাতিয়ার — branch prediction, বড় cache, আর flexible out-of-order execution ঠিক এই ধরনের irregular, data-dependent কাজের জন্য optimize করা।

সাধারণ নীতি: GPU-তে migrate করার আগে প্রশ্ন করুন — কাজটা কি বিশাল সংখ্যক, ছোট, uniform, branch-light, predictable-access একক-এ ভাঙা যায়? উত্তর “না” হলে, algorithm পুনর্গঠন ছাড়া GPU ব্যবহার প্রায়ই ক্ষতিকর (data transfer overhead + খারাপ utilization মিলে CPU-র চেয়েও ধীর হতে পারে)।

“GPU shared memory অনেকটা CPU-র L1 cache-এর মতো” — এই তুলনাটা কোন অর্থে সত্য, আর কোন অর্থে বিভ্রান্তিকর?

যুক্তি

সত্য অর্থে: দুটোই ছোট, on-chip, দ্রুত (২০–৩০ cycle পরিসরে), আর তাদের সংশ্লিষ্ট প্রসেসিং ইউনিটের (SM / CPU core) কাছাকাছি বসানো, যাতে বারবার-ব্যবহৃত ডেটা ধীর main memory-তে বারবার যেতে না হয়। উভয়ই memory hierarchy-র সেই একই মৌলিক নীতির প্রয়োগ: locality exploit করে গড় access latency কমানো।

বিভ্রান্তিকর অর্থে — এবং এটাই গুরুত্বপূর্ণ পার্থক্য:

	CPU L1 cache	GPU shared memory
নিয়ন্ত্রণ	Hardware automatic — programmer কিছু লেখে না	Programmer explicit — `__shared__` দিয়ে ঘোষণা, manual copy
Replacement	Hardware policy (LRU-এর কাছাকাছি)	কোনো policy নেই — programmer সিদ্ধান্ত নেয় কী রাখবে, কতক্ষণ
Coherence	Automatic (MESI protocol, আগের লেসন)	Explicit synchronization লাগে (`__syncthreads()`)
Scope	পুরো memory hierarchy-র transparent অংশ	শুধু একটা thread block-এর মধ্যে দৃশ্যমান
ভুল হলে	Cache miss — ধীর, কিন্তু সবসময় সঠিক ফলাফল	Race condition সম্ভব — ভুল ফলাফল, শুধু ধীর নয়

সবচেয়ে গুরুত্বপূর্ণ পার্থক্য শেষ সারিতে: CPU cache ভুল হলে শুধু performance এ ভুগবেন (miss, ধীর) — ফলাফল সবসময় সঠিক থাকবে, কারণ hardware coherence guarantee করে। GPU shared memory ভুলভাবে ব্যবহার করলে (যেমন __syncthreads() ছাড়া লেখা আর পড়া) আপনি ভুল উত্তর পেতে পারেন — এটা একটা correctness সমস্যা, শুধু performance সমস্যা নয়। এই মেকানিজম পরের লেসনে বিস্তারিত দেখব।

এরপর কী

আমরা এখন জানি GPU হার্ডওয়্যার কেমন গঠিত — SM, warp, SIMT, memory hierarchy, আর divergence-এর cost। কিন্তু এখনও দেখিনি একজন প্রোগ্রামার আসলে এই হার্ডওয়্যারকে কীভাবে address করেন — কীভাবে বলা হয় “এই ২০ লক্ষ pixel-এর প্রতিটার জন্য একটা thread চালাও”।

পরের লেসনে আমরা CUDA দেখব — একটা __global__ kernel function কীভাবে লেখা হয়, launch syntax দিয়ে কীভাবে thread/block/grid সংখ্যা নির্দিষ্ট করা হয়, threadIdx/blockIdx/blockDim দিয়ে প্রতিটা thread কীভাবে তার নিজের কাজের ভাগ বুঝে নেয়, আর কীভাবে __shared__ memory ব্যবহার করে এই লেসনে দেখা shared memory-কে সত্যিকারের কোডে রূপান্তর করা হয়। একটা সম্পূর্ণ, শুরু থেকে শেষ পর্যন্ত কাজ করা vector-addition প্রোগ্রাম লিখব — host code থেকে device memory allocation থেকে kernel launch পর্যন্ত।

GPU আর্কিটেকচার — হাজারো সরল কোরের জগৎ

এই লেসন শেষে আপনি পারবেন

আগে যা বোঝা থাকা দরকার

আগে এটা বুঝি

মূল ধারণা

Streaming Multiprocessor (SM)

Warp — SIMD-এর GPU নাম

SIMT — কেন GPU একে “SIMD” বলে না

ভেতরে কী ঘটছে

Warp Divergence — যেখানে illusion ভেঙে পড়ে

GPU-র memory hierarchy

Latency hiding — GPU কেন cache-এর উপর কম নির্ভর করে

উদাহরণ

একটা concrete উদাহরণ — ইমেজ ব্রাইটনেস বাড়ানো

নিজে চালিয়ে দেখুন

Warp Divergence Simulator — divergence-এর cost সরাসরি দেখুন

নিজে বানান

Occupancy Estimator

বাস্তব সিস্টেমে

GPU আর্কিটেকচার যেখানে সিদ্ধান্ত নেয়

যে ভুলগুলো সবাই করে

বুঝেছেন কি না দেখুন

এরপর কী

এরপর কী

আরও পড়ুন

এই লেসন শেষে আপনি পারবেন

আগে যা বোঝা থাকা দরকার

◔আগে এটা বুঝিWhy this matters

◆মূল ধারণাCore concept

Streaming Multiprocessor (SM)

Warp — SIMD-এর GPU নাম

SIMT — কেন GPU একে “SIMD” বলে না

⚙ভেতরে কী ঘটছেUnder the hood

Warp Divergence — যেখানে illusion ভেঙে পড়ে

GPU-র memory hierarchy

Latency hiding — GPU কেন cache-এর উপর কম নির্ভর করে

▶উদাহরণExample

একটা concrete উদাহরণ — ইমেজ ব্রাইটনেস বাড়ানো

⌘নিজে চালিয়ে দেখুনExperiment

Warp Divergence Simulator — divergence-এর cost সরাসরি দেখুন

⚒নিজে বানানBuild it yourself

Occupancy Estimator

◉বাস্তব সিস্টেমেReal-world usage

GPU আর্কিটেকচার যেখানে সিদ্ধান্ত নেয়

⚠যে ভুলগুলো সবাই করেMisconceptions

?বুঝেছেন কি না দেখুনCheck understanding

→এরপর কীNext step

এরপর কী

আরও পড়ুন

আগে এটা বুঝি

মূল ধারণা

ভেতরে কী ঘটছে

উদাহরণ

নিজে চালিয়ে দেখুন

নিজে বানান

বাস্তব সিস্টেমে

যে ভুলগুলো সবাই করে

বুঝেছেন কি না দেখুন

এরপর কী