II-1
BAB II DASAR TEORI Bab ini memuat beberapa dasar teori yang mendukung dalam pengerjaan Tugas Akhir, serta penjelasan mengenai berbagai metode yang digunakan untuk menyelesaikan permasalahan dalam Tugas Akhir.
2.1 Pergerakan Telepon Seluler Sebagai Sumber Data Umumnya, pengenalan kemacetan otomatis yang terjadi di jalan raya dapat dilakukan menggunakan data kecepatan pergerakan kendaraan. Salah satu metode untuk mendapatkannya adalah dengan menggunakan sensor yang dipasang di segenap penjuru kota. Dengan asumsi bahwa dalam setiap kendaraan terdapat minimal sebuah telepon seluler, maka dalam Tugas Akhir ini sendiri diusulkan penggunaan pergerakan telepon seluler sebagai sumber data. Alasan pemilihan pergerakan telepon seluler sebagai sumber data dapat dilihat pada Tabel II-1. Tabel II-1: Perbedaan pendeteksian kecepatan kendaraan dengan sensor dan telepon seluler Sensor Kecepatan
Telepon Seluler
Kecepatan kendaraan adalah kecepatan sesaat.
Kecepatan kendaraan adalah kecepatan rata-rata antara t1 dan t0.
Tidak dapat mengidentifikasi kecepatan kendaraan individual. Untuk dapat mengidentifikasi kecepatan individual kendaraan, dibutuhkan sensor tambahan.
Dapat mengidentifikasi kecepatan kendaraan individual dengan cara melihat ID dari telepon seluler.
Membutuhkan pembangunan infrastruktur baru di segenap penjuru kota.
Dapat menggunakan infrastruktur telepon seluler yang dimiliki operator telepon seluler.
Akurasi posisi dan kecepatan tinggi.
Akurasi posisi tergantung metode yang digunakan, untuk signal-based dapat mencapai 10 m, untuk LBS dapat mencapai 1 m. Akurasi kecepatan tinggi.
Pergerakan dan posisi setiap telepon seluler direpresentasikan sebagai sebuah vektor. Sebuah data telepon seluler memiliki enam properti, yakni: 1. Nomor telepon seluler (ID), dilambangkan dengan id. 2. Timestamp, dilambangkan dengan t(i).
3. Latitude pada timestamp ini, dilambangkan dengan xt(i). 4. Longitude pada timestamp ini, dilambangkan dengan yt(i). 5. Arah kecepatan terhadap Utara, dilambangkan dengan ΞΈt(i). Data ini diambil dengan rumus: ππ‘(π) = tanβ1
π¦π‘ π β π¦π‘(πβ1) π₯π‘(π) β π₯π‘(πβ1)
6. Besar kecepatan, dilambangkan dengan vt(i). Data ini diambil dengan rumus kecepatan rata-rata: (π₯π‘(π) β π₯π‘(πβ1) )2 + (π¦π‘(π) β π¦π‘(πβ1) )2 π£π‘
π
=
π‘ π β π‘(π β 1)
2.2 Teori Graf Dalam Tugas Akhir ini, Teori Graf digunakan untuk merepresentasikan jalanjalan yang terdapat dalam sebuah kota. Graf adalah sebuah struktur data yang terdiri
atas
simpul-simpul
(vertex)
dan
sisi-sisi
(edge).
Sebuah
sisi
menghubungkan dua buah simpul [MUN06]. 2.2.1 Jenis Graf Graf dapat dikelompokkan menjadi beberapa kategori (jenis) bergantung pada sudut pandang pengelompokannya. Pengelompokan graf dapat dipandang berdasarkan ada tidaknya sisi ganda atau gelang, berdasarkan jumlah simpul, atau berdasarkan orientasi arah pada sisi. Penjelasannya pada Tabel II-2. Tabel II-2: Jenis graf Sudut Pandang Pengelompokan Ada atau tidaknya sisi ganda
Jumlah simpul
Orientasi arah pada sisi
Jenis Graf
Keterangan
Graf sederhana (simple graph)
Tidak mengandung sisi ganda
Graf tak-sederhana (unsimple graph)
Mengandung sisi ganda
Graf berhingga
Jumlah simpul berhingga
Graf tak-berhingga
Jumlah simpul tidak berhingga
Graf tak-berarah
Sisi tidak memiliki orientasi arah
Graf berarah
Sisi memiliki orientasi arah
II-2
Graf yang digunakan dalam Tugas Akhir ini adalah graf tak-sederhana, berhingga, dan berarah. 2.2.2 Implementasi Graf Dalam memori komputer, graf dapat diimplementasikan dengan struktur data sebagai berikut: 1. adjacency list 2. adjacency matrix 3. edge list Dalam Tugas Akhir ini, representasi yang dipilih adalah adjacency list. Bahasa yang digunakan untuk merepresentasikan graf dalam sebuah file adalah GraphViz DOT. Sedangkan library graf yang dipakai untuk membaca bahasa GraphViz dan membangun reprsentasi graf dalam memori komputer adalah Boost Graph Library (BGL).
2.3 Stream Programming Stream programming adalah sebuah teknik pemrograman konkuren yang dapat digunakan untuk melakukan komputasi paralel. Dalam stream programming, data diorganisasi menjadi sekumpulan stream, dan sebuah program yang disebut kernel diaplikasikan pada setiap elemen stream [BUC06]. Dari Gambar II-1, dapat dilihat bahwa stream programming menggunakan paradigma Single Instruction, Multiple Data (SIMD). Artinya, program yang sama dieksekusi pada sekumpulan data secara paralel. Salah satu perangkat keras yang dapat digunakan untuk melakukan stream programming adalah Graphic Processing Unit (GPU). GPU sangat cocok untuk diprogram dengan teknik ini, karena GPU dirancang dari awal untuk melakukan proses rendering objek 3D, dimana proses tersebut membutuhkan hardware yang dirancang secara khusus untuk memproses data secara SIMD. Untuk mengimplementasikan stream programming pada GPU, ada banyak framework
II-3
yang dapat dipakai. Beberapa framework yang dapat dipakai dapat dilihat pada Tabel II-3: Tabel II-3: Framework untuk stream programming pada GPU Nama / Istilah Framework
Interface
Direct Programming
Direct3D 9
BrookGPU
Direct3D9
Tipe
Versi Terbaru
Isi
API
C++ Language
OpenGL 1.5
-
Cg / HLSL Bahasa
Brook Language
OpenGL 1.5
0.3 Beta
Brook Compiler Brook Runtime
Compute Unified Device Architecture (CUDA)
nVidia G80 Series Driver
Ekstensi Bahasa C
Nvcc Compiler
1.0
CUDA API CUDA Toolkit CUDA Basic Linear Algebra System CUDA Fast Fourier Transform
Parallel Programming
Single Instruction Multiple Data (SIMD)
Stream programming
On GPU (GPGPU)
Direct (via OpenGL)
Multiple Instruction Multiple Data (MIMD)
Concurrent programming
On Cell Broadband Engine
ATI CTM
nVidia CUDA
Gambar II-1: Stream programming dalam taksonomi parallel programming
II-4
Pada Tugas Akhir ini, stream programming dilakukan di atas sebuah GPU nVidia GeForce 8800GTS 320MB, yang merupakan GPU jenis G80, sedangkan framework yang dipilih adalah Compute Unified Device Architecture (CUDA). CUDA dipilih karena memiliki robustness yang tinggi (menggunakan ekstensi bahasa C) dan tingkat kesulitannya di antara framework yang lain tidak begitu rumit untuk diimplementasikan dalam Tugas Akhir ini. 2.3.1 Kemampuan Komputasi GPU Graphic Processing Unit (GPU) adalah salah satu peripheral komputer yang bertugas untuk melakukan proses rendering objek 3D ke layar. Karena rendering merupakan proses yang sangat paralel, maka evolusi arsitektur GPU pun mengikutinya: GPU berevolusi menjadi sebuah multiprosesor yang mampu menjalankan sebuah tugas secara paralel. GPU dapat disebut unit komputasi yang terpisah dari CPU karena GPU memiliki prosesor dan memori sendiri. GPU memiliki jumlah Arithmetic Logical Unit (ALU) yang jauh lebih besar dibandingkan komputer biasa. Oleh karena itu, kemampuan GPU dalam mengolah data yang besar dapat diandalkan. Pada Gambar II-2, dapat dilihat bahwa pertumbuhan kemampuan GPU jika diukur dengan Floating Point Operations per Second (FLOPS) jauh meninggalkan CPU, bahkan meninggalkan hukum Moore.
Gambar II-2: Peningkatan kecepatan GPU
II-5
2.4 Compute Unified Device Architecture (CUDA) CUDA adalah sebuah framework yang dapat digunakan untuk melakukan stream programming pada GPU nVidia seri G8x. Saat ini, CUDA sudah memasuki versi 1.0. GPU GeForce 8800GTS yang digunakan dalam Tugas Akhir ini merupakan anggota seri nVidia G80. Prosesor yang ada dalam GeForce 8800GTS berjumlah 96 buah [NVI05]. Prosesor tersebut dinamakan stream processor. Stream processor tersebut dikelompokkan menjadi 12 buah multiprocessor yang masing-masing berisi 8 stream processor. Setiap stream processor akan menjalankan sebuah thread dalam satu waktu. Gambar II-3 menerangkan tentang hirarki memori dalam CUDA. Setiap stream processor, yang dianalogikan dengan thread(x,y) dikelompokkan dalam block-block. Setiap block memiliki memori bersama yang dinamakan shared memory. Shared memory adalah memori berukuran kecil berkecepatan tinggi yang hanya dapat diakses stream processor dalam block yang sama. Bandwidth antara stream processor dan shared memory sangat besar, mencapai 70,7 GB/s. Ada lagi jenis memory lain yang bernama device memory. Memori ini dapat diakses oleh seluruh multiprocessor, akan tetapi dengan bandwidth yang lebih lambat. Adapun memori komputer diistilahkan dengan host memory. Implementasi konsep stream programming pada CUDA adalah sebagai berikut: a. Stream: Dalam CUDA, stream disimpan dalam sebuah memori milik GPU yang bernama global memory. Stream diorganisasi berdasarkan jumlah thread yang akan mengaksesnya secara paralel. Setelah itu, thread tersebut harus diorganisasikan menjadi thread block. b. Kernel: Dalam CUDA, kernel tetap merupakan sebuah program yang beroperasi pada data yang terdapat pada global memory.
II-6
Gambar II-3: Hirarki memori pada CUDA
Langkah-langkah umum untuk melakukan komputasi pada CUDA adalah sebagai berikut: 1. Tentukan jumlah thread yang akan beroperasi pada stream. Kemudian tentukan jumlah thread per block dan jumlah thread block yang efisien pada
hardware
CUDA.
Jumlah
tersebut
dinamakan
execution
configuration. Contoh execution configuration dapat dilihat pada Gambar II-4. Pada gambar itu, sebuah kernel dieksekusi dengan blok berukuran 2x3, dimana setiap blok terdiri atas 5x3 thread. 2. Pindahkan semua data dari host memory ke device memory. Ini harus dilakukan karena GPU tidak bisa membaca host memory. 3. Lakukan komputasi pada device memory beserta execution configurationnya. Data dalam device memory ini yang akan menjadi stream. 4. Pindahkan hasil komputasi pada device memory ke host memory.
II-7
Gambar II-4: Thread dan thread block
2.4.1 Contoh Komputasi pada CUDA: Parallel Scan Agar konsep CUDA dapat dimengerti, maka lebih baik disertakan contoh kasus yang menggunakan CUDA sebagai solusinya. Dalam sub-bab ini akan dijelaskan algoritma parallel prefix sum sebagai contoh penggunaan CUDA. Parallel prefix sum, juga dikenal sebagai parallel scan, adalah sebuah building block untuk banyak algoritma paralel seperti sorting dan pembuatan struktur data. Di sini akan dicontohkan implementasi parallel scan secara efisien pada CUDA beserta perbandingannya dengan metode sekuensial biasa. Contoh ini juga dimaksudkan untuk menjelaskan konsep pada CUDA yakni: 1. Thread 2. Thread block
II-8
3. Warp 4. Shared memory 5. Memory banks 6. Memory bank conflict 2.4.1.1 Definisi All-Prefix-Sum Operasi all-prefix-sum didefinisikan sebagai berikut: Sebuah operasi all-prefixsum membutuhkan sebuah operator biner asosiatif β dan sebuah array dari elemen [a0, a1, β¦, an] dan mengembalikan array [a0, (a0 β a1), β¦, (a0 β a1 β β¦ β an)] Contoh: Jika β adalah penambahan, maka sebuah operasi all-prefix-sum pada array [3 1 7 0 4 1 6 3] akan menghasilkan [3 4 11 11 15 16 22 25] Ada banyak kegunaan all-prefix-sum, meliputi βtetapi tidak terbatas padapengurutan (sorting), analisis leksikal, perbandingan string, evaluasi polinom, pemampatan stream, membuat histogram, dan struktur data (graf, pohon, dll) secara paralel [HAR07]. 2.4.1.2 Inclusive Scan dan Exclusive Scan All-prefix-sum pada sebuah array sering disebut scan. Terminologi ini akan dipakai sampai akhir bab 2. Seperti yang dilihat dari contoh di atas, sebuah scan pada array akan menghasilkan array baru dimana setiap elemen j adalah jumlah dari elemen sampai dan termasuk j. Ini disebut inclusive scan. Ada pula jenis scan lain dimana setiap elemen j pada array hasil berisi jumlah dari seluruh elemen sebelumnya, namun tidak termasuk j itu sendiri. Operasi ini biasa disebut exclusive scan (atau prescan). II-9
Definisi: Sebuah operasi exclusive scan membutuhkan operator asosiatif biner β dengan identitas I, dan sebuah array n-elemen [a0, a1, β¦, an] dan mengembalikan array [I, a0, (a0 β a1), β¦, (a0 β a1 β β¦ β an-2)] Contoh: jika β adalah sebuah penjumlahan, maka exclusive scan pada array [3 1 7 0 4 1 6 3] akan menghasilkan [0 3 4 11 11 15 16 22] Sebuah array exclusive scan bisa dibangkitkan dengan menggeser array inclusive scan ke kanan sebanyak 1 elemen kemudian menyisipkan identitas. Dengan hal yang sama, sebuah array inclusive scan bisa dibangkitkan dengan menggeser array exclusive scan ke kiri sebanyak 1 elemen kemudian menyisipkan elemen terakhir array hasil dan elemen terakhir array input pada akhir array hasil. Contoh komputasi CUDA yang akan dijelaskan akan memakai exclusive scan. 2.4.1.3 Scan Sekuensial Pengimplementasian versi sekuensial dari scan (contohnya yang berjalan pada thread tunggal pada 1 CPU) adalah trivial. Cukup dengan melakukan looping pada setiap elemen array-input dan menjumlahkan nilai dari elemen array-input yang sebelumnya dengan hasil penjumlahan yang telah dihitung untuk elemen array-output, dan menulis hasilnya pada elemen array-output yang sedang diiterasi. void scan (float *output, float *input, int length) { output[0] = 0; // karena ini adalah prescan, bukan scan for (int j = 1; j < length; j++) { output[j] = input[j-1] + output [j-1]; } } Listing Code II-1: Algoritma sekuensial scan
II-10
Listing Code II-1 menjelaskan penjumlahan sebanyak n-kali untuk array nelemen; ini adalah jumlah minimum penjumlahan yang diperlukan untuk menghasilkan array-output. Ketika versi paralel dibuat, algoritmanya harus workefficient. Artinya, algoritma ini tidak boleh melakukan lebih banyak penjumlahan dibandingkan versi sekuensialnya. Atau dengan kata lain, kedua implementasi harus memiliki kompleksitas yang paling tidak sama, yakni O(n). 2.4.1.4 NaΓ―ve Parallel Scan for d := 1 to log2n do forall k in parallel do if k >= 2d then x[k] := x[k-2d-1] + x[k] Algoritma II-1: Algoritma paralel scan yang tidak efisien
Pseudocode pada Algoritma II-1 menunjukkan implementasi naΓ―ve parallel scan. Sedangkan ilustrasinya dapat dilihat pada Gambar II-5. Permasalahan pada Algoritma II-1 akan terlihat apabila work-complexity-nya diperiksa. Algoritma ini melakukan
log 2 π π=1
π β 2πβ1 = O (n log2n) operasi penjumlahan, sedangkan
algoritma sekuensial melakukan O(n) penjumlahan. Faktor log2n dapat berpengaruh besar pada performa komputasi. Dalam kasus 1 juta elemen, perbedaan performa antara naΓ―ve parallel scan dengan algoritma yang workefficient dapat memiliki faktor 20. Algoritma II-1 mengasumsikan bahwa ada banyak prosesor sejumlah elemen data. Pada sebuah GPU yang menjalankan CUDA, hal ini tidak selalu berlaku. Keyword forall akan dipecah secara otomatis menjadi batch paralel kecil (dinamakan
warps)
yang
dieksekusi
secara
sekuensial
pada
sebuah
multiprocessor. Sebuah GPU G80 mengeksekusi sebuah warp yang terdiri atas 32 thread secara paralel. Karena tidak semua thread berjalan secara bersamaan untuk array yang lebih besar daripada ukuran warp, algoritma di atas tidak akan berjalan sebab ia menulis hasil scan pada array-input. Hasil dari warp yang satu akan dioverwrite oleh thread pada warp yang lain. Untuk menyelesaikan masalah ini, sebuah array double-buffer dibutuhkan untuk melakukan scan. Array temporer (temp[2][n]) dibutuhkan untuk melakukan ini. Pseudocode untuk algoritma ini diberikan pada Algoritma II-2, dan kode II-11
CUDA C untuk naΓ―ve scan diberikan pada Listing Code II-2. Perlu diperhatikan bahwa kode ini akan berjalan hanya pada thread block tunggal di GPU, sehingga ukuran array yang dapat diproses terbatas (hanya 512 elemen pada GPU G80). Ekstensi algoritma untuk array berukuran sangat besar akan dibahas selanjutnya. for d := 1 to log2n do forall k in parallel do if k >= 2d then x[out][k] := x[in][k-2d-1] + x[in][k] else x[out][k] := x[in][k] swap(in, out) Algoritma II-2: Versi double buffer dari parallel scan
Gambar II-5: Melakukan scan pada array 8 elemen dengan algoritma naive parallel scan
__global__ void scan(float *g_odata, float *g_idata, int n) { extern __shared__ float temp[]; // dialokasi saat invokasi int thid = threadIdx.x; int pout = 0, pin = 1; // load inputs into shared memory // this is exclusive scan, so shift right by one and set the first elmt to 0 temp[pout*n + thid] = (thid > 0)? g_idata[thid-1] : 0; __syncthreads(); for (int offset = 1; offset < n; offset *= 2) { pout = 1 - pout; // swap index double buffer
II-12
pin = 1 - pout; if (thid >= offset) temp[pout*n + thid] += temp[pin*n + thid offset]; else temp[pout*n + thid] = temp[pin*n + thid]; __syncthreads(); } g_odata[thid] = temp[pout*n+thid]; // write output } Listing Code II-2: Kode CUDA C untuk algoritma naΓ―ve scan. Versi ini hanya dapat menangani array yang besar maksimumnya sama dengan jumlah maksimum thread dalam sebuah thread block
2.4.1.5 Algoritma Parallel Scan yang Work-Efficient Dalam sub-bab ini, akan dijelaskan sebuah algoritma scan yang work-efficient, yang menghindari faktor ekstra log2n yang dihasilkan oleh naΓ―ve scan pada subbab sebelumnya. Untuk melakukan ini, akan digunakan sebuah algorithmicpattern yang sering digunakan dalam komputasi paralel: balanced-tree. Idenya adalah membangun balanced binary tree pada input data dan menyapunya ke dan dari akar untuk menghitung prefix sum. Dalam algoritma scan yang work-efficient ini, dua buah fase operasi dilakukan pada array di shared memory, yakni: fase reduksi (juga disebut fase up-sweep) dan fase down-sweep. Dalam fase reduksi, traversal dilakukan pada pohon dari daun untuk menghitung jumlah parsial pada node internal dalam pohon, seperti yang ditunjukkan pada Gambar II-6. Ini dikenal juga dengan nama reduksi paralel, karena sesudah fase ini node root (node terakhir pada array) akan berisi hasil dari seluruh node pada array. Pseudocode untuk fase reduksi diberikan pada Algoritma II-3. Dalam fase down-sweep, traversal balik dilakukan pada pohon menggunakan hasil penjumlahan parsial untuk membangun scan pada array dengan menggunakan hasil penjumlahan parsial pada fase reduksi. Fase down-sweep ditunjukkan pada Gambar II-7 dan pseudocode ditunjukkan pada Algoritma II-4. Karena ini adalah exclusive scan, di antara dua fase elemen terakhir diisi dengan nol. Angka nol ini terpropagasi balik ke elemen pertama array dalam fase down-sweep. Kode CUDA C untuk algoritma selengkapnya diberikan dalam Listing Code II-3. Seperti kode II-13
naΓ―ve scan pada sub-bab sebelumnya, kode pada Listing Code II-3 akan berjalan pada thread block tunggal. Karena ia memproses dua elemen array per thread, maka ukuran array maksimal pada kode ini adalah 1024 elemen pada G80. Scan dari array berukuran besar akan dijelaskan pada sub bab berikutnya. Algoritma scan ini melakukan O(n) operasi (yakni 2 * (n-1) penjumlahan dan n-1 swap); sehingga algoritma ini bekerja secara efisien pada array berukuran besar daripada algoritma naΓ―ve pada sub bab sebelumnya. Meskipun begitu, efisiensi algoritma belumlah cukup: algoritma tersebut harus berjalan pada hardware secara efisien. Jika operasi scan ini diperiksa pada GPU yang menjalankan CUDA, akan terlihat bahwa algoritma ini mengalami banyak konflik shared memory bank (memory bank conflict). Hal ini mengurangi performa pada setiap akses shared memory, dan berefek negatif secara signifikan pada performa keseluruhan. Dalam sub-bab berikutnya akan dibahas sedikit modifikasi sederhana yang dapat dilakukan untuk menghitung alamat memori pada GPU untuk mengembalikan banyak performa yang hilang. for d := 0 to log2(n-1) do for k from 0 to n-1 by 2d+1 in parallel do x[k+d2+1-1] = x[k+2d-1] + x[k+2d+1-1] Algoritma II-3: fase reduksi (up-sweep) dari algoritma scan yang work-efficient
Gambar II-6: Ilustrasi dari fase reduksi (up-sweep) algoritma scan yang efisien
II-14
x[n-1] := 0 for d := log2n down to 0 do for k from 0 to n-1 by 2d+1 in parallel do t := x[k+2d-1] x[k+2d-1] := x[k+2d+1-1] x[k+2d+1-1] := t + x[k+2d+1-1] Algoritma II-4: Fase down-sweep dari algoritma parallel-sum yang work-efficient
Gambar II-7: Ilustrasi fase down-sweep dari algoritma parallel-sum yang work-efficient. Perhatikan bahwa langkah pertama me-nol-kan elemen terakhir array
__global__ void prescan(float *g_odata, float *g_idata, int n) { extern __shared__ float temp[]; // dialokasi saat invokasi int thid = threadIdx.x; int offset = 1; // Blok A temp[2*thid] = g_idata[2*thid];
// load input to shared
memory temp[2*thid+1] = g_idata[2*thid+1]; for (int d = n>>1; d > 0; d >>= 1) // build sum in place up the tree { __syncthreads(); if (thid < d)
II-15
{ // Blok B int ai = offset*(2*thid+1)-1; int bi = offset*(t*thid+2)-1; temp[bi] += temp[ai]; } offset *= 2; } // Blok C if (thid == 0) temp[n-1] = 0; // set last element = 0 for (int d = 1; d < n; d *= 2) { offset >= 1; __syncthreads(); if (thid < d) { // Blok D int ai = offset*(2*thid+1)-1; int bi = offset*(2*thid+2)-1; float t = temp[ai]; temp[ai] = temp[bi]; temp[bi] += t; } } __syncthreads(); // Blok E g_odata[2*thid] = temp[2*thid]; // write result to shared memory g_odata[2*thid+1] = temp[2*thid+1]; } Listing Code II-3: Kode CUDA untuk algoritma parallel scan yang work-efficient
2.4.1.6 Menghindari Konflik Bank Algoritma scan pada sub-bab sebelumnya memiliki kompleksitas yang kira-kira sama dengan algoritma sekuensial. Akan tetapi, algoritma ini masih belum efisien pada nVidia CUDA karena pola akses memorinya. Seperti yang dijelaskan pada [NVI07], shared memory yang dimanfaatkan oleh algoritma scan ini terdiri atas banyak bank. Ketika banyak thread dalam satu warp mengakses bank yang sama, bank-conflict terjadi, kecuali seluruh thread dalam warp mengakses sebuah
II-16
alamat dalam 32-bit word yang sama. Jumlah thread yang mengakses sebuah bank dinamakan bank conflict degree. Konflik bank menimbulkan serialisasi dari akses bersamaan pada sebuah bank memory, sehingga sebuah akses shared memory dengan derajat n membutuhkan n kali cycle daripada sebuah akses tanpa konflik. Pada G80, yang mengeksekusi 16 thread secara paralel pada setengah warp, kasus terburuk adalah konflik berderajat 16. Algoritma pohon biner seperti algoritma scan yang work-efficient menggandakan stride antara memori akses pada setiap level dalam pohon, sehingga menggandakan pula jumlah thread yang mengakses bank yang sama. Untuk pohon yang dalam, ketika algoritma ini mencapai level tengah pohon, derajat konflik bank akan meningkat, dan menurun lagi di dekat akar ketika jumlah thread yang aktif berkurang (karena statement if pada Listing Code II-3). Sebagai contoh, bila scan dilakukan pada array 512-elemen, pembacaan dan penulisan shared memory dalam inner loop dari Listing Code II-3 akan menimbulkan konflik bank berderajat 16. Ini akan menimbulkan efek yang signifikan pada performa komputasi. Konflik bank dapat dihindari apabila array memory __shared__ diperhatikan. Sebagai contoh, dalam konvolusi, ini hanya masalah bagaimana melakukan padding pada array 2D dengan sebuah ukuran yang tidak habis dibagi oleh jumlah shared memory bank. Scan, karena pendekatannya yang memakai balanced-tree, memerlukan pendekatan yang agak rumit. Konflik bank pada scan dapat dihindari dengan menambahkan padding yang besarnya bervariasi tergantung pada index memory array yang akan dihitung. Secara spesifik, sebuah nilai ditambahkan pada index, yakni nilai dari index tersebut ditambah jumlah shared memory bank. Hal ini ditunjukkan pada Gambar II-8. Dengan melihat pada Listing Code II-3, blok yang disoroti dari A sampai E dimodifikasi. Untuk memudahkan perubahan kode, maka didefinisikan sebuah makro CONFLICT_FREE_OFFSET, ditunjukkan pada Listing Code II-4.
II-17
#define NUM_BANKS 16 #define LOG_NUM_BANKS 4 #ifdef ZERO_BANK_CONFLICTS #define CONFLICT_FREE_OFFSET(n) \ ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS)) #else #define CONFLICT_FREE_OFFSET(n) ((n) >> LOG_NUM_BANKS) #endif Listing Code II-4: Macro CONFLICT_FREE_OFFSET
Blok A sampai E pada Listing Code II-3 harus dimodifikasi menggunakan macro ini untuk menghindari konflik bank. Dua perubahan harus dilakukan pada blok A. Setiap thread me-load 2 elemen array dari __device__ array g_idata ke __shared__ array temp. Dalam kode yang asli, setiap thread me-load 2 elemen berdekatan, menghasilkan indexing yang berselang-seling pada shared memory array, menimbulkan konflik bank berderajat dua. Dengan me-load dua elemen dari setengah bagian array yang berlainan, sebuah cara untuk menghindari konflik bank telah dilakukan. Juga, untuk menghindari konflik bank saat melakukan traversal pohon, sebuah padding harus ditambahkan pada shared memory array setiap NUM_BANKS(16) elemen. Hal ini dilakukan dengan menggunakan macro pada Listing Code II-4 seperti yang terlihat pada kode blok A sampai E. Perhatikan bahwa offset disimpan pada index shared memory, sehingga dapat digunakan lagi ketika scan berakhir, yakni ketika menulis hasil scan ke output-array g_odata pada blok E. Blok A: int ai = thid; int bi = thid + (n/2); int bankOffsetA = CONFLICT_FREE_OFFSET(ai); int bankOffsetB = CONFLICT_FREE_OFFSET(bi); temp[ai + bankOffsetA] = g_idata[ai]; temp[bi + bankOffsetB] = g_idata[bi]; Listing Code II-5: Perubahan pada blok A
II-18
Gambar II-8: Padding sederhana diterapkan pada alamat memori dapat menghilangkan konflik bank berderajat tinggi pada algoritma berbasis tree seperti scan. Diagram bagian atas menunjukkan pengalamatan tanpa padding dan konflik bank yang dihasilkan. Bagian bawah menunjukkan pengalamatan dengan padding tanpa konflik bank.
II-19
Blok B dan D identik: int ai = offset * (2 * thid + 1) - 1; int bi = offset * (2 * thid + 2) - 1; ai += CONFLICT_FREE_OFFSET(ai); bi += CONFLICT_FREE_OFFSET(bi); Listing Code II-6: Perubahan pada blok B dan D
Blok C: if (thid == 0) { temp[n - 1 + CONFLICT_FREE_OFFSET(n-1)] = 0; } Listing Code II-7: Perubahan pada blok C
Blok E: g_odata[ai] = temp[ai + bankOffsetA]; g_odata[bi] = temp[bi + bankOffsetB]; Listing Code II-8: Perubahan pada blok E
2.4.1.7 Array dengan Ukuran Bebas Algoritma yang diberikan dalam sub-bab sebelumnya melakukan scan pada array dalam sebuah thread block. Hal ini tidak bermasalah pada array berukuran kecil, yang ukuran maksimalnya dua kali jumlah maksimum thread dalam sebuah thread block. Dalam GPU G80, hal ini membatasi hanya 1024-elemen. Jumlah elemen juga harus merupakan kelipatan 2. Dalam sub-bab ini akan dijelaskan bagaimana menyempurnakan algoritma untuk melakukan scan pada array berukuran bebas (tidak pangkat 2). Ide dasarnya sederhana: Sebuah array berukuran besar dibagi menjadi blok-blok, dimana setiap blok dapat di-scan oleh sebuah thread block tunggal, dan menulis sum dari setiap blok pada array-of-block-sum lain. Setelah itu, scan dilakukan pada block sum tersebut, membangkitkan sebuah array dari inkremen blok yang ditambahkan pada seluruh elemen pada bloknya yang sesuai.
II-20
Untuk lebih detail, perhatikan contoh berikut: Misalkan N adalah jumlah elemen array input IN, dan B adalah jumlah elemen yang diproses oleh sebuah blok. Oleh karena itu, N/B thread block akan dialokasi, dengan B/2 thread setiap thread block-nya (pada sub-bab ini, diasumsikan bahwa N adalah kelipatan B dan akan disempurnakan untuk jumlah elemen bebas pada paragraf selanjutnya). Pilihan umum untuk B adalah 512 pada G80. Algoritma scan pada sub-bab sebelumnya digunakan untuk melakukan scan pada setiap blok j secara independen, menyimpan hasil scan pada lokasi sekuensial array-output OUT.
Gambar II-9: Algoritma untuk melakukan scan pada array berukuran besar
Selain itu, sebuah modifikasi minor dilakukan pada algoritma scan. Sebelum menol-kan elemen terakhir dari blok j (label B), total sum dari blok j disimpan pada sebuah array bantu SUMS. Array ini kemudian dapat di-inkremen dengan cara yang sama, menuliskan jumlah pada array INCR. Selanjutnya, INCR(j) ditambahkan pada setiap elemen block j menggunakan sebuah uniform add kernel
II-21
pada N/B thread block yang berisi B/2 thread. Hal ini ditunjukkan pada Gambar II-9. Untuk menangani dimensi array yang tidak kelipatan pangkat 2, array tersebut tinggal dibagi menjadi kelipatan B elemen dan diproses dengan cara yang sama seperti di atas (menggunakan B/2 thread per block), dan memproses sisanya dengan kernel yang dimodifikasi untuk menangani elemen yang bukan kelipatan 2 dalam sebuah block. Kernel ini melakukan padding pada shared memory array dan menginisialisasi memori ekstra ini dengan nol ketika data di-load dari device memory ke shared memory.
II-22
