Foundationপ্রথম নীতি থেকে
LEVEL 11লেসন ১৩/১৭কঠিন১ ঘণ্টা

Parallel Patterns — Map, Reduce, Scan

Parallel Patterns: Map, Reduce, Scan

MAP (স্বাধীন, trivially parallel), REDUCE (associative combine, tree দিয়ে log n depth), SCAN (দেখতে সিরিয়াল, কিন্তু Hillis-Steele দিয়ে log n depth) — তিনটা reusable pattern যা দিয়ে বেশিরভাগ বাস্তব parallel প্রোগ্রাম বানানো হয়, single GPU থেকে পুরো data center পর্যন্ত।

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

  • MAP pattern-এর সংজ্ঞা দিতে ও চিনতে পারবেন কেন এটা trivially parallel
  • REDUCE-এর জন্য tree-structured combination আঁকতে ও তার work/depth জটিলতা বিশ্লেষণ করতে পারবেন
  • Scan (prefix sum) কেন সিরিয়াল মনে হয়েও parallel করা যায়, সেটা Hillis-Steele algorithm দিয়ে ব্যাখ্যা করতে পারবেন
  • কোনো operator reduce-এর জন্য উপযুক্ত কি না (associativity) যাচাই করতে পারবেন
  • কেন বাস্তব parallel প্রোগ্রাম হাতে-লেখা custom algorithm-এর বদলে এই তিনটা pattern কম্পোজ করে বানানো হয়, সেটা যুক্তি দিয়ে বলতে পারবেন

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

আগে এটা বুঝি

গত লেসনে দুইটা kernel লিখেছিলাম। vecAdd — প্রতিটা thread সম্পূর্ণ স্বাধীন, একজন আরেকজনের data স্পর্শ করে না। blockSum — থ্রেডদের একটা tree-আকৃতির সহযোগিতায় একটা সংখ্যায় একত্রিত হতে হয়।

এই দুটো একে অপরের থেকে গঠনগতভাবে আলাদা, আর এই পার্থক্যটা এত মৌলিক আর এত ঘন ঘন দেখা যায় যে computer science-এ এদের নিজস্ব নাম দেওয়া হয়েছে: map আর reduce। আজ আমরা তৃতীয় একটা pattern দেখব — scan — যেটা প্রথম দেখায় সিরিয়াল মনে হয় (প্রতিটা running total আগেরটার উপর নির্ভর করে বলে মনে হয়), কিন্তু আসলে একটা চমৎকার, পাল্টা-স্বজ্ঞাত (counter-intuitive) parallel algorithm আছে।

কেন এই তিনটা নাম করে শেখা জরুরি? কারণ বেশিরভাগ বাস্তব parallel প্রোগ্রাম নতুন algorithm হাতে-উদ্ভাবন করে বানানো হয় না — এই তিনটা pattern-কে কম্পোজ করে বানানো হয়। একবার আপনি এই তিনটা চিনতে পারলে, একটা নতুন সমস্যা দেখেই বুঝতে পারবেন কোন pattern (বা কয়েকটার সমন্বয়) কাজে লাগবে — নতুন করে সমান্তরালতা নিয়ে চিন্তা করতে হবে না প্রতিবার।

মূল ধারণা

MAP — কোনো নির্ভরতা নেই

সংজ্ঞা: একটা একই function f প্রতিটা input element-এ স্বাধীনভাবে প্রয়োগ করা, একটা সমান আকারের output তৈরি করতে।

map(f,[x1,x2,…,xn])=[f(x1),f(x2),…,f(xn)]\text{map}(f, [x_1, x_2, \ldots, x_n]) = [f(x_1), f(x_2), \ldots, f(x_n)]

কোনো element-এর হিসাব অন্য কোনো element-এর উপর নির্ভর করে না — এজন্যই একে বলা হয় embarrassingly parallel (বিব্রতকরভাবে সহজ, কারণ সমান্তরাল করার জন্য কোনো clever কৌশল লাগে না, শুধু nটা independent কাজকে nটা worker-এ ভাগ করে দিলেই হয়)।

আমরা এটা ইতিমধ্যে দুইবার দেখেছি, ভিন্ন স্কেলে:

  • Lesson ১০-এর SIMD — একটা vaddps instruction ৮টা float-এ একই + অপারেশন প্রয়োগ করে
  • Lesson ১২-এর vecAdd kernel — হাজার হাজার CUDA thread, প্রতিটা একটা element-এ একই + প্রয়োগ করে

দুটোই map, শুধু ভিন্ন hardware-এ, ভিন্ন “width”-এ প্রয়োগ করা।

REDUCE — associative combine, tree দিয়ে সমান্তরাল

সংজ্ঞা: একটা associative binary operator ⊕ ব্যবহার করে একটা সিকোয়েন্সের সব element-কে একটা একক ফলাফলে একত্রিত করা।

reduce(⊕,[x1,…,xn])=x1⊕x2⊕⋯⊕xn\text{reduce}(\oplus, [x_1, \ldots, x_n]) = x_1 \oplus x_2 \oplus \cdots \oplus x_n

সাধারণ উদাহরণ: sum (+), product (×), max, min, logical AND/OR। এদের সবার একটা কমন বৈশিষ্ট্য: associativity — (a⊕b)⊕c=a⊕(b⊕c)(a \oplus b) \oplus c = a \oplus (b \oplus c)। মানে গ্রুপিং পাল্টালেও ফলাফল একই থাকে (যদিও floating-point-এ এটা প্রায়-সত্য, সম্পূর্ণ সত্য নয় — misconception section-এ ফিরে আসছি)।

সরল সিরিয়াল কোড শুধু বাম থেকে ডানে accumulate করে:

total = 0
for x in data:
    total = total + x    # প্রতিটা ধাপ আগেরটার উপর নির্ভর করে — একটা লম্বা chain

এটা n-1টা + অপারেশন, কিন্তু সিরিয়ালি, একটার পর একটা — depth n-1।

কিন্তু associativity-র কারণে আমরা পুনর্গঠন করতে পারি — element-গুলোকে জোড়ায় জোড়ায় আগে যোগ করে, তারপর সেই জোড়ার ফলাফলগুলোকে আবার জোড়ায় যোগ করে — একটা tree:

Level 0 (input):     3   1   4   1   5   9   2   6
                      \ /     \ /     \ /     \ /
Level 1 (৪টা যোগ,      4       5      14       8      ← একসাথে, সমান্তরালে
সমান্তরালে):            \       /        \       /
                          \     /          \     /
Level 2 (২টা যোগ,          9               22          ← একসাথে
সমান্তরালে):                 \             /
                               \           /
Level 3 (১টা যোগ):                  31                 ← চূড়ান্ত ফলাফল

মোট কাজ (work)  = 4 + 2 + 1 = 7 = n − 1        (সিরিয়ালের সমান!)
মোট গভীরতা (depth) = 3 = log₂(8)                (সিরিয়ালের চেয়ে অনেক কম!)
Parallel reduction tree — n=8 element-এর যোগফল, sum উদাহরণ [3,1,4,1,5,9,2,6]।

সবচেয়ে গুরুত্বপূর্ণ পর্যবেক্ষণ: tree-র মোট work (কতগুলো + অপারেশন) সিরিয়ালের সমান — n-1টাই। যা কমেছে সেটা depth — কতগুলো ধাপ সিরিয়ালি করতে হবে। সিরিয়ালে n-1 ধাপ (প্রতিটা ধাপে ১টা যোগ, একের পর এক), tree-তে মাত্র log₂ n ধাপ (প্রতিটা ধাপে একাধিক যোগ, কিন্তু সেগুলো একে অপরের থেকে স্বাধীন বলে একসাথে হতে পারে)।

এটা ঠিক mathematics/asymptotic-notation লেসনের সেই পর্যবেক্ষণ — merge sort-এর recursion tree-র গভীরতা log n, কিন্তু প্রতিটা স্তরে কাজ Θ(n)। এখানেও তাই: একই মোট কাজ, কিন্তু গাছের আকারে সাজালে গভীরতা লগারিদমিক। যথেষ্ট সমান্তরাল প্রসেসর থাকলে (একটা GPU-তে হাজার হাজার lane), সেই log n গভীরতাই আসল সময় নির্ধারণ করে, n নয়।

SCAN (Prefix Sum) — দেখতে সিরিয়াল, আসলে নয়

সংজ্ঞা: input-এর প্রতিটা prefix-এর যোগফল (বা যেকোনো associative operator-এর combine) বের করা — শুধু চূড়ান্ত মোট নয়, প্রতিটা মধ্যবর্তী running total।

input:            [1, 2, 3, 4]
inclusive scan:   [1, 3, 6, 10]     ← 1, 1+2, 1+2+3, 1+2+3+4

এটা প্রথম দেখায় অবশ্যই সিরিয়াল মনে হয় — output[i] গণনা করতে তো output[i-1] লাগবেই বলে মনে হচ্ছে:

output[0] = input[0]
for i in range(1, n):
    output[i] = output[i-1] + input[i]    # আগেরটার উপর নির্ভর করছে বলে মনে হয়

কিন্তু এই নির্ভরতাটা আসল নয় — এটা একটা বিশেষ, সরল implementation-এর নির্ভরতা, সমস্যাটার নিজস্ব নির্ভরতা নয়। একটা সুপরিচিত, চতুর algorithm আছে যা Θ(\log n) depth-এ scan করে ফেলে — Hillis-Steele scan (১৯৮৬)।

Hillis-Steele scan — ধাপে ধাপে

মূল কৌশল: offset দ্বিগুণ করতে করতে (১, ২, ৪, ৮, ...), প্রতিটা ধাপে প্রতিটা position i (যদি i ≥ offset) তার বর্তমান মান আর offset দূরের position-এর মান যোগ করে।

                         offset=1                    offset=2
Position:    0   1   2   3        0   1   2   3        0   1   2   3
Initial:     1   2   3   4
                                   ↓   ↓   ↓   ↓
             i\<1: অপরিবর্তিত ──── 1
             i≥1: a[i]+a[i-1] ───────  3 = 2+1
                                              5 = 3+2
                                                     7 = 4+3
After step1:                      1   3   5   7
                                                       ↓   ↓   ↓   ↓
             i\<2: অপরিবর্তিত ───────────────────────  1   3
             i≥2: a[i]+a[i-2] ──────────────────────────────  6 = 5+1
                                                                     10 = 7+3
After step2:                                          1   3   6  10   ✓
Hillis-Steele scan — [1,2,3,4] থেকে [1,3,6,10], মাত্র ২ ধাপে (log₂ 4 = 2)।

প্রতিটা ধাপের সব position একসাথে, সমান্তরালে আপডেট হয় (পড়া হয় আগের ধাপের মান থেকে, একই ধাপের অন্য position-এর নতুন মান থেকে নয় — তাই একটা আলাদা output buffer লাগে প্রতিটা ধাপে, বা GPU-তে __syncthreads() দিয়ে সাবধানে সাজানো, লেসন ১২-এ যেমন দেখেছি)। মোট ধাপ সংখ্যা \log_2 n — n=4-এ মাত্র ২ ধাপ, n=1024-এ মাত্র ১০ ধাপ।

ভেতরে কী ঘটছে

GPU-তে বাস্তবায়ন — শেয়ার্ড মেমরির উপর দাঁড়িয়ে

গত লেসনের blockSum kernel আসলে এই lesson-এর reduction tree-র সরাসরি CUDA বাস্তবায়ন ছিল। আরেকবার দেখি, এবার reduce-pattern-এর লেন্স দিয়ে:

__shared__ float cache[THREADS_PER_BLOCK];
cache[tid] = (i \< n) ? input[i] : 0.0f;
__syncthreads();

for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
    if (tid \< stride) cache[tid] += cache[tid + stride];
    __syncthreads();     /* প্রতিটা tree-level-এর মাঝে বাধ্যতামূলক */
}

stride শুরু হয় blockDim.x / 2 থেকে, প্রতি ধাপে অর্ধেক হয়ে যায় — এটাই tree-র প্রতিটা level, উপরের Figure-এর ঠিক বিপরীত দিক থেকে (leaves থেকে root-এর দিকে)। stride অর্ধেক করে করে নামানোকে বলে sequential addressing — এটা ইচ্ছাকৃতভাবে করা হয় শুধু correctness-এর জন্য নয়, shared memory-র bank conflict এড়াতেও (shared memory একাধিক “bank”-এ ভাগ করা থাকে, ভুল stride pattern একই bank-এ একাধিক thread-কে ঠেলে দিতে পারে, যা access সিরিয়ালাইজ করে দেয় — এই বিস্তারিত পরের কোনো optimization-কেন্দ্রিক লেসনের বিষয়)।

Hillis-Steele scan-ও একইভাবে shared memory-তে বসানো যায়:

__global__ void hillisSteeleScan(float *data, int n) {
    __shared__ float temp[THREADS_PER_BLOCK];
    int tid = threadIdx.x;
    temp[tid] = data[tid];
    __syncthreads();

    for (int offset = 1; offset \< blockDim.x; offset *= 2) {
        float val = (tid >= offset) ? temp[tid - offset] : 0.0f;
        __syncthreads();          /* সবাই পড়া শেষ করুক, তারপর কেউ লিখুক */
        if (tid >= offset) temp[tid] += val;
        __syncthreads();          /* সবাই লেখা শেষ করুক, তারপর পরের ধাপ */
    }

    data[tid] = temp[tid];
}

লক্ষ্য করুন দুইটা __syncthreads() প্রতি ধাপে — একটা read-এর পরে (যাতে write শুরুর আগে সবাই পুরনো মান পড়ে ফেলে, নতুন মান নয়), একটা write-এর পরে (যাতে পরের ধাপ শুরুর আগে সব write সম্পন্ন হয়)। এটাই lesson ১২-এ দেখা race-condition সতর্কতার হাতে-কলমে প্রয়োগ — “in-place” scan লেখার সবচেয়ে সাধারণ bug হলো এই read-write আলাদা না করা।

উদাহরণ

সম্পূর্ণ উদাহরণ — reduce max, আর ৮-element scan হাতে-কলমে

Reduce max, [3, 7, 2, 9, 4, 1, 8, 5]-এ:

Level 0:  3   7   2   9   4   1   8   5
           \ /     \ /     \ /     \ /
Level 1:    7       9       4       8      ← max(3,7)=7, max(2,9)=9, max(4,1)=4, max(8,5)=8
             \       /        \       /
Level 2:        9                8         ← max(7,9)=9, max(4,8)=8
                  \              /
Level 3:              9                    ← max(9,8)=9  ✓ (আসলেই সর্বোচ্চ)

max associative (max(max(a,b),c) = max(a,max(b,c))), তাই তাকেও ঠিক sum-এর মতো tree-তে reduce করা যায় — একই depth log₂ 8 = 3, একই work n-1 = 7।

৮-element Hillis-Steele scan, [1, 1, 1, 1, 1, 1, 1, 1] (প্রতিটা position-এর প্রত্যাশিত ফলাফল i+1, running count):

Initial:        1  1  1  1  1  1  1  1
offset=1:       1  2  2  2  2  2  2  2     (i≥1: a[i]+a[i-1])
offset=2:       1  2  3  4  4  4  4  4     (i≥2: a[i]+a[i-2])
offset=4:       1  2  3  4  5  6  7  8     (i≥4: a[i]+a[i-4])   ✓

\log_2 8 = 3 ধাপে, প্রত্যাশিত [1,2,3,4,5,6,7,8] পাওয়া গেল। প্রতিটা ধাপ যাচাই করুন — offset=2 ধাপে position 4-এর মান 4 + a[2](=2) = 4+... — দাঁড়ান, হিসাব করি: offset=1-এর পরের array [1,2,2,2,2,2,2,2]। offset=2 ধাপে position ৪: a[4] + a[4-2] = 2 + 2 = 4। হ্যাঁ মেলে। position ৭: a[7] + a[5] = 2+2=4। মেলে। offset=4 ধাপে array ছিল [1,2,3,4,4,4,4,4]। position ৪: a[4]+a[0] = 4+1=5। position ৭: a[7]+a[3]=4+4=8। সবই মিলছে চূড়ান্ত [1,2,3,4,5,6,7,8]-এর সাথে।

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

EXPERIMENT

Map/Reduce/Scan সঠিকতা ও Work-Depth যাচাই

Python 3· ২০ মিনিট
import math
import random

def sequential_reduce(data, op):
    """সিরিয়াল reduce, আর operation গণনা (কাজ পরিমাপের জন্য)"""
    ops = 0
    total = data[0]
    for x in data[1:]:
        total = op(total, x)
        ops += 1
    return total, ops


def tree_reduce(data, op):
    """Recursive tree reduce — depth গণনার জন্য আলাদা রিটার্ন"""
    ops = 0
    def helper(lst):
        nonlocal ops
        if len(lst) == 1:
            return lst[0], 0
        mid = len(lst) // 2
        left, dl = helper(lst[:mid])
        right, dr = helper(lst[mid:])
        ops += 1
        return op(left, right), max(dl, dr) + 1
    result, depth = helper(data)
    return result, ops, depth


def hillis_steele_scan(data):
    """Inclusive scan, work গণনাসহ"""
    n = len(data)
    a = list(data)
    ops = 0
    offset = 1
    steps = 0
    while offset \< n:
        new_a = list(a)
        for i in range(offset, n):
            new_a[i] = a[i] + a[i - offset]
            ops += 1
        a = new_a
        offset *= 2
        steps += 1
    return a, ops, steps


def sequential_scan(data):
    a = [data[0]]
    ops = 0
    for x in data[1:]:
        a.append(a[-1] + x)
        ops += 1
    return a, ops


# ── Reduce: work একই, depth কমে ──────────────────────────────
print("REDUCE — sum, বিভিন্ন n-এ")
print(f"{'n':>6} {'seq work':>9} {'tree work':>10} {'seq depth':>10} {'tree depth':>11}")
for n in [4, 8, 16, 64, 1024]:
    data = [random.random() for _ in range(n)]
    _, seq_ops = sequential_reduce(data, lambda a, b: a + b)
    _, tree_ops, tree_depth = tree_reduce(data, lambda a, b: a + b)
    print(f"{n:>6} {seq_ops:>9} {tree_ops:>10} {n-1:>10} {tree_depth:>11}   "
          f"(log2(n)={math.log2(n):.0f})")

# ── Scan: work বাড়ে, depth কমে — বিপরীত trade-off ────────────
print("\nSCAN — inclusive prefix sum, বিভিন্ন n-এ")
print(f"{'n':>6} {'seq work':>9} {'HS work':>8} {'seq steps':>10} {'HS steps':>9}   ঠিক আছে?")
for n in [4, 8, 16, 1024]:
    data = [1] * n
    seq_result, seq_ops = sequential_scan(data)
    hs_result, hs_ops, hs_steps = hillis_steele_scan(data)
    correct = seq_result == hs_result
    print(f"{n:>6} {seq_ops:>9} {hs_ops:>8} {n-1:>10} {hs_steps:>9}   {correct}")

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

REDUCE — sum, বিভিন্ন n-এ
     n  seq work  tree work  seq depth  tree depth
     4         3          3          3           2   (log2(n)=2)
     8         7          7          7           3   (log2(n)=3)
    16        15         15         15           4   (log2(n)=4)
    64        63         63         63           6   (log2(n)=6)
  1024      1023       1023       1023          10   (log2(n)=10)

SCAN — inclusive prefix sum, বিভিন্ন n-এ
     n  seq work  HS work  seq steps  HS steps   ঠিক আছে?
     4         3        5          3         2   True
     8         7       17          7         3   True
    16        15       49         15         4   True
  1024      1023     9217       1023        10   True

দুইটা সম্পূর্ণ ভিন্ন প্যাটার্ন লক্ষ্য করুন:

Reduce-এ tree work সবসময় seq work-এর সমান (n-1), শুধু tree depth-ই log₂ n-এ নেমে যায় — বিনামূল্যে গভীরতা কমানো।

Scan-এ HS work সবসময় seq work-এর চেয়ে বেশি, ব্যবধানটা n বড় হওয়ার সাথে সাথে বাড়তে থাকে (n=1024-এ প্রায় ৯ গুণ) — depth কমানোর জন্য বাড়তি কাজের মূল্য দিতে হচ্ছে। উভয় ক্ষেত্রেই HS_result == seq_result (ঠিক আছে? True), তাই algorithm-টা সঠিক — শুধু resource trade-off ভিন্ন।

নিজে বাড়ান: sequential_reduce-কে max operator দিয়ে চালিয়ে দেখুন একই work/depth প্যাটার্ন হয় কি না। তারপর একটা non-associative operator (যেমন lambda a, b: a - b, বিয়োগ) দিয়ে tree_reduce চালান — sequential আর tree ফলাফল কি মেলে? (মিলবে না — এটাই misconception section-এর প্রথম বিষয়।)

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

Tree reduction সিরিয়ালের সমান কাজ করে কিন্তু log(n) গভীরতায়; Hillis-Steele scan সঠিক উত্তর দেয় কিন্তু n·log(n) কাজ করে, সিরিয়ালের n-এর চেয়ে বেশি — উভয় দাবিই measurement দিয়ে যাচাইযোগ্য।

নিজে বানান

BUILD IT

বহু-কোর CPU-তে সত্যিকারের Map/Reduce লাইব্রেরি

Python (multiprocessing) · ●●●○○
  1. multiprocessing.Pool ব্যবহার করে একটা parallel_map বানান যা প্রকৃতপক্ষে একাধিক CPU core ব্যবহার করে
  2. ডেটাকে chunk-এ ভাগ করে একটা parallel_reduce বানান — প্রতিটা worker তার chunk-এর আংশিক ফলাফল বের করবে
  3. সেই আংশিক ফলাফলগুলোকে চূড়ান্তভাবে একত্রিত করুন (এই ধাপটা সস্তা, ছোট)
  4. বিভিন্ন chunk size-এ wall-clock সময় মেপে সেরা chunk size খুঁজে বের করুন
  5. scan-এর জন্য numpy.cumsum ব্যবহার করুন এবং বুঝুন কেন multiprocessing দিয়ে scan সহজে সমান্তরাল করা যায় না

Python-এর GIL (Global Interpreter Lock) সত্যিকারের CPU-bound সমান্তরালতা থ্রেড দিয়ে আটকায় — তাই আমরা multiprocessing ব্যবহার করব, যা আলাদা process (আলাদা GIL) স্পন করে।

import time
import math
from multiprocessing import Pool


def cpu_heavy(x):
    """ইচ্ছাকৃতভাবে ভারী — নাহলে process spawn overhead-ই dominate করবে"""
    return sum(math.sin(x * i) for i in range(200))


def parallel_map(fn, data, n_workers=4):
    with Pool(n_workers) as pool:
        return pool.map(fn, data)


def chunk_partial_reduce(chunk):
    total = chunk[0]
    for x in chunk[1:]:
        total += x
    return total


def parallel_reduce(data, n_workers=4):
    n = len(data)
    chunk_size = (n + n_workers - 1) // n_workers
    chunks = [data[i:i+chunk_size] for i in range(0, n, chunk_size)]

    with Pool(n_workers) as pool:
        partials = pool.map(chunk_partial_reduce, chunks)

    # চূড়ান্ত ধাপ — মাত্র n_workers টা সংখ্যা, সিরিয়ালি যোগ করাই যথেষ্ট সস্তা
    final = partials[0]
    for p in partials[1:]:
        final += p
    return final


if __name__ == "__main__":
    data = list(range(200_000))

    t0 = time.perf_counter()
    result_seq = [cpu_heavy(x) for x in data]
    t_seq = time.perf_counter() - t0

    t0 = time.perf_counter()
    result_par = parallel_map(cpu_heavy, data, n_workers=4)
    t_par = time.perf_counter() - t0

    print(f"সিরিয়াল map:   {t_seq:.2f}s")
    print(f"parallel map:  {t_par:.2f}s   (speedup {t_seq/t_par:.2f}×)")
    assert result_seq == result_par

    big_data = [float(x) for x in range(2_000_000)]
    t0 = time.perf_counter()
    seq_total = sum(big_data)
    t_seq_r = time.perf_counter() - t0

    t0 = time.perf_counter()
    par_total = parallel_reduce(big_data, n_workers=4)
    t_par_r = time.perf_counter() - t0

    print(f"\nসিরিয়াল reduce:  {t_seq_r*1000:.2f}ms")
    print(f"parallel reduce: {t_par_r*1000:.2f}ms   (speedup {t_seq_r/t_par_r:.2f}×)")

সাধারণ ফলাফল (৪-কোর মেশিনে):

সিরিয়াল map:   4.85s
parallel map:  1.34s   (speedup 3.62×)

সিরিয়াল reduce:  18.32ms
parallel reduce: 24.61ms   (speedup 0.74×)   ← parallel আসলে ধীর!

Map-এ চমৎকার speedup (~৩.৬×, ৪ core-এর কাছাকাছি) — কারণ প্রতিটা cpu_heavy(x) কল যথেষ্ট ভারী যে process spawn/communicate overhead নগণ্য হয়ে যায়। Reduce-এ উল্টো ফল — কারণ sum()-এর মতো হালকা কাজে process তৈরি ও ফলাফল pickle করে ফেরত পাঠানোর overhead-ই আসল কাজের চেয়ে বেশি সময় নেয়।

এটাই একটা গুরুত্বপূর্ণ, ঘন ঘন ভুলে যাওয়া শিক্ষা: parallel pattern চেনা যথেষ্ট নয় — granularity (প্রতিটা কাজের ইউনিট কতটা ভারী) সিদ্ধান্ত নেয় parallelize করা লাভজনক কি না। খুব হালকা কাজে overhead জিতে যায়। এটা পরের লেসনের Amdahl’s Law-র সাথেও সরাসরি সম্পর্কিত — parallel overhead effectively (1-p) অংশটাই বাড়িয়ে দেয়।

Scan multiprocessing দিয়ে কেন কঠিন: Hillis-Steele-এর প্রতিটা ধাপে সব worker-কে synchronize হতে হয় (একটা ধাপ শেষ না হওয়া পর্যন্ত পরেরটা শুরু করা যায় না) — এই বারবার sync-এর overhead process-based multiprocessing-এ shared-memory GPU thread-এর __syncthreads()-এর চেয়ে বহুগুণ বেশি ব্যয়বহুল। তাই CPU-তে বাস্তবে scan-এর জন্য numpy.cumsum() ব্যবহার করা হয় — এটা ভেতরে সিরিয়াল, কিন্তু vectorized (SIMD-accelerated) C loop, যা ছোট overhead-এ O(n) কাজ করে ফেলে। GPU-তেই Hillis-Steele/Blelloch-ধরনের সত্যিকার সমান্তরাল scan লাভজনক, কারণ সেখানে __syncthreads() একটা hardware-সমর্থিত, সস্তা barrier — process spawn নয়।

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

এই তিনটা pattern যেখানে সবকিছু চালায়

Hadoop/Spark MapReduce। ঠিক এই নামেই — Google-এর ২০০৪-এর MapReduce paper পুরো ধারণাটাকেই single machine থেকে পুরো data center-এ স্কেল করেছিল। একটা map ধাপ হাজার হাজার সার্ভারে স্বাধীনভাবে ডেটা প্রসেস করে, একটা reduce ধাপ সেই ফলাফলগুলো একত্রিত করে — এই লেসনের GPU-র map/reduce-এর সাথে গঠনগতভাবে হুবহু একই ধারণা, শুধু “worker” এখন একটা CUDA thread নয়, একটা পুরো সার্ভার। distributed-systems module (Level ৯)-এ এই স্কেলের চ্যালেঞ্জ (network failure, straggler node) বিস্তারিত দেখা হয়েছে।

NVIDIA Thrust/CUB। thrust::reduce(), thrust::inclusive_scan() — production-grade, হাতে-optimize করা (bank-conflict-free, work-efficient Blelloch scan) library function, যা এই লেসনে শেখা concept-গুলোর industrial-strength বাস্তবায়ন। বাস্তব CUDA প্রোগ্রামাররা কখনো নিজে থেকে scan/reduce হাতে লেখেন না যদি না বিশেষ কারণ থাকে — Thrust/CUB কল করেন।

Radix sort। প্রতিটা digit/bit-এর জন্য bucket-এ ভাগ করার সময় প্রতিটা element-এর output position বের করতে scan ব্যবহার হয় (exclusive prefix sum — কতগুলো element তার আগে একই bucket-এ যাচ্ছে)। GPU radix sort-এর মূল বাধা এই scan ধাপের গতি — এজন্যই work-efficient scan এত গুরুত্বপূর্ণ।

Stream compaction। একটা array থেকে শুধু কিছু predicate মানা element রাখা (বাকি বাদ), কিন্তু ক্রম বজায় রেখে — physics engine-এ “শুধু জীবিত particle রাখো” এই কাজে ব্যবহৃত হয়। Exclusive scan দিয়ে প্রতিটা মানা element-এর নতুন output index হিসাব করা হয়, তারপর একটা map ধাপে সেই index-এ লেখা হয় — scan + map-এর composition।

Histogram equalization (image processing)। ছবির contrast বাড়ানোর একটা কৌশল CDF (cumulative distribution function) ব্যবহার করে — যা আসলে pixel-intensity histogram-এর একটা scan।

Database analytics/window function। SQL-এর SUM(...) OVER (ORDER BY date) (running total) সরাসরি একটা scan। Columnar database engine (ClickHouse, DuckDB) এই ধরনের aggregate query SIMD/GPU-accelerated reduce/scan দিয়ে চালায় — লেসন ১০-এর SIMD-র সাথে সরাসরি যোগসূত্র।

Functional programming। Haskell-এর foldr/foldl (reduce), scanl (scan), Python-এর functools.reduce(), itertools.accumulate() (scan) — এই একই তিনটা pattern প্রতিটা আধুনিক functional/multi-paradigm ভাষায় first-class citizen হিসেবে আছে, কারণ এরা শুধু GPU-নির্দিষ্ট নয় — এরা সমান্তরাল চিন্তার সার্বজনীন building block।

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

“Reduce যেকোনো binary operator দিয়েই সমান্তরাল করা যায়।”

শুধু associative operator-এর জন্য tree reduction নিরাপদ। Associativity মানে গ্রুপিং পাল্টালেও ফলাফল একই থাকে — (a⊕b)⊕c=a⊕(b⊕c)(a \oplus b) \oplus c = a \oplus (b \oplus c)।

+, ×, max, min, logical AND/OR — সবই associative। কিন্তু বিয়োগ (-) associative নয়: (10 - 3) - 2 = 5, কিন্তু 10 - (3 - 2) = 9 — সম্পূর্ণ ভিন্ন উত্তর! একটা tree-reduce - দিয়ে চালালে গ্রুপিংয়ের উপর নির্ভর করে ভিন্ন ভিন্ন (ভুল) উত্তর আসবে।

Floating-point addition-ও কঠোরভাবে associative নয় — rounding error-এর কারণে। ধরুন data = [1e16, 1, -1e16]। সিরিয়াল বাম-থেকে-ডান:

(1e16 + 1) + (-1e16)

Double precision-এ 1e16-এর কাছাকাছি সংখ্যার resolution (ULP) 1-এর চেয়ে বড় (কারণ 1e16 > 2^53 ≈ 9.007×10^15, IEEE-754 double precision-এর নিখুঁত-integer সীমা) — তাই 1e16 + 1 রাউন্ড হয়ে আবার 1e16-ই থাকে, 1-টা হারিয়ে যায়। ফলে:

(1e16+1)+(−1e16)=1e16+(−1e16)=0(1e16 + 1) + (-1e16) = 1e16 + (-1e16) = 0

কিন্তু প্রকৃত গাণিতিক যোগফল 1e16 + 1 - 1e16 = 1, 0 নয়! আরেকভাবে গ্রুপ করলে (আগে 1e16 আর -1e16 যোগ করে, তারপর 1):

1e16+(−1e16)+1=0+1=1✓সঠিক1e16 + (-1e16) + 1 = 0 + 1 = 1 \quad \checkmark \text{সঠিক}

এখানেই গুরুত্বপূর্ণ ব্যবহারিক প্রভাব: একটা GPU tree-reduce sum আর একটা CPU সিরিয়াল sum-এর ফলাফল সামান্য (বা adversarial ডেটায় বড়) ভিন্ন হতে পারে, কারণ তারা ভিন্ন ক্রমে যোগ করে — গাণিতিকভাবে “সঠিক” কোনো একক উত্তর নেই, শুধু ভিন্ন rounding। এটাই deep learning training-এ GPU non-determinism-এর একটা পরিচিত বাস্তব উৎস — একই মডেল, একই ডেটা, কিন্তু ভিন্ন GPU run-এ সামান্য ভিন্ন ফলাফল, কারণ reduction-এর অভ্যন্তরীণ ক্রম প্রতিবার একই নাও হতে পারে (thread scheduling-এর উপর নির্ভরশীল)। cudnn.deterministic-এর মতো flag ঠিক এই কারণেই থাকে।

“Scan সংজ্ঞা অনুযায়ীই সিরিয়াল — সমান্তরাল করা অসম্ভব।”

এই লেসনের কেন্দ্রীয় contra-উদাহরণ: Hillis-Steele algorithm \Theta(\log n) depth-এ সম্পূর্ণ সঠিক scan করে। বিভ্রান্তিটা আসে সবচেয়ে সরল, naive sequential implementation-কে সমস্যার অন্তর্নিহিত নির্ভরতা ভেবে ফেলা থেকে — কিন্তু naive implementation আর সমস্যার প্রকৃত dependency structure এক জিনিস নয়। এটা ঠিক সেই একই ভুল যা মানুষ প্রায়ই করে asymptotic analysis-এ (mathematics/asymptotic-notation লেসন) — একটা নির্দিষ্ট algorithm-এর বৈশিষ্ট্যকে সমস্যার নিজস্ব সীমা ভেবে ফেলা।

“সমান্তরাল algorithm-এ বেশি total operation করা মানেই সেটা খারাপ ডিজাইন।”

Hillis-Steele n \log n কাজ করে, সিরিয়ালের n-এর চেয়ে বেশি। তবুও এটা একটা valid, প্রায়ই বাস্তবে ব্যবহৃত algorithm — কারণ যথেষ্ট প্রসেসর থাকলে (GPU-র হাজার হাজার lane), বাড়তি “কাজ” সমান্তরালে প্রায় বিনামূল্যে হয়ে যায়, আর যা আসলে গুরুত্বপূর্ণ সেটা হলো wall-clock সময়, যা depth-এর সমানুপাতিক, মোট work-এর নয়।

সঠিক প্রশ্ন “কম কাজ করছে কি না” নয় — “আমার হার্ডওয়্যার বাজেটে (প্রসেসর সংখ্যা বনাম n) কোন algorithm দ্রুততম শেষ করবে”। যখন প্রসেসর সীমিত (CPU-র ৮–৬৪ core, n লক্ষ লক্ষ), work-efficient Blelloch scan (Θ(n) work) সাধারণত জেতে। যখন প্রসেসর প্রচুর (GPU-র হাজার হাজার lane, n তুলনামূলক ছোট প্রতি block-এ), সরল Hillis-Steele-ও যথেষ্ট ভালো, আর কোড সরল রাখার সুবিধা পাওয়া যায়।

“Map/reduce/scan শুধু GPU-নির্দিষ্ট ধারণা।”

এই তিনটা সাধারণ সমান্তরাল চিন্তার building block, GPU আবিষ্কার করেনি। Hadoop/Spark পুরো data center জুড়ে এই একই pattern চালায় (realworld section দ্রষ্টব্য)। Multi-core CPU-তে multiprocessing.Pool-ও একই pattern প্রয়োগ করে, যদিও ভিন্ন granularity ও overhead বৈশিষ্ট্যে (build exercise দ্রষ্টব্য)। এমনকি functional programming ভাষায় (Haskell, এমনকি সিরিয়াল Python) map/reduce/scan first-class function হিসেবে আছে, কারণ তাদের describe করার ক্ষমতা (কোন কাজ independent, কোনটা associative-combine, কোনটা running-total) hardware-নিরপেক্ষভাবেই মূল্যবান — সমান্তরাল করা হোক বা না হোক, এই pattern-এ কোড লেখা bug কমায় আর পড়া সহজ করে।

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

1

n = 1024 element-এর একটা array reduce (sum) করতে হবে। Tree reduction-এ depth কত, আর মোট কতগুলো + অপারেশন লাগবে? সিরিয়ালের সাথে তুলনা করুন।

প্রয়োগ

Depth: \log_2 1024 = 10। মাত্র ১০টা সিরিয়াল ধাপ লাগবে, n=1024-এর জন্যও।

মোট work: n - 1 = 1023টা + অপারেশন — ঠিক সিরিয়ালের সমান। Reduce-এ tree পুনর্গঠন কোনো বাড়তি কাজ যোগ করে না, শুধু সেই কাজকে সমান্তরাল-বান্ধব আকারে সাজায়।

তুলনা: সিরিয়াল — 1023 ধাপ, প্রতিটায় ১টা +, মোট সময় \propto 1023। Tree — 10 ধাপ, প্রথম ধাপে 512টা + সমান্তরালে, পরের ধাপে 256টা, … শেষ ধাপে ১টা। যথেষ্ট প্রসেসর (≥512 টা) থাকলে মোট সময় \propto 10 — প্রায় ১০০× দ্রুত, বিনা বাড়তি কাজে।

2

“Reduce-এ tree ব্যবহার করে depth কমানো বিনামূল্যে হয় (work একই থাকে), কিন্তু scan-এ Hillis-Steele দিয়ে depth কমাতে গেলে বাড়তি work দিতে হয়” — কেন এই পার্থক্যটা ঘটে? দুটোই তো associative operator ব্যবহার করছে।

যুক্তি

পার্থক্যটা কী গণনা করতে হচ্ছে তার গঠনে, শুধু operator-এর associativity-তে নয়।

Reduce-এ আমাদের দরকার শুধু একটা চূড়ান্ত ফলাফল। Tree-র প্রতিটা internal node একটা intermediate ফলাফল ধরে রাখে যা পরে আর কখনো আলাদাভাবে দরকার হয় না — শুধু তার parent-এর হিসাবে অবদান রাখে। তাই n element-কে n-1টা combine-এ (যেকোনো বাইনারি tree-শেপে) একটা ফলাফলে নামানো যায় — associativity গ্যারান্টি দেয় যেকোনো grouping একই উত্তর দেবে, তাই আমরা সবচেয়ে balanced, কম-depth grouping (একটা perfectly balanced tree) বেছে নিতে পারি, বিনা মূল্যে।

Scan-এ আমাদের দরকার প্রতিটা position-এর জন্য আলাদা ফলাফল (nটা ভিন্ন আউটপুট, শুধু একটা নয়)। position i-এর ফলাফল নির্ভর করে 0 থেকে i পর্যন্ত সব element-এর উপর। Hillis-Steele যা করে তা হলো প্রতিটা position-এর জন্য আলাদাভাবে এই আংশিক combination-গুলো তৈরি করা, কিন্তু বিভিন্ন position-এর জন্য এই আংশিক কাজগুলোর মধ্যে overlap থাকে (যেমন output[5] আর output[6] দুটোই input[0..4]-এর কিছু অংশ পুনরায় ব্যবহার করে, কিন্তু আলাদাভাবে recompute করে) — এই পুনরাবৃত্তিই বাড়তি work-এর উৎস।

Blelloch scan এই সমস্যা সমাধান করে দুই-ধাপে (প্রথমে reduce-এর মতো up-sweep করে partial সমষ্টি জমা করে, তারপর সেগুলো down-sweep-এ পুনর্বণ্টন করে) — এভাবে reduce-এর “বিনামূল্যে depth-কমানো” সুবিধাটা scan-এও Θ(n) work বজায় রেখে পাওয়া সম্ভব, কিন্তু এলগরিদমটা উল্লেখযোগ্যভাবে জটিল, তাই এই লেসনে শুধু সরল Hillis-Steele বিস্তারিত দেখানো হয়েছে।

3

আপনি একটা tree_reduce function লিখেছেন যা যেকোনো op নেয় (associative ধরে নিয়ে)। কেউ সেটা op = lambda a, b: (a - b) (বিয়োগ) দিয়ে কল করে। কী ভুল হবে, এবং কেন এই bug-টা মাঝেমধ্যে “কাকতালীয়ভাবে সঠিক” ফলাফলও দিতে পারে?

যুক্তি

বিয়োগ associative নয়, তাই tree-র shape-এর উপর নির্ভর করে ভিন্ন ভিন্ন উত্তর আসবে — আর সেই উত্তর সিরিয়াল বাম-থেকে-ডান বিয়োগের উত্তর থেকেও ভিন্ন হবে।

উদাহরণ: [10, 3, 2]।

সিরিয়াল (বাম-থেকে-ডান): (10 - 3) - 2 = 5।

Balanced tree (যদি tree_reduce প্রথমে [10,3] জোড়া বানায়, তারপর 2-এর সাথে combine করে): (10 - 3) - 2 = 5, একই — কারণ n=3-এ tree structure বাম-থেকে-ডানের কাছাকাছি পড়ে যেতে পারে।

কিন্তু [10, 3, 2, 1]-এ (n=4, perfectly balanced tree, mid = len//2 = 2): tree প্রথমে [10,3] আর [2,1]-কে আলাদাভাবে combine করবে: left = 10-3 = 7, right = 2-1 = 1, তারপর 7 - 1 = 6।

সিরিয়াল: ((10-3)-2)-1 = (7-2)-1 = 5-1 = 4।

6 ≠ 4 — ভিন্ন উত্তর, শুধু গ্রুপিং পাল্টানোর কারণে।

কেন মাঝেমধ্যে “কাকতালীয়ভাবে সঠিক” মনে হয়: ছোট n (২ বা ৩)-এ, অথবা এমন ডেটায় যেখানে element-গুলো ০ (identity-র কাছাকাছি), গ্রুপিং-জনিত পার্থক্য চোখে না-ও পড়তে পারে, বা টেস্ট কেস দুর্ঘটনাক্রমে মিলে যেতে পারে — যা এই bug-কে বিশেষভাবে বিপজ্জনক করে তোলে: ছোট পরীক্ষায় pass করে, কিন্তু বড় বা ভিন্ন ডেটায় ব্যর্থ হয়। নিয়ম: reduce ব্যবহারের আগে সবসময় যাচাই করুন operator সত্যিই associative কি না — গাণিতিকভাবে প্রমাণ করে, “মনে হচ্ছে ঠিক আছে” দিয়ে নয়।

4

Build exercise-এ দেখা গেছে হালকা কাজে (sum()) parallel_reduce আসলে সিরিয়ালের চেয়ে ধীর। কোন পরিবর্তন এই নির্দিষ্ট experiment-কে পাল্টে দিতে পারে যাতে parallel version জেতে?

প্রয়োগ

সমস্যাটা granularity — প্রতিটা worker-কে দেওয়া কাজ এত হালকা যে process spawn/communication overhead আসল কাজের চেয়ে বেশি সময় নেয়। তিনটা দিক থেকে এটা ঠিক করা যায়:

১. মোট ডেটা অনেক বড় করা। n = 2,000,000-এর বদলে n = 200,000,000 হলে প্রতিটা chunk-এর কাজ (এমনকি সরল sum-ও) যথেষ্ট ভারী হয়ে যাবে যে ৪-worker overhead আপেক্ষিকভাবে নগণ্য হয়ে পড়বে।

২. প্রতি element-এর কাজ ভারী করা (build exercise-এর cpu_heavy function-এর মতো) — map-এ ঠিক এটাই ঘটেছিল, তাই সেখানে ভালো speedup পাওয়া গেছে। যদি reduce-এর chunk_partial_reduce-এও প্রতিটা যোগের আগে একটা ভারী হিসাব থাকত, একই সুবিধা পাওয়া যেত।

৩. Worker সংখ্যা কমিয়ে, chunk বড় করা — কম কিন্তু বড় chunk মানে কম process-communication overhead (মোট n_workersটা spawn/pickle, ডেটার আকারের সাথে সরাসরি সম্পর্কিত নয়)।

সাধারণ নীতি: parallelize করার লাভ তখনই আসে যখন per-unit কাজ overhead-এর চেয়ে উল্লেখযোগ্যভাবে বেশি ভারী — এই “overhead-vs-কাজ” অনুপাতটাই আসলে পরের লেসনের Amdahl’s Law-র (1-p) অংশে (serial/overhead ভাগ) রূপান্তরিত হয়।

5

একটা লগ-প্রসেসিং সিস্টেমে, প্রতিদিনের কোটি কোটি লগ লাইন থেকে (ক) প্রতিটা লাইনকে parse করে একটা structured record বানাতে হবে, (খ) মোট error-count বের করতে হবে, (গ) প্রতি মিনিটে কতগুলো request এসেছে তার একটা running/cumulative count লাগবে dashboard-এর জন্য। এই তিনটা কাজের জন্য কোন pattern (map/reduce/scan) মানানসই, এবং কেন?

ডিজাইন

তিনটা আলাদা কাজ, তিনটা আলাদা pattern — একটা বাস্তব সিস্টেম কীভাবে এই তিনটাকে কম্পোজ করে তার ভালো উদাহরণ।

(ক) Parse করা → MAP। প্রতিটা লগ লাইন সম্পূর্ণ স্বাধীনভাবে parse করা যায় — একটা লাইনের parsing আরেকটার উপর নির্ভর করে না। log_line → structured_record একটা বিশুদ্ধ map, হাজার হাজার worker/thread/GPU-lane-এ trivially বিতরণযোগ্য।

(খ) মোট error-count → REDUCE। প্রতিটা রেকর্ড থেকে “এটা error কি না” (০ বা ১) বের করে, তারপর সবগুলো যোগ (+, associative) করে একটা সংখ্যায় নামাতে হবে। এটা classic reduce — প্রথমে (ক)-এর মতো map করে বুলিয়ান/সংখ্যায় রূপান্তর, তারপর reduce দিয়ে sum।

(গ) প্রতি-মিনিট cumulative count → SCAN। এখানে প্রতিটা মিনিটের জন্য পৃথক, ক্রমবর্ধমান ফলাফল দরকার (dashboard-এ একটা লাইন গ্রাফ আঁকতে) — শুধু একটা চূড়ান্ত সংখ্যা নয়। প্রতি-মিনিট count-এর array-তে exclusive/inclusive scan চালালে সরাসরি cumulative total-এর সিকোয়েন্স পাওয়া যায় — ঠিক এই লেসনের [1,2,3,4] → [1,3,6,10] উদাহরণের মতো, শুধু “element” এখন “এক মিনিটের request সংখ্যা”।

সম্পূর্ণ pipeline: MAP (parse) → MAP (is_error বা per-minute bucket-এ ফেলা) → REDUCE (মোট error) সমান্তরালে, আবার আলাদাভাবে সেই একই per-minute বাকেট থেকে SCAN (cumulative)। এই তিনটা ধাপ একটা Spark/Flink pipeline-এ সরাসরি .map(), .reduce(), .scan()/windowed-aggregation API call হিসেবে লেখা হয় — distributed-systems module-এ এই স্কেলের বাস্তবায়ন বিস্তারিত।

এরপর কী

এরপর কী

আমরা এখন তিনটা শক্তিশালী building block হাতে পেয়েছি — map, reduce, scan — যা দিয়ে বিশাল সংখ্যক বাস্তব সমস্যা সমান্তরাল করা যায়, single GPU থেকে পুরো data center পর্যন্ত। কিন্তু একটা কেন্দ্রীয় প্রশ্ন এখনও বাকি: কতটা দ্রুত হতে পারে, বাস্তবে?

প্রতিটা বাস্তব প্রোগ্রামের কিছু অংশ এই pattern-গুলোয় সুন্দরভাবে ভাঙা যায় (parallel), কিছু অংশ পারে না — file I/O, network call, setup/teardown, বা একটা lock-সুরক্ষিত shared state আপডেট (আগের লেসনগুলোর false sharing/NUMA/atomics মনে করুন)। এই অপরিবর্তনীয় সিরিয়াল অংশটাই নির্ধারণ করে দেয় আপনার প্রোগ্রাম আসলে কতটা দ্রুত হতে পারবে, যত প্রসেসরই যোগ করুন না কেন।

পরের লেসনে আমরা এই সীমাটাকে গণিতের ভাষায় সুনির্দিষ্ট করব — Amdahl’s Law — আর দেখব কেন এটা একটা genuinely সোজাসাপ্টা, সোবারিং ফলাফল। তারপর Gustafson’s Law দিয়ে একটা গুরুত্বপূর্ণ পাল্টা-প্রশ্ন করব: যদি সমস্যার আকারই বড় হতে পারে, তাহলে গল্পটা কি বদলে যায়?

আরও পড়ুন

  • Programming Massively Parallel Processors, 4th ed. — Chapters 10–11 — Kirk & Hwu · Reduction ও scan-এর CUDA implementation বিস্তারিতভাবে
  • Prefix Sums and Their Applications — Guy E. Blelloch (1990) · Work-efficient scan algorithm-এর মূল paper
  • Thrust/CUB Documentation — NVIDIA · বাস্তব production-grade reduce/scan primitive