Back to IF3130 Sistem Paralel dan Terdistribusi
Pola Algoritma Paralel: Konvolusi (Stencil Computation)
Questions/Cues
Apa itu komputasi konvolusi?
Apa itu stencil operation?
Bagaimana cara kerja konvolusi 1D?
Apa masalah kondisi batas (boundary condition)?
Apa itu ghost cells?
Bagaimana implementasi konvolusi 2D?
Mengapa tiling sangat penting untuk konvolusi?
Apa beda Input Tile vs Output Tile?
Bagaimana analisis performa Tiled Convolution?
Reference Points
7 - IF-3230-07-GPU-06-2023.pdfApa itu Komputasi Konvolusi?
Konvolusi adalah sebuah operasi matematis di mana setiap elemen output dihitung sebagai jumlah terbobot (weighted sum) dari elemen-elemen tetangganya di input. Pola komputasi ini juga dikenal sebagai operasi stensil (stencil operation), karena seolah-olah ada sebuah “stensil” atau “cetakan” yang digeser di atas data input untuk menghasilkan setiap elemen output.
Komponen:
Input Array: Data sumber (misalnya, sinyal audio 1D atau gambar 2D).
Mask/Kernel Konvolusi: Sebuah array kecil yang berisi bobot-bobot. Pola bobot ini menentukan efek dari konvolusi (misalnya, blurring, sharpening, deteksi tepi).
Output Array: Hasil dari operasi.
Konvolusi 1D dan Kondisi Batas
Dalam konvolusi 1D, mask 1D digeser di sepanjang array input. Untuk menghitung
P[i], kita menempatkan pusat mask di atasN[i]dan melakukan perkalian-akumulasi dengan elemen-elemen tetangganya.
Contoh:
P[2] = N[0]*M[0] + N[1]*M[1] + N[2]*M[2] + N[3]*M[3] + N[4]*M[4]Masalah Kondisi Batas: Apa yang terjadi ketika kita mencoba menghitung
P[0]? Mask akan membutuhkan elemen-elemen diN[-2]danN[-1]yang tidak ada. Elemen “imajiner” di luar batas ini disebut ghost cells (atau halo cells).Solusi Umum:
Zero Padding: Mengasumsikan semua ghost cells bernilai nol.
Clamping/Replication: Mengulang nilai elemen di tepi (misalnya,
N[-1]dianggap sama denganN[0]).Wrapping: Mengasumsikan array bersifat periodik (misalnya,
N[-1]dianggap sama denganN[end]).Implementasi kernel naif harus secara eksplisit memeriksa apakah indeks tetangga berada dalam rentang yang valid.
Konvolusi 2D
Konsepnya sama, tetapi diperluas ke dua dimensi. Sebuah mask 2D digeser di atas gambar input 2D. Untuk menghitung nilai piksel output di
(row, col), pusat mask ditempatkan di atas piksel input di(row, col), dan operasi weighted sum dilakukan pada area yang ditutupi oleh mask. Ini adalah operasi fundamental dalam computer vision dan deep learning (misalnya, Convolutional Neural Networks).Tiled Convolution: Kunci Performa
Kernel konvolusi naif sangat tidak efisien karena tingkat penggunaan kembali data (data reuse) yang masif. Perhatikan bahwa untuk menghitung
P[i]danP[i+1], keduanya menggunakan banyak elemen input yang sama. Kernel naif akan membaca elemen-elemen ini dari Global Memory berulang kali.Tiled Convolution memecahkan masalah ini dengan:
Memuat satu bagian besar dari input (sebuah input tile) ke Shared Memory.
Melakukan komputasi untuk menghasilkan satu bagian dari output (sebuah output tile) sepenuhnya dari Shared Memory.
Input Tile vs. Output Tile
Ini adalah konsep kunci dalam tiled convolution. Karena setiap elemen output membutuhkan tetangganya, input tile yang dibutuhkan untuk menghitung sebuah output tile harus lebih besar.
Output Tile: Bagian dari matriks output yang dihitung oleh satu thread block. Ukurannya, misalnya,
O_TILE_WIDTH x O_TILE_WIDTH.Input Tile: Bagian dari matriks input yang perlu dimuat ke Shared Memory. Ukurannya adalah
(O_TILE_WIDTH + MASK_WIDTH - 1) x (O_TILE_WIDTH + MASK_WIDTH - 1).Radius tambahan di sekitar input tile ini berisi ghost cells yang diperlukan untuk menghitung elemen-elemen di tepi output tile.
Alur Kernel Tiled:
Setiap thread dalam block memuat satu elemen dari input tile yang besar dari Global Memory ke Shared Memory.
__syncthreads()untuk memastikan seluruh input tile telah dimuat.Setiap thread menghitung nilai elemen outputnya dengan mengakses data hanya dari Shared Memory.
Setiap thread menulis hasil akhirnya ke Global Memory.
Analisis Pengurangan Bandwidth Memori
Keuntungan dari tiling sangat signifikan. Rasio pengurangan bandwidth Global Memory untuk konvolusi 2D dapat dihitung sebagai:
Pembilang: Jumlah total akses ke Shared Memory.
Penyebut: Jumlah total pemuatan dari Global Memory.
Untuk
O_TILE_WIDTH = 16danMASK_WIDTH = 5, rasio pengurangannya adalah 16x. Ini berarti lalu lintas ke Global Memory dikurangi sebanyak 16 kali, yang secara dramatis meningkatkan performa.
Konvolusi adalah operasi stensil di mana setiap elemen output adalah jumlah terbobot dari tetangganya di input, yang sangat penting dalam pemrosesan sinyal dan gambar. Karena adanya penggunaan kembali data yang sangat besar, implementasi naif yang mengakses Global Memory berulang kali sangat tidak efisien. Solusi berkinerja tinggi adalah dengan menggunakan Tiled Convolution, di mana satu thread block memuat input tile yang diperbesar (untuk mencakup ghost cells) ke Shared Memory, lalu menghitung output tile sepenuhnya dari memori on-chip yang cepat, sehingga secara drastis mengurangi bottleneck bandwidth memori.
Additional Information
Optimasi: Separable Convolution
Untuk mask tertentu (seperti Gaussian blur), konvolusi 2D dapat dipecah menjadi dua lintasan konvolusi 1D: satu secara horizontal, lalu satu lagi secara vertikal pada hasilnya.
Kompleksitas Naif 2D: Untuk mask
MxM, diperlukanM²perkalian-akumulasi per piksel.Kompleksitas Separable: Memerlukan M operasi untuk lintasan horizontal dan M operasi untuk lintasan vertikal, total 2*M operasi per piksel.
Untuk mask 5x5, ini mengubah 25 operasi menjadi hanya 10 operasi. Ini adalah penghematan komputasi yang signifikan di samping penghematan memori dari tiling.
Optimasi: Constant Memory untuk Mask
Mask konvolusi adalah kandidat yang sempurna untuk disimpan di Constant Memory. Alasannya:
Read-Only: Nilainya tidak pernah berubah selama eksekusi kernel.
Akses Seragam: Semua thread dalam satu warp akan mengakses elemen mask yang sama pada saat yang bersamaan saat mereka menghitung stensil.
Penggunaan Constant Memory memungkinkan hardware untuk men-cache nilai mask dan melakukan broadcast ke semua thread dalam warp dalam satu siklus, yang jauh lebih efisien daripada setiap thread membacanya dari Global atau bahkan Shared Memory.
Eksplorasi Mandiri
Coba analisis rasio pengurangan bandwidth untuk konvolusi 1D. Bagaimana formulanya berubah dari versi 2D?
Pikirkan tentang desain thread block untuk tiled convolution. Haruskah ukuran block sama dengan output tile atau input tile? Masing-masing memiliki implikasi. Jika ukuran block sama dengan input tile, beberapa thread hanya akan memuat data dan tidak menghitung output. Jika ukuran block sama dengan output tile, beberapa thread mungkin perlu memuat lebih dari satu elemen input.