PEMROSESAN PARALEL PADA NVIDIA CUDA
Oleh: Satya Kurnia NIM: 622007006
Skripsi Untuk melengkapi salah satu syarat memperoleh Gelar Sarjana Teknik Program Studi Sistem Komputer Fakultas Teknik Elektronika dan Komputer Universitas Kristen Satya Wacana
Salatiga April 2013
PEMROSESAN PARALEL PADA NVIDIA CUDA
oleh Satya Kurnia NIM: 622007006
Skripsi ini telah diterima dan disahkan Untuk melengkapi salah satu syarat memperoleh Gelar Sarjana Teknik dalam Konsentrasi Sistem Embedded Program Studi Sistem Komputer Fakultas Teknik Elektronika dan Komputer Universitas Kristen Satya Wacana Salatiga
Disahkan oleh:
Pembimbing I
Pembimbing II
Hartanto K.Wardana, M.T.
Darmawan Utomo,M.Eng.
Tanggal:
Tanggal:
PERNYATAAN BEBAS PLAGIAT
Saya, yang bertanda tangan di bawah ini: NAMA
: Satya Kurnia
NIM
: 622007006
JUDUL
: PEMROSESAN PARALEL PADA NVIDIA CUDA
SKRIPSI Menyatakan bahwa skripsi tersebut di atas bebas plagiat. Apabila ternyata ditemukan unsur plagiat di dalam skripsi saya, maka saya bersedia mendapatkan sanksi apa pun sesuai aturan yang berlaku.
Salatiga, 1 April 2013
Satya Kurnia
INTISARI
Nvidia CUDA adalah salah satu teknologi pemrosesan paralel yang menggabungkan Central Processing Unit (CPU) dan Graphics Processing Unit (GPU) untuk meningkatkan kinerja pemrosesan data secara paralel. Di Indonesia pemanfaatan Nvidia CUDA untuk pemrosesan paralel belum dikenal secara luas. Oleh karena itu, pada skripsi ini dibuat modul pedoman pembelajaran dan modul pedoman praktikum untuk matakuliah pemrosesan paralel. Modul pedoman pembelajaran terdiri dari 7 topik pedoman dan modul praktikum terdiri dari 8 topik pedoman. Pedoman pembelajaran dan pedoman praktikum telah melalui pengujian terhadap 43 mahasiswa yang mengambil matakuliah dasar pemrograman, pengujian yang dilakukan melalui pengisian kuesioner. Setelah dilakukan pengisian kuesioner, tahap selanjutnya melakukan analisis data melalui uji hipotesis. Uji validalitas dan uji reliabilitas belum dapat dilaksanakan. Hasil dari uji hipotesis adalah responden tidak puas dengan pedoman pembelajaran tetapi puas dengan pedoman praktikum. Selain itu juga dilakukan pengujian dalam bentuk praktikum terhadap mahasiswa dengan prasyarat matakuliah pemrosesan paralel. Hasil dari pengujian praktikum adalah 5 dari 6 topik praktikum yang diujikan telah melampaui standar penilaian, sedangkan 1 topik praktikum masih membutuhkan perbaikan. Selain itu tidak ada perbedaan nilai yang signifikan antara mahasiswa yang mempunyai ipk kurang dari 3 dengan mahasiswa yang mempunyai ipk lebih dari sama dengan 3.
i
ABSTRACT
The development of computer technology are increasingly demanding advanced data processing technology solutions faster and cheaper. Nvidia CUDA is a parallel processing technologies that combine the Central Processing Unit (CPU) and Graphics Processing Unit (GPU) to improve the performance of processing data in parallel. In Indonesia, the use of Nvidia CUDA parallel processing has not been widely known. Therefore, in this thesis is the module learning guidance and lab guidance module for parallel processing course. Module learning guidance consists of 7 topics guidelines and lab module consists of 8 topic guidelines. Guidelines learning and lab guidelines have gone through testing against 43 students taking the basic course programming, testing conducted through questionnaires. After filling out the questionnaire, the next stage of data analysis through hypotheses test. Validity and reliability tests can not be executed. The results of hypothesis testing are the respondents are not satisfied with the learning guidelines but satisfied with lab guidelines. There was also tested in a lab for students with prerequisite course parallel processing. The results of the testing lab is 5 of 6 topics lab tested has exceeded the standard appraisal practice and 1 topic still needs improvement. In addition there was no significant difference in score between the students who have less ipk than 3 with students who have ipk more than equal to 3.
ii
KATA PENGANTAR Puji Syukur ke hadirat Tuhan Yang Maha Esa, atas segala berkat dan karunia-Nya sehingga penulis dapat menyelesaikan studi dan skripsi di Fakultas Teknik Elektronika dan Komputer Universitas Kristen Satya Wacana. Penulis ingin mengucapkan terima kasih sebesar-besarnya kepada: 1. Papa, mama, semua kakak-kakak dan adik-adik yang telah memberikan dukungan materiil dan moril selama penulis menjalani kuliah di FTEK UKSW. 2. Bapak Hartanto K.W selaku pembimbing I. 3. Bapak Darmawan Utomo selaku pembimbing II. 4. Semua teman-teman yang membantu dalam pengujian, Ricky, Budhi, Vincent, Peg-peg, Lundy, Mima dan
Daniel, terimakasih atas semua dukungan dan
bantuan selama pengujian skripsi. 5. Teman-teman Lab XT , mbahe, Pakko, Heri, Black, Penda,
Theo, Codot,
Vincent, Patria, Onne, Widji dan semua teman-teman lab XT yang tidak bisa penulis sebutkan satu per satu. 6. Teman-teman Elektro 07, Eko, Rhino, Dwi, Codot, Evan, Putu, Agus, Indra, Tama, Dede dan semua teman -teman 07 yang tidak bisa penulis sebutkan satu per satu. 7. Semua teman-teman peserta kelas dasprog. 8. Mbak Yeti dan Mbak Ida. 9. Semua anggota jemaat MK, para tetangga rumah, orang-orang yang selalu menanyakan “kapan lulus?” dan semua orang yang tidak bisa penulis sebutkan satu per satu. Terimakasih atas dukungan dan doanya selama ini. Walaupun masih banyak kekurangan dalam pembuatan skripsi ini, penulis berharap skripsi ini dapat berguna dan dapat disempurnakan suatu hari kelak. Salatiga, 1 April 2013 Penulis,
Satya Kurnia iii
DAFTAR ISI
INTISARI .............................................................................................................................. i ABSTRACT .......................................................................................................................... ii KATA PENGANTAR.........................................................................................................iii DAFTAR ISI ....................................................................................................................... iv DAFTAR GAMBAR........................................................................................................... vi DAFTAR TABEL .............................................................................................................. vii DAFTAR SINGKATAN...................................................................................................viii BAB I PENDAHULUAN..................................................................................................... 1 1.1. Tujuan ........................................................................................................................ 1 1.2. Latar Belakang Masalah............................................................................................. 1 1.3. Spesifikasi .................................................................................................................. 4 1.4. Sistematika Penulisan ................................................................................................ 6 BAB II DASAR TEORI ...................................................................................................... 7 2.1. Sejarah Nvidia CUDA ............................................................................................... 7 2.2. Arsitektur Nvidia CUDA ........................................................................................... 8 2.3. GPU Sebagai Mesin Pemrosesan Paralel ................................................................. 12 2.4. Statistika ................................................................................................................... 15 2.4.1. Skala Likert ....................................................................................................... 15 2.4.2. Uji Chi-Square .................................................................................................. 16 2.4.3. Uji t ................................................................................................................... 16 BAB III PERANCANGAN ............................................................................................... 18 3.1. Pedoman Pembelajaran Nvidia CUDA .................................................................... 18 3.1.1. Pedoman pembelajaran topik 1: Konsep pemrosesan paralel, sejarah dan pengenalan Nvidia CUDA sebagai mesin pemrosesan paralel.........................19 3.1.2. Pedoman pembelajaran topik 2 : Nvidia CUDA programming model ............ 20 3.1.3. Pedoman pembelajaran topik 3: Nvidia CUDA Threading .............................. 20 3.1.4. Pedoman pembelajaran topik 4: Nvidia CUDA Memory ................................. 21 3.1.5. Pedoman pembelajaran topik 5: Nvidia CUDA Compilation and API ............ 22 3.1.6. Pedoman pembelajaran topik 6: Nvidia CUDA Library .................................. 23 iv
3.1.7. Pedoman pembelajaran topik 7: Nvidia CUDA Optimization.......................... 24 3.2. Pedoman Praktikum Nvidia CUDA ......................................................................... 26 3.2.1. Pedoman praktikum topik 0 : Pengenalan Nvidia CUDA dan Instalasi Nvidia CUDA Toolkit, Nvidia CUDA SDK, Nvidia Nsight dan Integrasi Nvidia CUDA dengan Visual Studio 2008 dan Visual Studio 2010 ................ 27 3.2.2. Pedoman praktikum topik 1: Nvidia CUDA Basic ........................................... 28 3.2.3. Pedoman praktikum topik 2 : Nvidia CUDA Threading .................................. 29 3.2.4. Pedoman praktikum topik 3 : Nvidia CUDA Memory ..................................... 30 3.2.5. Pedoman praktikum topik 4 : Nvidia CUDA API ............................................ 32 3.2.6. Pedoman praktikum topik 5 : Nvidia CUDA Library ...................................... 33 3.2.7. Pedoman praktikum topik 6 : Nvidia CUDA Optimization .............................. 35 3.2.8. Pedoman praktikum 7 : Tugas Rancang ........................................................... 36 3.3. Rancangan Pengujian ............................................................................................... 37 3.3.1. Pengujian oleh dosen ........................................................................................ 37 3.3.2. Pengujian oleh mahasiswa ................................................................................ 37 3.4. Pelaksanaan Pengujian Pedoman Pembelajaran dan Pedoman Praktikum ............. 38 BAB IV PENGUJIAN PEDOMAN .................................................................................. 41 4.1. Pengujian oleh Dosen .............................................................................................. 41 4.2. Pengujian oleh mahasiswa ....................................................................................... 41 4.2.1. Uji Hipotesis ..................................................................................................... 43 4.2.2. Analisis Pengujian Pedoman Praktikum ........................................................... 47 BAB V KESIMPULAN DAN SARAN ............................................................................. 50 5.1. Kesimpulan .............................................................................................................. 50 5.2. Saran Pengembangan ............................................................................................... 50 DAFTAR PUSTAKA ......................................................................................................... 51 LAMPIRAN A.................................................................................................................... 53 LAMPIRAN B .................................................................................................................... 63 LAMPIRAN C.................................................................................................................... 74 LAMPIRAN D.................................................................................................................... 93
v
DAFTAR GAMBAR
Gambar 1.1. Core GPU.....................................................................................................2 Gambar 1.2. Floating-Point Operations per second CPU dan GPU.................................2 Gambar 1.3. Memory Bandwidth CPU dan GPU..............................................................3 Gambar 2.1. Arsitektur komputer dengan Nvidia CUDA.................................................8 Gambar 2.2. GPU pada Nvidia CUDA............................................................................11 Gambar 2.3. SM pada arsitektur Fermi...........................................................................11 Gambar 2.4. Perbandingan Kemampuan Floating-Point CPU dengan GPU.................12 Gambar 2.5. Perbandingan Memory Bandwidth CPU dengan GPU...............................13 Gambar 2.6. Diagram Alir Kerja Nvidia CUDA.............................................................15 Gambar 3.1. Diagram Alir Pedoman Pembelajaran........................................................18 Gambar 3.2. Diagram Alir Pedoman Praktikum.............................................................26
vi
DAFTAR TABEL
Tabel 1.1. Pedoman Pembelajaran.....................................................................................4 Tabel 1.2. Pedoman Praktikum..........................................................................................5 Tabel 2.1. Sejarah Perkembangan GPU............................................................................7 Tabel 2.2. Generasi Arsitektur Nvidia CUDA.................................................................10 Tabel 2.3. Perbedaan Filosofi Desain CPU dengan GPU................................................14 Tabel 3.1. Pelaksanaan Pengujian Pedoman Pembelajaran dan Pedoman Praktikum....38 Tabel 4.1. Pernyataan Kuesioner Pedoman Pembelajaran..............................................42 Tabel 4.2. Pernyataan Kuesioner Pedoman Praktikum...................................................43 Tabel 4.3. Tingkat Kesetujuan Responden Pedoman Pembelajaran...............................44 Tabel 4.4. Test Chi-Square Pedoman Pembelajaran.......................................................45 Tabel 4.5. Tingkat Kesetujuan Responden Pedoman Praktikum....................................46 Tabel 4.6. Test Chi-Square Pedoman Praktikum............................................................46 Tabel 4.7. Nilai Rata-Rata Pengujian..............................................................................47 Tabel 4.8. Nilai Rata-Rata Pengujian berdasarkan IPK..................................................47 Tabel 4.9. Hasil Uji t Pedoman Praktikum.....................................................................49
vii
DAFTAR SINGKATAN
Singkatan ALU API CPU CUDA DRAM FPU GDDR GFLOPS GPGPU GPU GHz IPK ISA MATLAB Mhz PC PFLOPS SDK SFU SM SP TFLOPS VGA 2D 3D
Kepanjangan Arithmetic And Logic Unit Application Programming Interface Central Processing Unit Compute Unified Device Architecture Dynamic Random Access Memory Floating Point Unit Graphics Double Data Rate Giga Floating Point Operations Per Second General Purpose Graphics Processing Unit Graphics Processing Unit Gigahertz Indeks Prestasi Kumulatif Instruction Set Architecture Matrix Laboratory Megahertz Personal Computer Peta Floating Point Operations Per Second Software Development Kit Special Function Unit Streaming Multiprocessor Streaming Processor Tera Floating Point Operations Per Second Video Graphics Array 2 Dimensional 3 Dimensional
viii
BAB I PENDAHULUAN
Pada bab ini akan dijelaskan tujuan, latar belakang masalah, spesifikasi tugas dan sistematika penulisan skripsi.
1.1. Tujuan Mempelajari arsitektur dan pemrograman Nvidia CUDA yang merupakan teknologi pemrosesan paralel yang menggabungkan Central Processing Unit (CPU) dengan Graphics Processing Unit (GPU) untuk meningkatkan kinerja pemrosesan data dan membuat pedoman pengajaran dalam bentuk diktat dan slide serta membuat pedoman praktikum untuk matakuliah Pemrosesan Paralel.
1.1. Latar Belakang Masalah Seiring dengan kemajuan teknologi, pemanfaatan komputer menjadi sangat luas dan meliputi segala aspek. Komputer menjadi alat bantu manusia untuk menyelesaikan berbagai bentuk permasalahan dari permasalahan yang sederhana sampai dengan permasalahan yang kompleks. Semakin majunya kekuatan pemrosesan komputer, menjadi alasan mengapa permasalahan yang kompleks bisa diselesaikan dalam waktu yang singkat. CPU merupakan
unit pemrosesan utama di
dalam
sebuah
komputer.
Perkembangan CPU dalam 30 tahun terakhir telah meningkat pesat dari clock speeds 1 MHz pada awal tahun 1980 hingga mencapai clock speeds 4 GHz pada saat ini. Walaupun clock speeds CPU meningkat, tetapi masalah boros daya dan panas yang dihasilkan CPU akibat tingginya clock speeds menyebabkan clock speeds sulit untuk ditingkatkan lagi, sehingga merupakan masalah besar yang harus dihadapi tiap pabrikan CPU untuk meningkatkan kinerja CPU [12, h.2-3]. Pada tahun 2006 Nvidia mengenalkan teknologi pemrosesan paralel yang cepat, murah dan efisien yaitu Compute Unified Device Architecture (CUDA) yang berbasis General Purpose Graphics Processing Unit (GPGPU). Arsitektur Nvidia CUDA dapat meningkatkan kinerja pemrosesan data di dalam komputer dengan cara menggabungkan kekuatan pemrosesan CPU dan GPU untuk meningkatkan kinerja pemrosesan data.
1
GPU pada Nvidia CUDA mempunyai puluhan hingga ratusan core, sehingga dapat melakukan pemrosesan data secara bersama-sama dalam satu waktu. Gambar 1.1 merupakan gambar banyaknya core di dalam GPU pada Nvidia CUDA.
Gambar 1.1. Core GPU [1] Dengan banyaknya core di dalam sebuah GPU maka kemampuan Floating-point Operations per Second GPU jauh di atas CPU yang hanya mempunyai beberapa unit core. Saat ini kemampuan Floating-point Operations per Second GPU telah mencapai lebih dari 1,5 TFLOPS jauh di atas CPU yang baru mencapai angka 100 GFLOPS. Gambar 1.2 adalah gambar grafik perbandingan CPU dengan GPU untuk kemampuan Floating-point Operations per Second.
Gambar 1.2. Floating-Point Operations per second CPU dan GPU [8, h.2] 2
Fungsi utama GPU adalah untuk rendering grafik, oleh karena itu diperlukan memory bandwith yang lebar. Jenis memory yang digunakan pada GPU adalah Graphic Double Data Rate (GDDR) yang pada saat ini telah mencapai generasi GDDR5 yang mempunyai kecepatan 6000 MHz [16], melebihi kecepatan main memory pada CPU yang baru mencapai generasi DDR3 dengan kecepatan tertinggi 2000 MHz. Lebarnya memory bandwidth GPU maka merupakan salah satu keunggulan GPU dibanding dengan CPU dalam hal kecepatan akses memory. Gambar 1.3 adalah gambar grafik perbandingan memory bandwidth CPU dengan GPU.
Gambar 1.3. Memory Bandwidth CPU dan GPU [8, h.2]
Keunggulan yang dimiliki CUDA yaitu: 1. Antarmuka CUDA application menggunakan standar bahasa C sehingga mudah dipelajari. 2. Data transfer antara CPU dan GPU sangat cepat sehingga dapat mengurangi latency pengiriman data dari CPU ke GPU ataupun sebaliknya. 3. Mempunyai shared memory pada setiap multiprosesor dengan kecepatan akses tinggi. 4. Hemat biaya dan daya listrik.
3
Oak Ridge National Laboratory telah menggunakan 18.688 16 core AMD Opteron CPU dan juga 18000 Nvidia CUDA GPU pada superkomputer Cray XK6 untuk menghasilkan performa 20 PFLOPS. Superkomputer Cray XK6 2 kali lebih cepat dan 3 kali lebih hemat daya daripada K superkomputer buatan Fujitsu Jepang yang menggunakan 88.128 8 core SPARC64 VIIIfx CPU yang hanya menghasilkan performa 10 PFLOPS [14]. Nvidia CUDA telah diperkenalkan pada universitas-universitas di luar negeri seperti Universitas Harvard [4] dan Institut Teknologi Tokyo [11] sejak tahun 2009. Di Indonesia teknologi Nvidia CUDA belum dikenal secara luas dikarenakan belum ada matakuliah yang yang memanfaatkan teknologi ini. Oleh karena itu, melalui skripsi ini akan dibuat pedoman pengajaran dan pedoman praktikum Nvidia CUDA yang akan diterapkan untuk matakuliah Pemrosesan Paralel. 1.3.
Spesifikasi Berdasarkan surat keputusan nomor: 21/I.3/FTEK/VI/2012 dan surat keputusan
nomor: 02/Kep/D/FTEK/III/2013 perincian tugas yang dikerjakan adalah sebagai berikut: 1. Materi yang diberikan terbagi menjadi 7 pedoman untuk pedoman pembelajaran dan 7 pedoman untuk pedoman praktikum.
Tabel 1.1. Pedoman Pembelajaran Judul
Pedoman 1
Konsep pemrosesan paralel, sejarah dan pengenalan Nvidia CUDA sebagai mesin pemrosesan paralel
2
Nvidia CUDA Programming Model
3
Nvidia CUDA Threading
4
Nvidia CUDA Memory
5
Nvidia CUDA Compilation and API
6
Nvidia CUDA Library
7
Nvidia CUDA Optimization
4
Tabel 1.2. Pedoman Praktikum Pedoman 1
Judul Pengenalan dan instalasi Nvidia CUDA Toolkit, Nvidia CUDA SDK, Nvidia Nsight dan integrasi Nvidia CUDA dengan Visual Studio 2008 dan Visual Studio 2010
2
Nvidia CUDA Basic
3
Nvidia CUDA Threading
4
Nvidia CUDA Memory
5
Nvidia CUDA API
6
Nvidia CUDA Library
7
Nvidia CUDA Optimization
8
Tugas Rancang
2. Format untuk pedoman pembelajaran adalah sebagai berikut: 1. Judul 2. Tujuan 3. Dasar Teori 4. Summary 5. Latihan Soal 6. Daftar Pustaka
Format untuk pedoman praktikum adalah sebagai berikut: 1. Judul 2. Tujuan 3. Dasar Teori 4. Contoh Program 5. Soal Praktikum 6. Set Dosen 7. Daftar Pustaka
5
3. Pembuatan program Nvidia CUDA menggunakan Visual Studio 2008 dan Visual Studio 2010. 4. Pengujian untuk mahasiswa berjumlah 6 orang dengan prasyarat matakuliah pemrosesan paralel. 5. Penilaian kemampuan mahasiswa rata-rata 70. 6. Standar penilaian kemampuan mahasiswa dengan menggunakan patokan yang telah telah ditentukan. 7. Keluaran berupa : pedoman, slide dan rekomendasi prasyarat. 8. Metode pengujian yang digunakan adalah uji Chi-Square, uji t dan uji z sedangkan skala pengujian yang digunakan adalah skala likert.
1.4. Sistematika Penulisan Penulisan skripsi ini terdiri dari 5 bab yaitu: Bab I
Pendahuluan Berisi tujuan, latar belakang masalah, spesifikasi dan sistematika penulisan.
Bab II Dasar Teori Berisi dasar teori Nvidia CUDA yang meliputi sejarah, arsitektur, pengembangan GPU sebagai mesin pemrosesan paralel dan dasar teori statistika yang akan digunakan untuk pengujian. Bab III Perancangan Sistem Berisi perancangan pedoman kuliah, pedoman praktikum dan perancangan pengujian pedoman. Bab IV Pengujian dan Analisis Berisi hasil pengujian pedoman. Bab V Kesimpulan dan Saran Berisi kesimpulan dan saran-saran yang dapat digunakan untuk pengembangan Nvidia CUDA ke tahap selanjutnya.
6
BAB II DASAR TEORI Bab ini berisi dasar teori yang berhubungan dengan perancangan skripsi antara lain sejarah Nvidia CUDA, arsitektur Nvidia CUDA, GPU sebagai mesin pemrosesan paralel dan dasar teori untuk statistika 2.1. Sejarah Nvidia CUDA 2006-Nvidia mengenalkan teknologi pemrosesan paralel berbasis General Purpose Graphics Processing Unit (GPGPU) yaitu Compute Unified
Device
Architecture (CUDA). Nvidia bertujuan membuat Graphics Processing unit (GPU) tidak hanya digunakan untuk mengolah aplikasi grafis tetapi juga dapat digunakan untuk memproses aplikasi non-grafis. GPU pada mulanya hanya digunakan untuk rendering grafis 2D maupun 3D, namun sekarang telah banyak aplikasi yang memanfaatkan kinerja dari GPU untuk aplikasi perhitungan matematika seperti MATLAB. Pada awal perkembangan komputer,
CPU masih memegang peranan
penting dalam pemrosesan grafis namun setelah perkembangan grafis yang makin meluas, kinerja CPU menjadi menurun karena harus melakukan rendering grafis. Pada awal tahun 1980an dibuatlah hardware khusus untuk mengolah grafis yang terdiri dari Geometry unit, Rasterization unit dan Fragment unit yang kemudian dinamakan GPU. Tabel 2.1 menunjukkan sejarah perkembangan GPU. Tabel 2.1. Sejarah Perkembangan GPU [17]
Tahun 1980an
Sejarah •
Intel mengembangkan graphics controller yang pertama berbasis multibus.
•
Commodore Amiga merupakan PC pertama yang menggunakan GPU. (bersambung)
7
Tabel 2.1. Sejarah Perkembangan GPU [17] (lanjutan)
1990an
•
Aplikasi yang menggunakan grafis 3D bermunculan terutama game.
•
Munculnya beberapa pabrikan pembuat chip GPU seperti ATI, Nvidia dan Matrox.
•
OpenGL menjadi standar grafis API.
•
Microsoft mengembangkan DirectX untuk game Microsoft yang menggunakan GPU.
2000an
•
Nvidia menerapkan programmable shading pada produknya yaitu geforce 3 yang memungkinkan geometric vertex dapat diproses ke dalam program sederhana sebelum ditampilkan ke layar.
2006
•
Nvidia mengenalkan Compute Unified Device Architecture (CUDA) yang berbasis General Purpose Graphics Processing Unit (GPGPU).
2.2. Arsitektur Nvidia CUDA Nvidia CUDA mempunyai arsitektur berbasis General Purpose Computing Graphics Processing Units (GPGPU). Nvidia CUDA dapat berjalan dengan platform CPU berbasis x86 dan x64 dari platform manapun dan chipset motherboard yang mendukung PCI-Express Bus. Gambar 2.1 merupakan gambar arsitektur komputer dengan Nvidia CUDA.
Gambar 2.1. Arsitektur komputer dengan Nvidia CUDA
Pada Gambar 2.1 Nvidia CUDA dihubungkan melalui interface PCI-Express Bus yang mendukung bandwidth sampai 16 GB/s. PCI-Express dihubungkan oleh
8
Northbridge pada motherboard yang berfungsi menjembatani CPU, DRAM, Southbridge dan PCI-Express.
Nvidia CUDA terdiri dari beberapa komponen yaitu[7, h.2]: 1. Unit pemrosesan di dalam GPU. 2. OS kernel - level yang mendukung inisialisasi dan konfigurasi hardware. 3. User mode driver yang menunjukkan device level API. 4. PTX instruction set architecture (ISA) untuk komputasi paralel kernel dan fungsi, kernel merupakan potongan program yang akan dieksekusi secara paralel oleh GPU [9, h.7].
Saat ini Nvidia CUDA telah mencapai 4 generasi arsitektur. Tabel 2.2 menunjukkan generasi arsitektur Nvidia CUDA .
9
Tabel 2.2. Generasi Arsitektur Nvidia CUDA [3]
Arsitektur
G80/G92
Jumlah
SP/ CUDA
Jumlah SP
Shared
Maksimum
maksimum
Dimensi
maksimum SM
cores
maksimum
memory/ cache
memory
thread per
block
bandwidth
block
86,4 GB/s
512
512 x 512 x 64
141,7 GB/s
512
512x 512 x 64
177,4 GB/s
1024
1024 x 1024 x 64
192,4 GB/s
1024
1024 x 1024 x 64
16
8
128
246 KB shared memory, 512 KB register
GT 200
30
8
240
480KB shared memory, 1920 KB register
Fermi
16
32
512
L1 cache 1024 KB, L2 cache 768 KB, 2048 KB register
Kepler
8
192
1536
L1 cache 512 KB, L2 cache 1536 KB, 2048 KB register
10
GPU pada Nvidia CUDA terdiri beberapa unit Streaming Multiprocessors (SM) seperti pada Gambar 2.2.
Gambar 2.2. GPU pada Nvidia CUDA [1]
Setiap SM terdiri dari beberapa Streaming Processors (SP)/CUDA cores, Setiap arsitektur mempunyai jumlah SP/CUDA cores yang berbeda-beda. Gambar 2.3 adalah contoh gambar SM pada arsitektur Fermi.
Gambar 2.3. SM pada arsitektur Fermi [10, h.8]
11
Bagian yang terdapat pada setiap SM arsitektur Fermi adalah [10, h.5-8]: • 32 CUDA core yang terdiri dari Arithmetic Logic Unit (ALU) dan Floating Point Unit (FPU). ALU pada arsitektur Fermi mendukung sampai 32 bit precision unit untuk semua instruksi sedangkan FPU mendukung sampai 64 bit. • 16 unit Load Store Unit yang digunakan untuk pengalamatan resource dan destination yang dapat menampung sampai 16 thread per clock. • 4 Special Function Unit (SFU) yang berfungsi mengeksekusi instruksi transcendental seperti sin,cos dan bilangan kuadrat. • 2 Warp scheduler untuk penjadwalan 32 thread yang bekerja secara paralel. • 64 KB Shared memory dan L1 cache. • Register serbaguna dengan ukuran 32.768 x 32 bit.
2.3. GPU Sebagai Mesin Pemrosesan Paralel Sejak tahun 2003 kemampuan Graphics Processing Unit (GPU) telah melampaui kemampuan Central Processing Unit (CPU) dalam hal kemampuan pemrosesan floating point, baik single precision maupun double precision. Gambar 2.4 menunjukkan perbandingan kemampuan GPU Nvidia CUDA dengan CPU.
Gambar 2.4. Perbandingan Kemampuan Floating-Point CPU dengan GPU [9, h.2]
Gambar 2.4 memperlihatkan perbedaan yang sangat mencolok antara kemampuan GPU Nvidia CUDA dengan CPU. CPU tercepat buatan Intel pada tahun 2011 yaitu 12
Intel Sandy Bridge kemampuannya berbeda jauh dengan GPU NVIDIA GTX 680 yang juga keluar pada tahun yang sama. Memory Bandwidth juga mempengaruhi perbedaan kemampuan CPU dengan GPU Nvidia CUDA. Gambar 2.5 menunjukkan perbandingan memory bandwidth antara GPU Nvidia CUDA dengan CPU.
Gambar 2.5. Perbandingan Memory Bandwidth CPU dengan GPU [9, h.2]
Gambar 2.5 memperlihatkan GPU Nvidia CUDA sangat unggul dalam memory bandwith, hal ini terlihat jelas dari GPU Nvidia GTX 680 yang mempunyai memory bandwith 192.2 GB/s sedangkan CPU tercepat intel Sandy Bridge hanya memiliki memory bandwidth sekitar 50 GB/s. GPU mempunyai akses memory yang cepat karena pada dasarnya GPU memang diperuntukkan untuk rendering grafis yang membutuhkan memory bandwidth yang tinggi untuk mengurangi bottleneck. Dengan perkembangan game 3D yang semakin lama membutuhkan GPU yang semakin cepat, pabrikan chip GPU pun saling berlomba-lomba membuat GPU yang mempunyai kemampuan tinggi tetapi dapat dijual dengan harga yang murah. Hal ini menyebabkan hampir setiap tahun 2 pabrikan besar pembuat chip GPU yaitu ATI dan Nvidia mengeluarkan GPU baru dengan arsitektur yang semakin baik dalam segi performa dan semakin hemat daya listrik. 13
Setelah melihat perbedaan signifikan antara CPU dengan GPU, maka dapat ditarik kesimpulan bahwa kemampuan CPU tidak sepadan dengan GPU. Hal ini disebabkan CPU dengan GPU mempunyai filosofi yang berbeda dalam desainnya. Tabel 2.3 menunjukkan perbedaan filosofi desain CPU dengan GPU.
Tabel 2.3. Perbedaan Filosofi Desain CPU dengan GPU [5, h.39-41]
CPU
GPU
Mempunyai sedikit unit pemrosesan
Mempunyai banyak unit pemrosesan
Akses ke cache memory sangat cepat
Akses ke onboard memory sangat cepat
CPU sangat baik untuk task parallelism
GPU sangat baik untuk data parallelism
Kemampuan tinggi untuk single thread
Kemampuan tinggi untuk tugas paralel
execution
Pada Tabel 2.3 memperlihatkan bahwa GPU mempunyai desain yang berbeda dengan CPU. GPU dapat melakukan apa yang CPU tidak bisa lakukan, sebaliknya CPU juga dapat melakukan apa yang GPU tidak bisa lakukan. Inilah alasan Nvidia CUDA menggabungkan CPU dengan GPU. Dengan banyaknya unit pemrosesan di dalam GPU maka data dalam jumlah banyak dapat diproses secara bersamaan dalam satu waktu. CPU akan mengeksekusi kode sequential yang berisi instruksi yang diperlukan GPU untuk melakukan pemrosesan paralel, setelah itu GPU akan mengeksekusi data yang diberikan oleh CPU secara paralel. Gambar 2.6 menunjukkan diagram alir kerja Nvidia CUDA.
14
Gambar 2.6. Diagram Alir Kerja Nvidia CUDA [15]
Penjelasannya adalah sebagai berikut: 1. Alur kerja yang pertama yaitu CPU akan menyalin semua data yang dibutuhkan GPU untuk eksekusi dari main memory ke GPU memory onboard. 2. Setelah semua data yang dibutuhkan GPU untuk proses eksekusi disalin. CPU akan mengerjakan kode sequential dan kemudian akan memberikan instruksi yang akan dikerjakan oleh GPU. 3. GPU akan mengeksekusi secara paralel. 4. Hasil dari pemrosesan paralel akan disalin ke main memory.
2.4. Statistika [6, h.70-89] [13, h.15-18] 2.4.1. Skala Likert Skala likert adalah skala psikometrik yang sering digunakan untuk kuesioner di dalam riset berupa survei. Disediakan lima pilihan skala dengan format sebagai berikut: 1. Sangat tidak setuju 2. Tidak setuju 3. Tidak pasti 4. Setuju 5. Sangat setuju 15
Langkah skala likert 1.Menentukan dan memahami apa yang akan diukur. 2.Menyusun perancangan.
2.4.2. Uji Chi-Square Uji Chi-Square bertujuan untuk menguji perbedaan proporsi dua atau lebih kelompok. Perhitungan uji Chi-Square menggunakan persamaan 2.1.
=
∑(
)
(2.1)
Dengan: Fh = Frekuensi harapan Fo = Frekuensi observasi
(2.2)
Fh = Dengan: = Total baris = Total kolom N = Total 2.4.3. Uji t
Uji t biasa digunakan untuk sampel kecil. Uji t yang digunakan adalah Uji t independen berbeda varian sesuai persamaan 2.3.
=
(
).
!(
).
16
× #
.
.(
!
!
)
(2.3)
$=
∑(% − %̅ ) (
Dengan: S12 = Standar Deviasi 1 S22 = Standar Deviasi 2 n1 = Jumlah responden kelompok 1 n2 = Jumlah responden kelompok 2 %1 = Nilai rata-rata kelompok 1
%2 = Nilai rata-rata kelompok 2
17
BAB III PERANCANGAN Pada bab ini akan dijelaskan deskripsi pedoman pengajaran dan pedoman praktikum Nvidia CUDA. 3.2. Pedoman Pembelajaran Nvidia CUDA Diagram alir pedoman pembelajaran dapat dilihat pada Gambar 3.1.
Gambar 3.1. Diagram Alir Pedoman Pembelajaran
18
Gambar 3.1 menunjukkan perancangan pedoman pembelajaran. Berikut adalah deskripsi dari masing-masing pedoman : 3.2.1. Pedoman pembelajaran topik 1: Konsep pemrosesan paralel, sejarah dan pengenalan Nvidia CUDA sebagai mesin pemrosesan paralel. Melalui pedoman pembelajaran topik 1 diharapkan mahasiswa dapat : • Mengingat kembali konsep - konsep dasar pemrosesan paralel. • Mengerti konsep Nvidia CUDA sebagai mesin pemrosesan paralel. Pemilihan topik pedoman ini didasarkan : • Perlunya mahasiswa mengingat kembali konsep - konsep dasar pemrosesan paralel. • Perlunya pengenalan konsep Nvidia CUDA sebagai mesin pemrosesan paralel.
Perancangan pedoman pengajaran : 1. Tujuan. 2. Materi 2.1. Konsep pemrosesan paralel 2.1.1. Konsep pemrosesan paralel secara umum. 2.1.2. Hukum Amdahl. 2.1.3. Platform memory. 2.1.4. Taksonomi Flyn. 2.1.5. Dekomposisi. 2.2. Pengenalan Nvidia CUDA 2.2.1. Sejarah perkembangan GPU. 2.2.2. GPU sebagai mesin pemrosesan paralel. 3. Ringkasan. 4. Soal-Soal Latihan. 5. Daftar Pustaka.
19
3.2.2. Pedoman pembelajaran topik 2 : Nvidia CUDA programming model. Melalui pedoman pembelajaran topik 2 diharapkan mahasiswa dapat: • Memahami dan mengerti arsitektur Nvidia CUDA. • Menguasai dasar - dasar programming model Nvidia CUDA. Pemilihan topik pedoman ini didasarkan : • Mahasiswa perlu mengetahui konsep dasar arsitektur dan pemrograman Nvidia CUDA sebagai dasar untuk melakukan
pemrosesan paralel
menggunakan Nvidia CUDA.
Perancangan pedoman pengajaran : 1. Tujuan. 2. Materi 2.1. Arsitektur Nvidia CUDA. 2.2. Programming Model Nvidia CUDA 2.2.1. Kernel. 2.2.2. Thread. 2.2.3. Heterogeneous Programming. 2.2.4. Management memory Nvidia CUDA. 3. Ringkasan. 4. Soal-Soal Latihan. 5. Daftar Pustaka.
3.2.3. Pedoman pembelajaran topik 3: Nvidia CUDA Threading. Melalui pedoman pembelajaran topik 3 diharapkan mahasiswa dapat: • Memahami dan mengerti konsep dasar dari thread. • Menguasai dasar-dasar pemrograman thread. • Mengerti manfaat dari penjadwalan thread.
20
Pemilihan topik pedoman ini didasarkan : • Thread merupakan kunci utama Nvidia CUDA untuk melakukan pemrosesan paralel, thread berfungsi mengeksekusi data secara paralel.
Perancangan pedoman pengajaran : 1. Tujuan. 2. Materi 2.1. Konsep Thread pada Nvidia CUDA. 2.2. Pemrograman Thread. 2.3. Penjadwalan Thread. 2.4. Transparent Scalability. 3. Ringkasan. 4. Soal-Soal Latihan. 5. Daftar Pustaka.
3.2.4. Pedoman pembelajaran topik 4: Nvidia CUDA Memory. Melalui pedoman pembelajaran topik 4 diharapkan mahasiswa dapat: • Memahami dan mengerti konsep dasar memory model pada Nvidia CUDA. • Mengerti jenis - jenis memory yang terdapat pada Nvidia CUDA beserta pemanfaatannya.
Pemilihan topik pedoman ini didasarkan : • Penggunaan memory pada Nvidia CUDA sangat penting karena penggunaan memory dapat mempengaruhi unjuk kerja program. • Perlunya mahasiswa mengetahui cara penggunaan memory pada Nvidia CUDA.
Perancangan pedoman pengajaran: 1. Tujuan. 2. Materi 2.1. Nvidia CUDA memory model. 21
2.2. GPU memory 2.2.1. Register. 2.2.2. Local memory. 2.2.3. Shared memory. 2.2.4. Constant memory. 2.2.5. Texture memory. 2.2.6. Global memory. 3. Ringkasan. 4. Soal-Soal Latihan. 5. Daftar Pustaka.
3.2.5. Pedoman pembelajaran topik 5: Nvidia CUDA Compilation and API. Melalui pedoman pembelajaran topik 5 diharapkan mahasiswa dapat: • Memahami dan mengerti Nvidia CUDA Compilation dan Nvidia CUDA API. • Menguasai 5 fungsi dasar Nvidia CUDA Runtime API yaitu device management, memory management, stream management, event management dan error handling. • Menguasai dasar - dasar dari Nvidia Driver API. Pemilihan topik pedoman ini didasarkan: • Pemanggilan fungsi API merupakan dasar pemrograman Nvidia CUDA. • Mengenalkan 5 fungsi dasar Nvidia CUDA API Runtime.
Perancangan pedoman pengajaran : 1. Tujuan. 2. Materi 2.1. Nvidia CUDA Compilation. 2.2. Nvidia CUDA Language. 2.2.1. Declaration Qualifier. 2.2.2. Built in Variables. 22
2.2.3. Built in Types. 2.2.4. Execution Configuration. 2.3. Nvidia CUDA API 2.3.1. Runtime API 2.3.1.1. Device Management. 2.3.1.2. Memory Management. 2.3.1.3. Event Management. 2.3.1.4. Stream Management. 2.3.1.5. Error Handling. 2.3.2. Driver API. 3. Ringkasan. 4. Soal-Soal Latihan. 5. Daftar Pustaka.
3.2.6. Pedoman pembelajaran topik 6: Nvidia CUDA Library. Melalui pedoman pembelajaran topik 6 diharapkan mahasiswa dapat: • Memahami penggunaan Nvidia Library untuk mempermudah user dalam memproses data secara paralel. • Menguasai penggunaan Library CUBLAS untuk operasi perhitungan vektor dan matrik. • Menguasai penggunaan Library Thrust untuk transformasi, reduksi, prefixsums dan sorting.
Pemilihan topik ini didasarkan: • Nvidia CUDA Library dapat mempermudah user dalam memproses data secara paralel menggunakan Nvidia CUDA. Perancangan pedoman pengajaran: 1. Tujuan. 2. Materi 2.1. CUBLAS 23
2.1.1. CUBLAS level 1. 2.1.2. CUBLAS level 2. 2.1.3. CUBLAS level 3. 2.2. Thrust 2.2.1. Containers. 2.2.2. Algoritma. 2.2.3. Iterators. 3. Ringkasan. 4. Soal-Soal Latihan. 5. Daftar Pustaka. 3.2.7. Pedoman pembelajaran topik 7: Nvidia CUDA Optimization. Melalui pedoman pembelajaran topik 7 diharapkan mahasiswa dapat: • Mengerti dasar-dasar optimasi Nvidia CUDA secara software dan hardware.
Pemilihan topik ini didasarkan: • Perlunya mahasiswa mengetahui cara meningkatkan unjuk kerja Nvidia CUDA melalui software dan hardware.
Perancangan pedoman pengajaran: 1. Tujuan. 2. Materi 2.1. Optimasi Software 2.1.1. Optimasi memory 2.1.1.1. Efisiensi transfer memory. 2.1.1.2. Coalescing memory. 2.1.1.3. Menggunakan shared memory secara efisien. 2.1.2. Optimasi thread. 2.1.3. Optimasi instruksi. 2.2. Optimasi Hardware. 3. Ringkasan.
24
4. Soal-Soal Latihan. 5. Daftar Pustaka.
25
3.2. Pedoman Praktikum Nvidia CUDA Diagram alir pedoman praktikum dapat dilihat pada Gambar 3.2.
Gambar 3.2. Diagram Alir Pedoman Praktikum.
26
Gambar 3.2 menunjukkan perancangan pedoman praktikum. Berikut adalah deskripsi dari masing-masing pedoman: 3.2.1. Pedoman praktikum topik 0 : Pengenalan Nvidia CUDA dan Instalasi Nvidia CUDA Toolkit, Nvidia CUDA SDK, Nvidia Nsight dan Integrasi Nvidia CUDA dengan Visual Studio 2008 dan Visual Studio 2010. Melalui pedoman praktikum topik 0 diharapkan mahasiswa menguasai konsep: • Mahasiswa dapat melakukan instalasi Nvidia CUDA Toolkit, Nvidia CUDA SDK dan Nvidia Nsight. • Mahasiswa dapat melakukan integrasi Nvidia CUDA dengan Visual Studio 2008 dan 2010. • Mahasiswa dapat mengetahui properties GPU yang digunakan melalui program GPU-Z. • Mahasiswa dapat menguji contoh program pada Nvidia CUDA SDK.
Pemilihan topik pedoman ini didasarkan : • Perlunya mahasiswa mengingat kembali konsep-konsep dasar pemrosesan paralel. • Perlunya pengenalan konsep Nvidia CUDA sebagai mesin pemrosesan paralel.
Perancangan pedoman praktikum: 1. Tujuan. 2. Mengecek jenis GPU dan versi driver pada PC atau laptop. 3. Instalasi Nvidia CUDA Toolkit. 4. Instalasi Nvidia CUDA SDK. 5. Integrasi Nvidia CUDA dengan Visual Studio 2008. 6. Integrasi Nvidia CUDA dengan Visual Studio 2010. 7. Instalasi Nvidia Nsight. 8. Kompilasi source CUDA dengan Visual Studio 2010 command prompt 9. GPU Computing SDK test. 27
10. Pemrograman Nvidia CUDA dengan Visual Studio 2008 atau Visual Studio 2010. 11. Soal Praktikum 11.1. Soal praktikum set 1: Analisa program. 11.2. Soal praktikum set 2: Analisa program. 12. Set Dosen 12.1. Set Dosen 1. 12.2. Set Dosen 2. 13. Daftar Pustaka.
3.2.2. Pedoman praktikum topik 1: Nvidia CUDA Basic. Melalui pedoman praktikum topik 1 diharapkan mahasiswa menguasai konsep: • Pembuatan program Nvidia CUDA sederhana. • Penyalinan memory, eksekusi kernel dan menampilkan hasil eksekusi. • Eksekusi thread pada kernel. Pemilihan topik pedoman ini didasarkan : • Mengenalkan CUDA template. • Dasar dari penyalinan memory, pemanggilan kernel dan eksekusi thread pada kernel. • Mengenalkan kode host dan kode device.
Perancangan pedoman praktikum: 1. Tujuan. 2. Dasar Teori 2.1. Kernel. 2.2. Thread. 2.3. CUDA template. 2.4. Alokasi memory. 28
3. Contoh Program 3.1. Contoh program 1: menghitung kuadrat dari thread. 4. Soal Praktikum set 1 4.1. Soal 1 : Pembatasan thread dari thread 0 sampai thread 3. 4.2. Soal 2 : Penjumlahan paralel. 5. Soal Praktikum set 2 5.1. Soal 1 : Pembatasan thread dari thread 7 sampai thread 4. 5.2. Soal 2 : Perkalian paralel. 6. Set Dosen 6.1. Set Dosen 1. 6.2. Set Dosen 2. 7. Daftar Pustaka.
3.2.3. Pedoman praktikum topik 2 : Nvidia CUDA Threading. Melalui pedoman praktikum topik 2 diharapkan mahasiswa menguasai konsep: • Menggunakan thread dalam 3 sumbu vektor x,y dan z. • Menggabungkan semua thread di dalam block. • Menggunakan timer untuk mengetahui waktu eksekusi GPU. Pemilihan topik pedoman ini didasarkan : • Konfigurasi thread yang digunakan dalam mengeksekusi kernel mempengaruhi unjuk kerja GPU. • Penggunaan timer penting untuk mengetahui unjuk kerja GPU. Perancangan pedoman praktikum : 1. Tujuan. 2. Dasar Teori 2.1. Konfigurasi thread. 2.2. Penggabungan thread antar block. 29
3. Contoh Program 3.1. Contoh program 1: Eksekusi kernel dengan menggunakan thread vektor x,y dan z. 3.2. Contoh program 2: Eksekusi kernel dengan menggunakan gabungan thread antar block. 3.3. Contoh program 3: Penggunaan timer untuk menghitung waktu eksekusi. 4. Soal Praktikum set 1 4.1. Soal 1: Penjumlahan koordinat dengan menggunakan vektor x dan y. 4.2. Soal 2: Perkalian matrik secara paralel. 5. Soal Praktikum set 2 5.1. Soal 1: Perkalian koordinat dengan menggunakan vektor x dan y. 5.2. Soal 2: Perkalian matrik secara paralel. 6. Set Dosen 6.1. Set Dosen 1. 6.2. Set Dosen 2. 7. Daftar Pustaka.
3.2.4. Pedoman praktikum topik 3 : Nvidia CUDA Memory. Melalui pedoman praktikum topik 3 diharapkan mahasiswa menguasai konsep: • Memory model Nvidia CUDA. • Jenis - jenis memory Nvidia CUDA. • Menentukan jenis memory yang tepat untuk mengeksekusi program Nvidia CUDA. • Penggunaan shared memory, constant memory dan global memory dalam program.
30
Pemilihan topik pedoman ini didasarkan : • Perlunya mengetahui jenis memory Nvidia CUDA. • Perlunya mengetahui cara penggunaan memory di dalam program Nvidia CUDA.
Perancangan pedoman praktikum: 1. Tujuan. 2. Dasar Teori 2.1. Memory model Nvidia CUDA. 2.2. Tipe Memory Nvidia CUDA. 3. Contoh Program 3.1. Contoh program 1: Contoh penggunaan global memory, shared memory dan constant memory kemudian membandingkan kecepatan shared memory dan global memory dalam memproses data. 4. Soal Praktikum set 1 4.1. Soal 1: Perkalian matrik secara paralel dengan menggunakan shared memory. 4.2. Soal 2: Bitonic sort. 5. Soal Praktikum set 2 5.1. Soal 1: Perkalian matrik secara paralel dengan menggunakan shared memory, data yang digunakan random. 5.2. Soal 2: Bitonic sort dengan menggunakan data random. 6. Set Dosen 6.1. Set Dosen 1. 6.2. Set Dosen 2. 7. Daftar Pustaka.
31
3.2.5. Pedoman praktikum topik 4 : Nvidia CUDA API. Melalui pedoman praktikum topik 4 diharapkan mahasiswa menguasai konsep : • Menguasai dan menerapkan 5 fungsi dasar Nvidia CUDA Runtime API yaitu device management, memory management, stream management, event management dan error handling. • Pembuatan aplikasi dengan menggunakan Nvidia CUDA API Runtime. Pemilihan topik pedoman ini didasarkan : •
Mengenalkan 5 fungsi dasar Nvidia CUDA API Runtime.
Perancangan pedoman praktikum : 1. Tujuan. 2. Dasar Teori 2.1. Nvidia CUDA Device Management. 2.2. Nvidia CUDA Memory Management. 2.3. Nvidia CUDA Event Management. 2.4. Nvidia CUDA Stream Management. 2.5. Nvidia CUDA Error Handling. 3. Contoh Program 3.1. Contoh program 1: Contoh penggunaan Nvidia CUDA Device Management. 3.2. Contoh program 2: Contoh penggunaan Nvidia CUDA Memory Management untuk penjumlahan vektor menggunakan mapping memory. 3.3. Contoh program 3: Contoh penggunaan Nvidia CUDA Event, Stream dan Error handling untuk mengecek pemanggilan API, jika ada error maka program akan menampilkan pesan error.
32
4. Soal Praktikum set 1 4.1. Soal 1: Membuat Menu untuk Searching bilangan via Runtime API secara paralel, paralel,
Mencari bilangan prima via Runtime API secara
Mencari bilangan prima dengan menggunakan mapping
memory dan memproses data menggunakan CUDA stream. 5. Soal Praktikum set 2 5.1. Soal 1: Membuat Menu untuk Searching bilangan via Runtime API secara paralel, Mencari bilangan bukan prima via Runtime API secara paralel, Mencari kelipatan bilangan dengan menggunakan mapping memory dan memproses data menggunakan CUDA stream. 6. Set Dosen 6.1. Set Dosen 1. 6.2. Set Dosen 2. 7. Daftar Pustaka.
3.2.6. Pedoman praktikum topik 5 : Nvidia CUDA Library. Melalui pedoman praktikum topik 5 diharapkan mahasiswa menguasai konsep : • Penggunakan fasilitas library Nvidia CUDA. • Penggunaan library CUBLAS dan Thrust untuk melakukan pemrosesan paralel.
Pemilihan topik pedoman ini didasarkan : • Penggunaan Nvidia CUDA Library dapat mempermudah user dalam memproses data secara paralel.
Perancangan pedoman praktikum : 1. Tujuan. 2. Dasar Teori 2.1. CUBLAS 2.1.1. CUBLAS level 1.
33
2.1.2. CUBLAS level 2. 2.1.3. CUBLAS level 3. 2.2. Thrust 2.2.1. Containers. 2.2.2. Algoritma. 2.2.3. Iterator. 3. Contoh Program 3.1. Contoh program 1: Menggunakan CUBLAS level 1 untuk menghitung operasi vektor dengan vektor menggunakan fungsi cublasSaxpy, cublasSasum dan cublasScoopy. 3.2. Contoh program 2: Menggunakan CUBLAS level 2 untuk menghitung operasi matrik dengan vektor menggunakan fungsi cublasSgemv. 3.3. Contoh program 3: Menggunakan CUBLAS level 3 untuk menghitung operasi matrik dengan matrik menggunakan fungsi cublasSgemm. 3.4. Contoh program 4: Menggunakan Thrust untuk transformasi, reduksi, prefix-sums dan sorting. 4. Soal Praktikum set 1 4.1. Soal 1: Melakukan pembacaan text file dengan format yang telah ditentukan kemudian membuat menu unuk melihat data list barang , melihat barang dari harga termurah, melihat barang dari harga termahal dan membeli barang. Semua perhitungan menggunakan Thrust Library. 4.2. Soal 2: Melakukan pembacaan data melalui text file kemudian melakukan perkalian matrik dengan matrik itu sendiri yang telah ditransformasi negasi dan dilakukan
transpose matrik. Perkalian
matrik menggunakan CUBLAS level 3. 5. Soal Praktikum set 2 5.1. Soal 1: Melakukan pembacaan text file dengan format yang telah ditentukan kemudian membuat menu unuk melihat data list barang , melihat barang dari harga termurah, melihat barang dari harga termahal dan membeli barang. Semua perhitungan menggunakan Thrust Library. 34
5.2. Soal 2: Melakukan pembacaan data melalui text file kemudian melakukan perkalian matrik dengan matrik itu sendiri yang telah ditransformasi modulo dan dilakukan
transpose matrik. Perkalian
matrik menggunakan CUBLAS level 3. 6. Set Dosen 6.1. Set Dosen 1. 6.2. Set Dosen 2. 7. Daftar Pustaka.
3.2.7. Pedoman praktikum topik 6 : Nvidia CUDA Optimization. Melalui pedoman praktikum topik 6 diharapkan mahasiswa menguasai konsep : • Optimasi Nvidia CUDA secara software . • Optimasi Nvidia CUDA dengan menggunakan coalescing memory. • Memaksimalkan penggunaan thread untuk mengoptimalkan unjuk kerja Nvidia CUDA.
Pemilihan topik pedoman ini didasarkan : • Perlunya optimasi Nvidia CUDA secara software untuk meningkatkan unjuk kerja Nvidia CUDA dalam memproses data secara paralel. Perancangan pedoman praktikum: 1. Tujuan. 2. Dasar Teori 2.1. Coalescing memory. 2.2. Optimasi thread. 3. Contoh Program 3.1. Contoh program 1: Membandingkan unjuk kerja coalescing memory dengan uncoalescing memory. 3.2. Contoh program 2: Membandingkan waktu eksekusi Nvidia CUDA dengan konfigurasi block dan thread yang berbeda. 35
4. Soal Praktikum set 1 4.1. Soal 1: Membuat program untuk melakukan transpose matrik dengan uncoalescing memory, coalescing memory, shared memory dan optimasi thread kemudian membandingkan waktu eksekusinya. 4.2. Soal 2: Analisa potongan program kernel, menentukan apakah potongan program kernel merupakan coalescing memory atau bukan. 5. Soal Praktikum set 2 5.1. Soal 1: Membuat program untuk melakukan transpose matrik dengan uncoalescing memory, coalescing memory, shared memory dan optimasi thread kemudian membandingkan waktu eksekusinya. Inputan berasal dari text file. 5.2. Soal 2: Analisa potongan program kernel, menentukan apakah potongan program kernel merupakan coalescing memory atau bukan. 6. Set Dosen 6.1. Set Dosen 1. 6.2. Set Dosen 2. 7. Daftar Pustaka.
3.2.8. Pedoman praktikum 7 : Tugas Rancang. Melalui pedoman praktikum topik 7 diharapkan mahasiswa dapat: • Menerapkan pedoman praktikum dari topik 1 sampai topik 6. • Membuat aplikasi CUDA untuk menyelesaikan permasalahan yang membutuhkan waktu eksekusi yang lama apabila dikerjakan oleh CPU. • Membandingkan unjuk kerja Nvidia CUDA dengan CPU. Tugas Rancang yang diberikan adalah menemukan jalur terpendek menggunakan algoritma Floyd-Warshall, mahasiswa akan diberikan algoritma Floyd-Warshall secara sequential kemudian mengubahnya ke algoritma paralel. Selain itu mahasiswa juga membandingkan unjuk kerja Nvidia CUDA dengan CPU dalam memproses data menggunakan algoritma Floyd-Warshall.
36
Spesifikasi tugas rancang : 1. Algoritma menemukan jalur terpendek menggunakan Floyd-Warshall. 2. Menggunakan shared memory dalam pemrosesan data. 3. Penghitungan waktu eksekusi GPU menggunakan timer CUDA event. 4. Penghitungan waktu eksekusi CPU menggunakan timer dari library ctime. 5. Data dibuat dalam bentuk matrik. 6. Ukuran matrik data minimal 50x50. 7. Transformasi data menggunakan Nvidia CUDA library.
3.3. Rancangan Pengujian Pengujian yang dilakukan melalui tahapan sebagai berikut: 3.3.1. Pengujian oleh dosen Dosen akan mengoreksi kekurangan di dalam pedoman dan akan memberikan masukan untuk memperbaiki pedoman.
3.3.2. Pengujian oleh mahasiswa a) Pengisian kuesioner Mahasiswa akan mengisi kuesioner tentang kualitas pedoman, dari hasil pengisian kuesioner akan dilakukan uji hipotesis. Pada akhir uji hipotesis akan disimpulkan hasil pengujian pedoman. b) Penilaian Mahasiswa akan mengerjakan soal-soal praktikum yang terdapat pada pedoman. Dari hasil penilaian dilakukan: Perhitungan nilai rata-rata. • Pedoman dinyatakan dapat membantu mahasiswa memahami materi apabila rata-rata penilaian di atas 70. Perhitungan perbedaan nilai kelompok. • Pedoman dinyatakan baik apabila
tidak ada perbedaan yang
signifikan antara kelompok yang terdiri dari mahasiswa yang memiliki ipk lebih dari sama dengan 3 dengan kelompok yang terdiri dari mahasiswa yang memiliki ipk kurang dari 3.
37
3.4. Pelaksanaan Pengujian Pedoman Pembelajaran dan Pedoman Praktikum. Tabel 3.1 menunjukkan pelaksanaan pengujian pedoman pembelajaran dan pedoman praktikum.
Tabel 3.1. Pelaksanaan Pengujian Pedoman Pembelajaran dan Pedoman Praktikum No.
1
Pelaksanaan Pengujian
Menjelaskan
kembali
konsep
Pedoman
Pedoman
Pembelajaran
Praktikum
topik 1
topik 0
topik 2
topik 1
topik 3
topik 2
pemrosesan paralel, menjelaskan konsep dasar Nvidia CUDA. Instalasi Nvidia Toolkit, Nvidia SDK, Nvidia Nsight dan integrasi dengan Visual Studio 2008 dan 2010. 2
Menjelaskan
konsep
dasar
programming model Nvidia CUDA yang meliputi thread, kernel dan memory
kemudian
melakukan
pengujian soal praktikum dengan standar nilai yang telah ditentukan. 3
Menjelaskan konsep dasar thread Nvidia
CUDA
pengajaran.
pada
pedoman
Kemudian
mecoba
contoh program pada pedoman praktikum. menerima dilakukan
Setelah semua
mahasiswa materi
baru
pengujian
soal
praktikum dengan standar nilai yang telah ditentukan. (Bersambung)
38
Tabel 3.1. Pelaksanaan Pengujian Pedoman Pembelajaran dan Pedoman Praktikum (Lanjutan) No.
4
Pelaksanaan Pengujian
Menjelaskan konsep dasar memory Nvidia
CUDA pada
Pedoman
Pedoman
Pembelajaran
Praktikum
topik 4
topik 3
topik 5
topik 4
topik 6
topik 5
pedoman
pengajaran. Kemudian mencoba contoh program pada pedoman praktikum. menerima
Setelah semua
dilakukan
mahasiswa materi
baru
pengujian
soal
praktikum dengan standar nilai yang telah ditentukan. 5
Menjelaskan dasar teori CUDA Compilation
kemudian
menjelaskan
jenis-jenis
CUDA
Runtime beserta contoh program pada pedoman praktikum. Setelah mahasiswa
menerima
semua
materi, baru dilakukan pengujian soal praktikum dengan standar nilai yang telah ditentukan. 6
Menjelaskan CUDA
contoh-contoh Library
beserta
keunggulannya mempermudah mempelajari
dalam user,
kemudian
CUBLAS
Library
dan Thrust Library secara spesifik. Setelah semua
mahasiswa materi,
baru
menerima dilakukan
pengujian soal praktikum. (Bersambung)
39
Tabel 3.1. Pelaksanaan Pengujian Pedoman Pembelajaran dan Pedoman Praktikum (Lanjutan) No.
7
Pelaksanaan Pengujian
Menjelaskan dasar teori tentang CUDA
Optimization
kemudian
mencoba contoh program pada pedoman
praktikum.
Setelah
mahasiswa menerima semua materi baru
dilakukan
pengujian
soal
praktikum dengan standar nilai yang telah ditentukan.
40
Pedoman
Pedoman
Pembelajaran
Praktikum
topik 7
topik 6
BAB IV PENGUJIAN PEDOMAN Pada bab ini akan dijelaskan mengenai pengujian pedoman beserta analisis hasil penilaian pedoman praktikum.
4.1. Pengujian oleh Dosen Setiap pedoman yang telah selesai dikerjakan, pedoman akan diperiksa oleh dosen pembimbing. Dosen pembimbing akan memberi masukan dan memeriksa kekurangan yang terdapat pada pedoman, masukan yang diterima yaitu: 1. Penambahan flowchart pada pedoman praktikum. 2. Penambahan penjelasan pada contoh program. 3. Perbaikan contoh program.
4.2. Pengujian oleh Mahasiswa Pengujian oleh mahasiswa dilakukan dengan 2 cara yaitu melalui kuesioner dan penilaian pedoman praktikum. Mahasiswa yang mengisi kuesioner adalah mahasiswa yang mengambil kelas dasar pemrograman sedangkan mahasiswa yang mengikuti pengujian penilaian pedoman praktikum adalah mahasiswa yang telah mengambil matakuliah pemrosesan paralel. Skala yang digunakan dalam pembuatan kuesioner adalah skala likert, skala likert terdiri dari lima pilihan yaitu sangat setuju, setuju, tidak pasti, tidak setuju dan sangat tidak setuju, pernyataan tersebut diberi skor 5, 4, 3, 2 dan 1. Tetapi pada skripsi ini hanya akan diberikan 4 pilihan yaitu sangat setuju, setuju, tidak setuju dan sangat tidak setuju, pilihan tidak pasti dihilangkan sehingga pernyataan diberi skor 4, 3, 2 dan 1. Kuesioner untuk pedoman pembelajaran dan pedoman praktikum masing-masing diberi pernyataan sejumlah 10 pernyataan. Tabel 4.1 merupakan pernyataan kuesioner pedoman pembelajaran.
41
Tabel 4.1. Pernyataan Kuesioner Pedoman Pembelajaran No. 1
Pernyataan Judul topik pedoman pembelajaran sesuai dengan tujuan dan materi pedoman pembelajaran.
2
Materi dalam pedoman pembelajaran sesuai dengan tujuan pedoman pembelajaran
3
Contoh-contoh program dalam pedoman dapat membantu memahami materi.
4
Kesesuaian contoh-contoh program dengan materi.
5
Kesesuaian ringkasan dengan dengan materi.
6
Ringkasan dapat membantu mahasiswa memahami materi.
7
Kesesuaian soal-soal latihan dengan materi.
8
Soal-soal latihan dapat menguji pemahaman mahasiswa tentang materi.
9
Pedoman pembelajaran dapat diajarkan sesuai dengan alokasi jam kuliah.
10
Format pedoman pembelajaran keseluruhan dapat membantu mahasiswa memahami materi.
42
Tabel 4.2 merupakan kuesioner untuk pedoman praktikum. Tabel 4.2. Pernyataan Kuesioner Pedoman Praktikum Pernyataan
No. 1
Materi dalam pedoman praktikum sudah sesuai dengan tujuan pedoman praktikum.
2
Dasar teori dalam pedoman praktikum dapat membantu memahami materi.
3
Kesesuaian dasar teori pedoman praktikum dengan materi pedoman kuliah.
4
Contoh-contoh program sesuai dengan dasar teori.
5
Contoh-contoh program dapat membantu memahami dasar teori.
6
Contoh-contoh program dapat membantu membuat aplikasi.
7
Kesesuaian soal praktikum dengan dasar teori.
8
Kesesuaian soal praktikum dengan contoh program.
9
Soal praktikum dapat selesai dikerjakan sesuai dengan alokasi jam praktikum.
10
Format pedoman praktikum keseluruhan dapat membantu mahasiswa memahami materi.
Setelah kuesioner diisi oleh mahasiswa maka dilakukan uji hipotesis untuk mendapatkan kesimpulan pengujian pedoman. 4.2.1. Uji Hipotesis Uji hipotesis bertujuan untuk mendapatkan kesimpulan pengujian pedoman berdasarkan data kuesioner.
43
Langkah-langkah uji hipotesis: 1. Menentukan Ho dan H1 Ho: responden tidak puas dengan kualitas pedoman H1: responden puas dengan kualitas pedoman 2. Responden dapat dikatakan puas dengan kualitas pedoman apabila mengisi kuesioner dengan skor lebih dari sama dengan 3. 3. Responden dapat dikatakan tidak puas dengan kualitas pedoman apabila mengisi kuesioner dengan skor kurang dari 3. 4. Menentukan α = 0.05 5. Menentukan derajat bebas (db) = (k-1)(b-1) =(5-1)(4-1) = 12 2
6. Menentukan nilai dari tabel x
Tabel x2 = x2(α, db) maka didapat 21.02
Tabel 4.3 merupakan tabel tingkat kesetujuan responden terhadap pedoman pembelajaran.
Tabel 4.3. Tingkat Kesetujuan Responden Pedoman Pembelajaran Pernyataan Sangat setuju Setuju Tidak setuju Sangat tidak setuju Total
Tujuan
Materi
Ringkasan
Soal
Format
Total
19 63 3 1
24 88 15 2
20 57 9 0
25 49 12 0
11 25 7 0
99 282 46 3
86
129
86
86
43
430
44
Tabel 4.4 merupakan tabel test Chi-Square pedoman pembelajaran. Tabel 4.4. Test Chi-Square Pedoman Pembelajaran Pernyataan
Fo
Fh
Fo-Fh
(Fo-Fh)2
(Fo-Fh)2 /Fh
Sangat setuju Sangat setuju Sangat setuju Sangat setuju Sangat setuju Setuju Setuju Setuju Setuju Setuju Tidak setuju Tidak setuju Tidak setuju Tidak setuju Tidak setuju Sangat tidak setuju Sangat tidak setuju Sangat tidak setuju Sangat tidak setuju Sangat tidak setuju Total
Tujuan Materi Ringkasan Soal Latihan Format Tujuan Materi Ringkasan Soal Latihan Format Tujuan Materi Ringkasan Soal Latihan Format Tujuan
19 24 20 25 11 26 88 57 49 25 3 15 9 12 7 1
19.8 29.7 19.8 19.8 9.9 56.4 84.6 56.4 56.4 28.2 9.2 13.8 9.2 9.2 4.6 0.6
-0.8 -5.7 0.2 5.2 1.1 6.6 3.4 0.6 -7.4 -3.2 -6.2 1.2 -0.2 2.8 2.4 0.4
0.64 32.49 0.04 27.04 1.21 43.56 11.56 1.36 54.76 10.24 38.44 1.44 0.04 7.84 5.76 0.16
0.0323 1.0939 0.002 1.3657 0.1222 0.7723 0.1366 0.0064 0.9709 0.3631 4.1783 0.1043 0.0043 0.8522 1.2522 0.2667
Materi
2
0.9
1.1
1.21
1.344
Ringkasan
0
0.6
-0.6
0.36
0.6
Soal Latihan
0
0.6
-0.6
0.36
0.6
Format
0
0.3
-0.3
0.09
0.3 14.368
Nilai Chi-Square dari Tabel 4.4 untuk tingkat signifikansi 5% dan derajat kebebasan 12 adalah 21.02 karena Ho < H1 maka Ho diterima dan H1 ditolak, dapat disimpulkan responden tidak puas dengan kualitas pedoman pembelajaran. Tabel 4.5 merupakan tabel tingkat kesetujuan responden pedoman praktikum.
45
Tabel 4.5. Tingkat Kesetujuan Responden Pedoman Praktikum Pernyataan
Sangat setuju Setuju Tidak setuju Sangat tidak setuju Total
Tujuan
Dasar
Contoh
Teori
Program
Soal
Format
Total
7 36 0 0
10 69 7 0
49 66 14 0
26 87 14 2
10 27 6 0
102 285 41 2
43
86
129
129
43
430
Tabel 4.6 merupakan tabel test Chi-Square pedoman praktikum. Tabel 4.6. Test Chi-Square Pedoman Praktikum Pernyataan
Fo
Fh
Fo-Fh
(Fo-Fh)2
(Fo-Fh)2 /Fh
Sangat setuju Sangat setuju Sangat setuju Sangat setuju Sangat setuju Setuju Setuju Setuju Setuju Setuju Tidak setuju Tidak setuju Tidak setuju Tidak setuju Tidak setuju Sangat tidak setuju Sangat tidak setuju Sangat tidak setuju Sangat tidak
Tujuan
7
10.2
-3.2
10.24
1.04
Dasar teori
10
20.4
-10.4
108.16
5.302
Contoh program
49
30.6
18.4
338.56
11.064
Soal praktikum
26
30.6
-4.6
21.16
0.6915
Format
10
10.2
-0.2
0.04
0.0039
Tujuan Dasar teori Contoh program Soal praktikum Format Tujuan Dasar teori Contoh program Soal praktikum Format Tujuan
36 69 66 87 27 0 7 14 14 6 0
28.5 57 85.5 85.5 28.5 4.1 8.2 12.3 12.3 4.1 0.2
7.5 12 -19.5 1.5 -1.5 -4.1 -1.2 1.7 1.7 1.9 -0.2
56.25 144 380.25 2.25 2.25 16.81 1.44 2.89 2.89 3.61 0.04
1.9737 2.5263 4.4474 0.0263 0.0789 4.1 0.1756 0.235 0.235 0.8805 0.2
Dasar teori
0
0.4
-0.4
0.16
0.4
Contoh program
0
0.6
-0.6
0.36
0.6
Soal praktikum
2
0.6
1.4
1.96
3.2667
46
setuju Sangat tidak setuju Total
Format
0
0.2
-0.2
0.04
0.2 37.411
Nilai Chi-Square dari Tabel 4.6 untuk tingkat signifikansi 5% dan derajat kebebasan 12 adalah 21.02 karena Ho > H1 maka Ho ditolak dan H1 diterima, dapat disimpulkan responden puas dengan kualitas pedoman praktikum. 4.2.2. Analisis Pengujian Pedoman Praktikum Hasil pengujian praktikum dapat dilihat pada tabel 4.7.
Tabel 4.7. Nilai Rata-Rata Pengujian
Praktikum
Jumlah Peserta 1 2 3 4 5 6
Nilai Rata-Rata 5 5 4 5 5 5
82 74 50 74 79 78
Dari 6 topik praktikum yang telah diujikan 5 topik pedoman praktikum telah melampaui standar penilaian rata-rata 70, tetapi ada satu topik praktikum yaitu topik praktikum nomor 3 yang gagal melampaui standar penilaian rata-rata karena soal yang diberikan terlalu sulit sehingga peserta tidak dapat mengerjakan soal praktikum. Selain itu jumlah peserta tidak bisa memenuhi standar jumlah peserta sesuai dengan spesifikasi dikarenakan sulitnya mencari peserta untuk mengikuti pengujian. Tabel 4.8 menunjukkan nilai rata-rata pengujian berdasarkan IPK. Kelompok 1 merupakan kelompok dengan ipk lebih dari sama dengan 3 sedangkan kelompok 2 merupakan kelompok dengan IPK kurang dari 3. Tabel 4.8. Nilai Rata-Rata Pengujian berdasarkan IPK Kelompok
Nilai rata-rata topik 1
Nilai rata-rata topik 3
Nilai rata-rata topik 2 47
Nilai rata-rata topik 4
Nilai rata-rata topik 5
Nilai rata-rata topik 6
1 2
100 70
86.6 55
50 50
80 70
85 75
80 76.6
Selanjutnya akan dilakukan uji t untuk mengetahui apakah ada perbedaan signifikan antara dua kelompok tersebut. 1. Langkah pertama adalah menentukan Ho dan H1. Ho = Tidak terdapat perbedaan nilai yang signifikan antara dua kelompok (µ1=µ2). H1 = Terdapat perbedaan nilai yang signifikan antara dua kelompok (µ1≠µ2). 2. Langkah kedua adalah menentukan kriteria penolakan. Df = (n1+n2)-2 = (3+2)-2 = 3. Maka di dapat nilai dari ttabel = 5.841dengan α=0.01 3. Langkah ketiga adalah menghitung nilai thitung dengan persamaan 2.3.
((1 − 1)$1 + ((2 − 1)$2 (1 + (2 , (1 + (2 − 2 (1(2
=
$
Dengan standar deviasi ∑(% − %̅ ) (
$=
Contoh menghitung uji t untuk penilaian topik 1 $1 = $2 = $
=
( --
--) !( -.
--)
=0
(/- /-) !(/- /-) !(/- /-) .
=0
(2 − 1)0 + (3 − 1)0 2 + 3 , 2+3−2 2.3 48
$
=0 = =
%1 − %2 $
100 − 70 0
= tidak terdefinisi Tabel 4.9 menunjukkan hasil uji t untuk semua topik pedoman praktikum Tabel 4.9. Hasil Uji t Pedoman Praktikum Topik
3453678
339:;<
Keterangan
Kesimpulan
5.841 5.841
t hitung < t tabel
Tidak ada perbedaan
1 2
tidak terdefinisi 4.52
3 4
tidak terdefinisi 3.87
9.925 5.841
t hitung < t tabel
Tidak ada perbedaan
5 6
tidak terdefinisi 0.79
5.841 5.841
t hitung < t tabel
Tidak ada perbedaan
Hasil uji t menunjukkan bahwa t
hitung
tabel
maka Ho diterima dan H1 ditolak
sehingga dapat disimpulkan tidak ada perbedaan signifikan antara 2 kelompok untuk topik praktikum nomor 2, 4 dan 6. Tetapi hasil uji t untuk topik praktikum nomor 1, 3 dan 5 tidak dapat disimpulkan karena hasil perhitungan standar deviasi adalah 0 sehingga hasil t hitung menjadi tidak bisa didefinisikan. Kesimpulan dari analisis pengujian pedoman praktikum adalah tidak ada perbedaan signifikan nilai antara kelompok dengan ipk lebih dari sama dengan 3 dengan ipk kurang dari 3. Walaupun melalui uji t masih ada 3 topik praktikum yang tidak terdefinisi sehingga 5 topik praktikum dinyatakan dapat membantu mahasiswa dalam memahami materi sedangkan 1 topik praktikum masih memerlukan perbaikan agar dapat lolos dari standar penilaian rata-rata.
49
BAB V KESIMPULAN DAN SARAN Pada bab ini akan dijabarkan kesimpulan yang didapatkan selama pengerjaaan skripsi yang meliputi perancangan, realisasi dan pengujian. Pada bab ini juga akan dipaparkan beberapa saran yang dapat dipertimbangkan untuk pengembangan skripsi ini. 5.1. Kesimpulan Berdasarkan perancangan, realisasi dan pengujian skripsi didapat beberapa kesimpulan yaitu: 1. Pemrosesan paralel dapat dilakukan pada Nvidia CUDA dan pemrosesan paralel pada Nvidia CUDA dapat meningkatkan kinerja CPU dalam pemrosesan data. 2. Hasil uji hipotesis untuk pedoman pembelajaran dan pedoman praktikum yaitu responden tidak puas terhadap pedoman pembelajaran dan responden puas terhadap pedoman praktikum. 3. Pengujian 6 topik pedoman praktikum melampaui standar rata-rata 70 untuk 5 topik praktikum, sedangkan 1 topik masih membutuhkan perbaikan untuk melampaui standar rata-rata. 5.2. Saran Pengembangan Beberapa saran yang penulis dapat berikan untuk pengembangan skripsi ini adalah sebagai berikut:
1. Meningkatkan jumlah peserta untuk pengujian agar mendapatkan hasil uji beda yang valid. 2. Menambah fasilitas VGA Nvidia CUDA di BB-5 agar praktikum dapat berjalan dengan lancar dan jumlah peserta praktikum dapat ditambah. 3. Membuat aplikasi CUDA dengan menggabungkan beberapa unit PC (Cluster CUDA) dapat dilakukan untuk pengembangan aplikasi Nvidia CUDA. 4. Apabila ingin melakukan uji validalitas dan uji reabilitas sebaiknya dengan cara membandingkan dua kelompok mahasiswa yang diberi perlakuan berbeda.
50
DAFTAR PUSTAKA
[1]
Benchmark reviews, “ Nvidia GF 100 GPU Fermi Architecture”, diakses dalam http://benchmarkreviews.com/index.php?option=com_content&task=view&id =518&Itemid=72&limit=1&limitstart=2 pada 22 Agustus 2012.
[2]
Farber, Rob, ”CUDA Application Design And Development”, MK, 2011.
[3]
Guru 3D, “Geforce GTX 680 review”, diakses dalam http://www.guru3d.com/articles_pages/geforce_gtx_680_review,3.html pada 22 Agustus 2012.
[4]
Harvard, “Harvard recognized by NVIDIA as a CUDA Center of Excellence”, dikases dalam http://www.seas.harvard.edu/news-events/newsarchive/2009/harvard-recognized-by-nvidia-as-a-cuda-center-of-excellencenvidia/?searchterm=None pada 10 April 2012.
[5]
Kirk David B., Hwu Wen-mei W., ”Programming Massively Parallel Processors”, MK, 2010.
[6]
Mardalis, “Metode Penelitian Suatu Pendekatan Proposal”, PT. Bumi Aksara, Jakarta, 2008.
[7]
Nvidia, “NVIDIA CUDA Architecture Introduction & Overview”, Nvidia, 2009.
[8]
Nvidia, ”NVIDIA CUDA C Programming Guide version 4”, Nvidia, 2011.
[9]
Nvidia, ” NVIDIA CUDA C Programming Guide version 4.2”, Nvidia, 2012.
[10] Nvidia, ”NVIDIA Fermi Compute Architecture”, Nvidia. [11] Nvidia, “Tokyo Institute of Technology Selected as Japan’s First CUDA Center of Excellence”, diakses dalam http://www.nvidia.com/object/io_1270716327243 pada 10 April 2012. [12] Sanders J., Kandrot E., “CUDA By Example”, Addison-Wesley, 2010. [13] Susanto, Stephanus H., “Sistem Operasi Android pada Kuliah Mobile Programming”, FTEK-UKSW, Salatiga, 2011. (Laporan Tugas Akhir untuk meraih gelar sarjana di FTEK-UKSW).
51
[14] Tom’s hardware, “Nvidia Powering World’s Most Poweful Supercomputer”, diakses dalam http://www.tomshardware.com/news/Tesla-Titan-Cray-XK6Supercomputer-Steve-Scott,13669 pada 10 April 2012. [15] Wikipedia, “CUDA”, diakses dalam http://en.wikipedia.org/wiki/CUDA pada 8 Februari 2012. [16] Wikipedia, ”GDDR5”, diakses dalam http://en.wikipedia.org/wiki/GDDR5 pada 1 Mei 2012. [17] Wikipedia, ” Graphics processing unit”, diakses dalam http://en.wikipedia.org/wiki/Graphics_processing_unit pada 1 Mei 2012.
52
LAMPIRAN A PEDOMAN PEMBELAJARAN TOPIK 2 NVIDIA CUDA PROGRAMMING MODEL 1. Tujuan Melalui pedoman pembelajaran topik 2 diharapkan mahasiswa dapat: • Memahami dan mengerti arsitektur Nvidia CUDA. • Menguasai dasar-dasar programming model Nvidia CUDA.
2. Materi 2.1. Arsitektur Nvidia CUDA Nvidia CUDA mempunyai arsitektur berbasis General Purpose Computing Graphics Processing Units (GPGPU). Nvidia CUDA dapat berjalan dengan platform prosesor berbasis x86 dan x64 dari platform manapun dan chipset motherboard yang mendukung PCI-Express Bus. Gambar 1 merupakan gambar arsitektur komputer dengan CUDA.
Gambar 1. Arsitektur komputer dengan CUDA Pada gambar diatas CUDA dihubungkan melalui interface PCI-Express Bus yang mendukung bandwidth sampai 16 GB/s. PCI-Express dihubungkan oleh Northbridge pada motherboard yang berfungsi menjembatani CPU, DRAM , Southbridge dan PCI-Express. 53
Nvidia CUDA terdiri dari beberapa komponen yaitu [4, h.2]: 5. Unit pemrosesan di dalam GPU. 6. OS kernel-level yang mendukung inisialisasi dan konfigurasi hardware. 7. User mode driver yang menunjukkan device level API. 8. PTX instruction set architecture (ISA) untuk komputasi paralel kernel dan fungsi.
Saat ini Nvidia CUDA
telah mencapai 4 generasi arsitektur. Tabel 1
menunjukkan generasi arsitektur Nvidia CUDA .
Tabel 1. Generasi Arsitektur Nvidia CUDA [2]
Arsitektur
G80/G92
Jumlah
SP/
Jumlah SP
Shared memory/
Maksimum
maksimu
Dimensi
maksimum
CUDA
maksimum
cache
memory
m thread
block
SM
cores
bandwidth
per block
16
8
86,4 GB/s
512
512 x 512 x 64
141,7 GB/s
512
512x 512 x 64
177,4 GB/s
1024
1024 x 1024 x 64
192,4 GB/s
1024
1024 x 1024 x 64
128
246 KB shared
memory, 512 KB register GT 200
30
8
240
480KB
shared
memory,
1920
KB register Fermi
16
32
512
L1 cache 1024 KB, L2 cache 768 KB, 2048 KB register
Kepler
8
192
1536
L1 cache
512
KB, L2 cache 1536 KB, 2048 KB register
GPU pada Nvidia CUDA terdiri beberapa unit Streaming Multiprocessors (SM) seperti yang terlihat pada Gambar 2.
54
Gambar 2. GPU pada Nvidia CUDA [1]
Setiap generasi arsitektur Nvidia CUDA mempunyai SM yang berbeda-beda. Gambar 3 adalah gambar SM pada arsitektur Fermi.
Gambar 3. SM pada arsitektur Fermi [6, h.8]
Bagian yang terdapat pada setiap SM arsitektur Fermi adalah [6, h.5-8]: • 32 CUDA core yang terdiri dari Arithmetic Logic Unit (ALU) dan Floating Point Unit (FPU). ALU pada arsitektur Fermi mendukung
55
sampai 32 bit precision unit untuk semua instruksi sedangkan FPU mendukung sampai 64 bit. • 16 unit Load Store unit yang digunakan untuk pengalamatan resource dan destination yang dapat menampung sampai 16 threads per clock. • 4 Special Function Unit (SFU) yang berfungsi mengeksekusi instruksi transcendental seperti sin,cos dan bilangan kuadrat. • 2 Warp scheduler untuk penjadwalan 32 thread yang bekerja secara paralel. • 64 KB Shared memory dan L1 cache. • Register serbaguna dengan ukuran 32.768 x 32 bit.
2.2. Programming Model Nvidia CUDA 2.2.1. Kernel [5, h.7] Kernel merupakan fungsi yang dipanggil dari program host dan akan dieksekusi didalam GPU. Sewaktu kernel dipanggil, kernel akan mengeksekusi thread secara paralel. Kernel didefinisikan didalam deklarasi __global__ .
Format penulisan kernel adalah sebagai berikut:
namakernel<<< jumlah block, jumlah threads dalam satu block >>> (Nilai yang dipassingkan); Contoh pemanggilan fungsi kernel:
1. __global__ void tambah(float* A, float* B, float* C) 2. { 3. int i = threadIdx.x; 4. C[i] = A[i] + B[i]; 5. } 6. int main() 7. { 8. ... 9. tambah<<1, N>>>(A, B, C); 10. ... 11. }
56
Penjelasan program: Pertama kali program akan berjalan dari fungsi main pada baris ke 6. Kemudian program akan memanggil fungsi kernel tambah pada baris ke 9. Program akan mengeksekusi kernel pada fungsi dengan deklarasi __global__ pada baris ke 1. Pada baris ke 3 Thread akan disimpan didalam array i dengan tipe integer (1 dimensi). Kemudian pada baris ke 4 vektor A akan ditambah dengan vektor B dan hasilnya akan disimpan dalam vektor C.
2.2.2. Thread [3, h.60] Thread adalah sekumpulan instruksi yang dapat dijadwalkan. Susunan thread akan mengeksekusi kernel secara bergantian. Setiap thread mempunyai ID yang disebut threadIdx, begitu juga dengan block mempunyai ID yang disebut blockIdx. ThreadIdx dan BlockIdx dapat dinyatakan dalam 3 dimensi vektor yaitu x, y dan z. Contoh: ThreadIdx.x , ThreadIdx.y, ThreadIdx.z, BlockIdx.x, BlockIdx.y,BlockIdx.z
2.2.3. Heterogenous Programming [5, h.13] Pada pemrograman Nvidia CUDA, program dieksekusi di dua device yang terpisah. Prosesor sebagai host sequential
berfungsi mengeksekusi program C secara
dan GPU sebagai device berfungsi mengeksekusi thread secara
paralel. Host dan device
juga mempunyai memory masing-masing yang
terpisah. Gambar 4 menjelaskan tentang heterogenous programming:
57
Gambar 4. Heterogeneous programming Program berjalan dari host code secara serial. Isi dari kode host ini adalah inisialisasi memory dan penyalinan memory dari CPU ke GPU. Setelah itu kernel akan dipanggil dan kode paralel akan dieksekusi didalam device. Satu kernel akan dieksekusi pada satu waktu sehingga hanya ada satu instruksi paralel yang dikerjakan pada satu waktu. Setelah kode paralel selesai dikerjakan oleh device, hasilnya dikembalikan pada host dan apabila masih ada instruksi paralel yang harus dikerjakan maka kernel dapat dipanggil lagi dan dieksekusi didalam device.
58
2.2.4. Management Memory Nvidia CUDA[3, h.46-50] Management memory Nvidia CUDA : •
CPU dan GPU mempunyai memory yang terpisah
•
CPU memanajemen GPU memory: 1. Mengalokasi dan membebaskan memory 1.1. cudaMalloc(void ** pointer, size_t nbytes) •
memesan lokasi dan kapasitas memory
1.2. cudaMemset(void *pointer, int value, size_t count) •
mengeset kapasitas dan isi memory
1.3. cudaFree(void * pointer) •
membebaskan memory
2. Menyalin data ke devices 2.1. cudaMemcpy(void *dst, void *src, size_t nbytes, cudaMemcpy direction) •
menyalin memory
2.2. cuda Memcpy direction: •
cudaMemcpyHostToDevice o menyalin memory dari host ke device
•
cudaMemcpyDeviceToHost o menyalin memory dari device ke host
•
cudaMemcpyDeviceToDevice o menyalin memory dari device ke device
3. Menggunakan Global memory
59
Contoh program manajemen memory pada Nvidia CUDA: 1. __global__ void VecAdd(float* A, float* B, float* C, int N) 2. { 3. int i = blockDim.x * blockIdx.x + threadIdx.x; 4. if (i < N) 5. C[i] = A[i] + B[i]; 6. } 7. int main() 8. { 9. int N = ..; 10. size_t size = N * sizeof(float); 11. float* h_A = (float*)malloc(size); 12. float* h_B = (float*)malloc(size); 13. float* d_A; 14. cudaMalloc(&d_A, size); 15. float* d_B; 16. cudaMalloc(&d_B, size); 17. float* d_C; 18. cudaMalloc(&d_C, size); 19. cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); 20. cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); 21. int threadsPerBlock = 256; 22. int blocksPerGrid = (N + threadsPerBlock – 1) / threadsPerBlock; 23. VecAdd<<
>>(d_A, d_B, d_C, N); 24. cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); 25. cudaFree(d_A); 26. cudaFree(d_B); 27. cudaFree(d_C); 28. }
Penjelasan program: Program dimulai dari baris ke 7 saat pemanggilan fungsi main. Pada baris ke 10 program memesan memory dengan tipe data float sebanyak N dikali 4 byte kemudian pada baris ke 11 program mengalokasikan vektor A dan B didalam host memory. Baris ke 13 sampai 17 adalah alokasi vektor didalam device memory dengan menggunakan cudaMalloc. Baris ke 19 sampai 20 program menyalin vektor dari host memory ke device memory. Baris ke 21 adalah inisialisasi jumlah thread didalam block sebanyak 256. Baris ke 22 adalah menentukan jumlah block yang dipakai misalnya: N data yang kita pakai adalah 60
1500 maka jika kita masukkan ke dalam persamaan pada baris ke 22 adalah ( 1500 + 256 – 1) / 256 = 6,85 karena tipe data yang dipakai adalah integer dan jumlah block yang dipakai adalah bilangan bulat maka jumlah block yang dipakai adalah 6. Karena jumlah block yang dipakai adalah 6 maka total thread yang dapat dipakai adalah 1536 didapat dari perkalian block dengan thread. Selanjutnya pada baris ke 23 adalah pemanggilan kernel VecAdd dan program akan meloncat ke baris nomor 1. Baris ke 3 program akan memberikan ID yang berbeda
pada setiap threadId.x, kemudian pada baris ke 4 program akan
mengecek ID dari thread. Jika ID dari thread kurang dari N maka thread yang mempunyai ID kurang dari N yang akan mengeksekusi penjumlahan vektor pada baris ke 5. Setelah kernel selesai dieksekusi program akan kembali pada baris ke 24, vektor C akan disalin dari device ke host kemudian baris ke 25 sampai ke 27 adalah membebaskan memory untuk vektor A, B dan C.
3. Ringkasan
1. GPU pada Nvidia CUDA merupakan gabungan dari
beberapa Streaming
Multiprocessors(SM), SM merupakan gabungan dari beberapa CUDA core. 2. Kernel merupakan fungsi yang dipanggil dari program host dan akan dieksekusi di dalam GPU. 3. Program pada Nvidia CUDA dieksekusi di 2 devices yang berbeda yaitu CPU dan GPU. 4. Execution model pada Nvidia CUDA menggunakan 2 devices yang terpisah yang mempunyai memory masing-masing, CPU sebagai host mengontrol manajemen memory pada GPU.
61
4. Soal-Soal Latihan
a) Jelaskan fungsi utama kernel! b) Apa yang terjadi jika cudaMalloc() tidak dideklarasikan sebelum menyalin memory dari CPU ke GPU? c) Jelaskan perbedaan kode serial dan kode paralel! d) Apa yang akan terjadi apabila jumlah thread yang dideklarasikan pada kernel melebihi dari jumlah data yang akan dieksekusi, jelaskan! e) Jelaskan perbedaan kode device dengan kode host!
5. Daftar pustaka 1. Benchmark reviews, “ Nvidia GF 100 GPU Fermi Architecture”, diakses dalam http://benchmarkreviews.com/index.php?option=com_content&task=view&id=51 8&Itemid=72&limit=1&limitstart=2 pada 22 Agustus 2012. 2. Guru 3D, “Geforce GTX 680 review”, diakses dalam http://www.guru3d.com/articles_pages/geforce_gtx_680_review,3.html pada 22 Agustus 2012. 3. Kirk David B., Hwu Wen-mei W., ”Programming Massively Parallel Processors”, MK, 2010. 4. Nvidia, “NVIDIA CUDA Architecture Introduction & Overview”, Nvidia, 2009. 5. Nvidia, ” Nvidia CUDA C Programming Guide version 4.2”, Nvidia.2011. 6. Nvidia, ”NVIDIA Fermi Compute Architecture”, Nvidia.
62
LAMPIRAN B PEDOMAN PEMBELAJARAN TOPIK 3 NVIDIA CUDA THREADING
1. Tujuan Melalui pedoman pengajaran topik 3 diharapkan mahasiswa dapat: • Memahami dan mengerti konsep dasar dari thread. • Menguasai dasar-dasar pemrograman thread. • Mengerti manfaat dari penjadwalan thread.
2. Materi 2.1. Konsep Thread pada Nvidia CUDA Thread adalah sekumpulan instruksi yang dapat dijadwalkan. Pada setiap Streaming Multiprocessor (SM) GPU Nvidia CUDA, jumlah thread yang dapat ditampung adalah 1536 untuk arsitektur Fermi, sedangkan pada arsitektur G80 dan GT200 jumlah maksimal thread yang dapat ditampung pada setiap SM adalah 768. Thread digunakan untuk mengeksekusi kernel, setiap susunan thread akan mengeksekusi kernel secara bergantian. Thread mempunyai unique ID yang digunakan sebagai alamat dan kontrol, Susunan thread mengeksekusi kernel dapat dilihat pada gambar 1.
Gambar 1. Susunan thread mengeksekusi kernel [3, h.5] 63
Karena jumlah thread yang digunakan untuk eksekusi sangat banyak, maka Nvidia CUDA disebut Single Instruction Multiple Thread (SIMT). SIMT merupakan gabungan dari Single Instruction Multiple Data (SIMD) dan Single Program Multiple Data (SPMD). Tabel 1 menjelaskan definisi dari SIMD dan SPMD.
Tabel 1. Definisi SIMD dan SPMD [2, h.61-62] Nama
Keterangan
Single Instruction Multiple
Setiap thread mempunyai kontrol yang
Data (SIMD)
sama tetapi dapat menangani banyak data sekaligus
Single Program Multiple Data (SPMD)
Terdapat
banyak
data
yang
harus
ditangani oleh satu program
Pada Tabel 1 dapat terlihat bahwa konsep dari SIMT yang merupakan gabungan dari SIMD dan SPMD adalah setiap thread akan menjalankan instruksi yang sama tetapi ada banyak data yang bisa ditangani. Hirarki dari thread dapat dilihat pada Gambar 2. Thread
Block
Block merupakan gabungan dari beberapa thread
64
Grid
Grid merupakan gabungan dari beberapa block
Gambar 2. Hirarki dari Thread [3, h.13]
Thread hanya dapat bekerjasama dengan thread yang ada didalam satu block yang sama. Thread di dalam satu block yang sama dapat berbagi resource dan dapat mengetahui apa yang sedang dikerjakan oleh thread lain di dalam satu block yang sama. Thread dalam CUDA dapat saling sinkronisasi dalam satu block dengan menggunakan fungsi syncthreads().
Tujuan pemanggilan fungsi
syncthreads
adalah untuk memastikan semua thread dalam block telah menyelesaikan eksekusi sebelum mengerjakan instruksi selanjutnya. Untuk execution model pada thread dapat dilihat pada Gambar 3.
65
Gambar 3. Execution model [3, h.13]
Gambar 3 menunjukkan execution model pada Nvidia CUDA. o
Satu prosesor akan mengeksekusi satu thread.
o
Gabungan dari beberapa prosesor (multiprosesor) mengeksekusi satu Block .
o
Gabungan dari beberapa multiprosesor (Device) mengeksekusi Grid.
66
2.2.Pemrograman Thread [2, h.8-10] Setiap thread mempunyai ID yang disebut threadIdx, begitu juga dengan block mempunyai ID yang disebut blockIdx. ThreadIdx dan blockIdx dapat dinyatakan dalam 3 dimensi vektor yaitu x, y dan z. Contoh: ThreadIdx.x , ThreadIdx.y, ThreadIdx.z, BlockIdx.x, BlockIdx.y,BlockIdx.z Contoh program:
1. __global__ void tambah(float A[N][N], float B[N][N], float C[N][N]) 2. { 3. int i = threadIdx.x; 4. int j = threadIdx.y; 5. C[i][j] = A[i][j] + B[i][j]; 6. } 7. int main() 8. { 9. ... 10. tambah<<1,4,4>>>(A, B, C); 11. ... 12. }
Penjelasan program: Program diatas merupakan modifikasi dari program
yang pertama.
Perbedaannya adalah menggunakan 2 dimensi vektor untuk eksekusinya yaitu pada baris ke 3 dan baris ke 4. Thread akan disimpan pada array i dan array j, sehingga dapat diilustrasikan seperti pada Gambar 4.
67
Gambar 4. Block 2 dimensi vektor
Pada baris ke 10 fungsi kernel memanggil 1 block yang berisi 8 thread. Kemudian Program akan mengeksekusi fungsi tambah pada deklarasi __global__ pada baris ke 1. Pada baris ke 3 thread ID vektor x akan disimpan pada array i dan pada baris ke 4 thread ID vektor y akan disimpan pada array j. Kemudian pada baris ke 5 vektor A akan ditambah dengan vektor B secara 2 dimensi dan hasilnya akan disimpan pada vektor C.
Pada gambar diatas program akan mengeksekusi thread pada baris pertama terlebih dahulu baru kemudian setelah thread pada baris pertama selesai, thread pada baris kedua akan dieksekusi.
Thread pada sebuah block dapat juga digabungkan dengan thread pada block lain contohnya ada pada Gambar 5.
68
Gambar 5. Grid [3, h.39]
Terdapat sebuah grid yang memiliki 3 block didalamnya masing-masing block mempunyai 5 buah thread Yang mempunyai thread ID 0, 1, 2, 3 dan 4 apabila semua thread ingin digabungkan untuk mengeksekusi kernel maka caranya adalah dengan instruksi sebagai berikut:
blockIdx.x *blockDim.x + threadIdx.x
misal: blockIdx.x (0) * blockDim.x (5) + threadIdx.x (0) = 0 blockIdx.x (0) * blockDim.x (5) + threadIdx.x (1) = 1 blockIdx.x (1) * blockDim.x (5) + threadIdx.x (0) = 5 blockIdx.x (2) * blockDim.x (5) + threadIdx.x (4) = 14
blockidx.x adalah ID dari block. blockDim.x adalah banyaknya thread di dalam block. threadIdx.x adalah ID dari thread. Maka thread ID nya menjadi adalah 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 Maka terdapat 15 thread dengan thread ID yang berbeda-beda.
69
2.3. Penjadwalan Thread [1, h.71-74] Penjadwalan thread adalah bagian yang penting dan harus diperhatikan dalam pemrograman Nvidia CUDA . Penjadwalan thread dalam Nvidia CUDA disebut dengan warp. Pada setiap warp berisi 32 unit thread yang hanya bisa dijadwalkan didalam 1 Streaming Multiprocessor (SM). Semisal terdapat 256 thread didalam sebuah block maka jumlah warp yang bisa dijadwalkan yaitu 256/32 = 8 warp. Jika pada setiap SM terdapat 3 block maka jumlah warp menjadi 8 *3 = 24 warp yang bisa dijadwalkan. Pada arsitektur Fermi maksimal thread pada setiap SM adalah 1024 maka jumlah warp yang bisa dijadwalkan adalah 1024/32 = 32 warp pada setiap SM arsitektur Fermi. Contoh penjadwalan warp dapat dilihat pada gambar 6.
Gambar 6. Penjadwalan Warp [4, h.10] Gambar 6 merupakan gambar penjadwalan warp. Pada penjadwalan pertama thread block 1 melakukan eksekusi tetapi intruksi nomor 7 memerlukan waktu yang lama untuk menunggu hasil dari instruksi sebelumnya kemudian thread block 1 akan masuk dalam waiting mode. Kemudian thread block 2 warp 1 akan memulai mengerjakan instruksi tetapi karena instruksi nomor 3 membutuhkan waktu yang lama maka thread block 3 akan mengambil alih. Setelah instruksi nomor 1 dan 2 selesai namun instruksi nomor 3 thread block 2 belum selesai. Maka thread block 3 kembali mengerjakan instruksi pada warp 2 namun kembali intruksi nomor 3 pada thread block 3 memerlukan waktu untuk menunggu hasil maka thread block 2 mengambil alih karena instruksi nomor 3 siap dikerjakan. 70
Setelah thread block 2 selesai mengerjakan, thread block 1 mengambil alih karena instruksi nomor 7 siap untuk dieksekusi. Setelah thread block 1 selesai maka kembali melakukan penjadwalan lagi pada thread block 1 tetapi dengan instruksi yang berbeda, setelah selesai thread block 3 kembali mengambil alih untuk mengerjakan instruksi nomor 3 dan 4. Dengan penjadwalan warp yang cukup maka hardware akan bekerja maksimal karena hardware akan selalu mengeksekusi intruksi sehingga selalu tetap sibuk. 2.4. Transparent Scalability [1, h.68-69] Kemampuan untuk menyelesaikan suatu aplikasi dengan perbedaan kecepatan dalam eksekusinya disebut dengan transparent scalability. Setiap block dapat dieksekusi dalam urutan apapun tergantung dari susunan block. Transparent scalability ditunjukkan oleh Gambar 7.
Gambar 7. Transparent scalability Pada gambar diatas block yang akan dieksekusi pertama merupakan block pada baris pertama terlebih dahulu. Device sebelah kanan akan lebih cepat mengeksekusi karena dapat mengeksekusi 4 block sekaligus sedangkan pada
71
device sebelah kiri akan lebih lama karena block yang dieksekusi hanya 2 block pada satu waktu. 3. Ringkasan a) Thread berfungsi untuk mengeksekusi kernel. b) Nvidia CUDA menggunakan SIMT yang merupakan gabungan dari SIMD dan SPMD. c) Gabungan dari beberapa thread membentuk block, gabungan dari beberapa block membentuk grid. d) Thread dapat dinyatakan dalam 3 dimensi vektor. e) Warp digunakan untuk menjadwalkan thread. f) Transparent scalability adalah kemampuan untuk menyelesaikan suatu aplikasi dengan perbedaan kecepatan dalam eksekusinya.
4. Soal - Soal Latihan a) Jelaskan mengapa Nvidia CUDA bisa disebut SIMT! b) Terdapat 7 SM pada Nvidia GTX 460 (Fermi) setiap SM mempunyai 48 core hitunglah berapa jumlah warp yang bisa dijadwalkan! c) Analisalah program dibawah ini ! 1. __global__ void kernel(int a[n][n], int b[n][n], int c[n][n]) 2. { 3. int i = blockIdx.x *blockDim.x + threadIdx.x 4. int j = blockIdx.y *blockDim.y+ threadIdx.y 5. if (i < N && j < N) 6. C[i][j] = A[i][j] + B[i][j]; 7. } 8. int main() 9. { 10. N= 1000; 11. ... 12. kernel<<1,2,8>>>(A, B, C); 13. ... 14.}
• • •
Gambarlah bentuk block dan thread dari program diatas! Apa isi dari int C pada program diatas? Manakah konfigurasi block yang lebih cepat kernel<<2,4,4>> atau kernel<<1,16>>?? 72
5. Daftar Pustaka 1. Kirk David B., Hwu Wen-mei W., ”Programming Massively Parallel Processors”, MK, 2010. 2. Nvidia, ”Nvidia CUDA C Programming Guide version 4”, Nvidia, 2011. 3. Nvidia, “Getting Started With CUDA, “Nvidia, 2008. 4. Kirk David B., Hwu Wen-mei W., ”CUDA threads”, Illinois, 2007.
73
LAMPIRAN C PEDOMAN PRAKTIKUM TOPIK 1 NVIDIA CUDA BASIC 1. Tujuan Melalui pedoman praktikum topik 1 diharapkan mahasiswa menguasai konsep: • Pembuatan program Nvidia CUDA sederhana. • Penyalinan memory, eksekusi kernel dan menampilkan hasil eksekusi program Nvidia CUDA sederhana. • Eksekusi thread pada kernel.
2. Dasar teori Nvidia CUDA merupakan mesin pemrosesan paralel yang menggunakan CPU dan GPU
untuk melakukan eksekusi program secara paralel. Nvidia CUDA
menggunakan CPU untuk mengeksekusi kode host secara sequential dan GPU untuk mengeksekusi kode device secara paralel. Gambar 1 menunjukkan diagram alir eksekusi Nvidia CUDA.
Gambar 1. Diagram Alir Eksekusi Nvidia CUDA [5]
74
Penjelasannya adalah sebagai berikut [5]: 5. Alur kerja yang pertama yaitu CPU akan menyalin semua data yang dibutuhkan GPU untuk eksekusi dari main memory ke GPU memory onboard. 6. Setelah semua data yang dibutuhkan GPU untuk proses eksekusi disalin. CPU akan mengerjakan kode sequential dan kemudian akan memberikan instruksi yang akan dikerjakan oleh GPU. 7. GPU akan mengeksekusi secara paralel. 8. Hasil dari pemrosesan paralel akan disalin ke main memory.
2.1. Kernel [2, h.42] Kernel merupakan fungsi yang dipanggil dari program host dan dieksekusi didalam GPU. Format pemanggilan kernel: Namakernel <<< jumlah block, jumlah threads dalam satu block >>> (Nilai yang dipassingkan); 2.2. Thread [2, h.60] Thread adalah sekumpulan instruksi yang dapat dijadwalkan. Gabungan dari thread disebut dengan block, gabungan dari beberapa block disebut dengan grid. Susunan thread akan mengeksekusi kernel secara bergantian. Thread mempunyai ID yang digunakan sebagai alamat dan kontrol, ilustrasi dapat dilihat pada Gambar 2.
. Gambar 2. Susunan Thread mengeksekusi Kernel [3, h.5] 75
2.3. CUDA Template CUDA
template
adalah
kode
minimum
yang
digunakan
untuk
mengeksekusi GPU. CUDA template terdiri dari kode host dan kode device. Potongan kode di bawah adalah kode template untuk aplikasi sederhana:
1. 2. 3. 4. 5. 6.
#include <stdio.h> #include <cuda.h> __global__ void square_array(float * a, int N) { }
7. int main(void) 8. { 9. //Inisialisasi memory CPU 10. float *a_h, *a_d; 11. const int N = 10; 12. size_t size = N * sizeof(float); 13. a_h = (float *)malloc(size); 14. //Inisialisasi memory GPU 15. cudaMalloc((void **) &a_d, size); 16. //Penyalinan data dari CPU ke GPU 17. cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice); 18. //pemanggilan kernel 19. kernel<<< jumlah block, jumlah thread >>> (a_d, N); 20. //Penyalinan data dari GPU ke CPU 21. cudaMemcpy(a_h,a_d,sizeof(float)*NcudaMemcpyDevi ceTo Host); 22. }
2.4. Alokasi Memory[2, h.46-50] CUDA mempunyai alokasi memory yang berbeda antara CPU dengan GPU. Beberapa contoh alokasi memory pada Nvidia CUDA adalah sebagai berikut: 4.
Mengalokasi dan membebaskan memory 1.4. cudaMalloc(void ** pointer, size_t nbytes) 76
• memesan lokasi dan kapasitas memory 1.5. cudaMemset(void *pointer, int value, size_t count) • mengeset kapasitas dan isi memory 1.6. cudaFree(void * pointer) • membebaskan memory 5.
Menyalin data ke devices •
cudaMemcpy(void *dst, void *src, size_t nbytes, cudaMemcpy direction) • menyalin memory
o
cuda Memcpy direction: •
cudaMemcpyHostToDevice o menyalin memory dari host ke device
•
cudaMemcpyDeviceToHost o menyalin memory dari device ke host
•
cudaMemcpyDeviceToDevice o menyalin memory dari device ke device
3. Contoh Program 3.1. Contoh program 1 [4]: Contoh program 1 merupakan contoh program untuk mengkuadratkan isi thread. Kode Sumber: 1. #include <stdio.h> 2. #include <cuda.h>
3. __global__ void kernelku( float *a) 4. { 5. int idx = threadIdx.x; 6. a[idx]= idx * idx ; 7. __syncthreads(); 8. } 9. int main(void) 10. { 77
11. cudaError_t cudaStatus; 12. float *a_h, *a_d; 13. const int N = 10; 14. size_t size = N * sizeof(float); 15. a_h = (float *)malloc(size); 16. cudaStatus = cudaMalloc((void**)&a_d, size); 17. if (cudaStatus != cudaSuccess) 18. { 19. printf( "alokasi gagal\n"); 20. } 21. else 22. { 23. printf("alokasi berhasil!\n"); 24. } 25. for (int i=0; i>> (a_d); 28. cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); 29. for (int i=0; i<10; i++) printf("%f\n",a_h[i]); 30. free(a_h); 31. cudaFree(a_d); 32. system("pause"); 33. }
Gambar 3. Printscreen contoh program 2
78
Gambar 4. Flowchart program kernel contoh program 1 (baris ke 3 sampai baris ke 8)
Gambar 5. Flowchart main program contoh program 2 (baris ke 9 sampai baris ke 33) 79
Penjelasan program per baris : 1. Inisialisasi pustaka <stdio.h>. 2. Inisialisasi pustaka <cuda.h>. 3. Inisialisasi program kernel dengan nama kernelku dengan passing parameter float a . 4. Kurung buka tanda awal program kernel. 5. Inisialisasi integer idx sebagai threadIdx.x dengan vektor sumbu x. 6. Mengalikan idx dengan dirinya sendiri kemudian disimpan pada array a. Pada proses ini thread akan mengalikan dirinya sendiri secara paralel dalam satu waktu. 7. __syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya. 8. Kurung tutup tanda akhir program kernel. 9. Deklarasi program utama main, program akan dimulai dari main program. 10. Kurung buka program. 11. Inisialisasi cudaerror_t dengan nama cudaStatus, cudaerror_t berfungsi menampilkan status error. 12. Deklarasi pointer a_h dan a_d. 13. Deklarasi constant integer N dengan nilai 10. 14. Mengisi size dengan tipe data float sejumlah 10. 15. Alokasi a_h pada memory CPU dengan tipe data float sejumlah 10. 16. Inisialisasi a_d pada GPU memory apabila inisialisasi sukses maka akan mendapatkan status cudaSuccess pada cudaStatus. 17. Jika cudaStatus mendapat status cudaSuccess maka alokasi memory berhasil sebaliknya jika cudaStatus mendapat status mendapat status selain cudaSuccess maka alokasi memory gagal. 18. Kurung buka. 19. Menampilkan tampilan alokasi memory gagal. 20. Kurung tutup. 21. Else. 22. Kurung buka. 80
23. Menampilkan tampilan alokasi memory berhasil. 24. Kurung tutup. 25. Mengisi array di a_h sejumlah 10 dengan tipe data float. 26. Menyalin data dari memory CPU ke memory GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a_h disalin ke a_d dengan arah pengiriman host to device. 27. Memanggil kernel dengan nama kernelku dengan jumlah block 1 dan jumlah thread 50 . 28. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a_d disalin ke a_h dengan arah pengiriman device to host. 29. Mencetak isi a_h. 30. Membebaskan alokasi memory a_h. 31. Membebaskan alokasi memory a_d. 32. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 33. Kurung tutup tanda program berakhir.
81
4. Soal Praktikum set 1 4.1. Isilah array dengan data [10,12,14,16,18,20,22,24,26,28] kemudian ubahlah program pada contoh diatas seperti printscreen menjadi seperti gambar 6.
Gambar 6. Printscreen soal praktikum 1
4.2. Isilah array dengan data [10,12,14,16,18,20,22,24] kemudian jumlahkan semua array seperti printscreen pada gambar 7.
Gambar 7. Printscreen soal praktikum 2
82
5. Soal Praktikum Set 2 5.1. Isilah array dengan data [10,12,14,16,18,20,22,24,26,28] kemudian ubahlah program pada contoh diatas seperti printscreen menjadi seperti gambar 6.
Gambar 6. Printscreen soal praktikum 1
5.2. Isilah array dengan data [10,12,14,16,18,20,22,24] kemudian kalikan semua array seperti printscreen pada gambar 7.
Gambar 7. Printscreen soal praktikum 2
83
6. Set Dosen 6.1. Set Dosen 1 Soal 4.1. Kode Sumber: 1. #include <stdio.h> 2. #include <cuda.h> 3. __global__ void square_array(float *bb) 4. { 5. int idx = threadIdx.x; 6. bb [idx]= bb [idx % 4]; 7. __syncthreads(); 8. } 9. int main(void) 10. { 11. float *a_h,*bb; 12. const int N = 10; 13. float x[10]={10,12,14,16,18,20,22,24,26,28}; 14. size_t size = N * sizeof(float); 15. a_h = (float *)malloc(size); 16. cudaMalloc((void **) &bb, size); 17. for (int i=0; i>> (bb); 20. cudaMemcpy(a_h,bb,sizeof(float)*N,cudaMemcpyDevi ceToHost); 21. for (int i=0; i<10; i++) printf("%f\n",a_h[i]); 22. free(a_h); 23. cudaFree(bb); 24. system("pause"); 25. }
Penjelasan program per baris: 1. inisialisasi pustaka <stdio.h>. 2. inisialisasi pustaka <cuda.h>. 3. Inisialisasi program kernel dengan nama square_array dengan passing parameter float bb . 4. Kurung buka tanda awal program kernel. 5. Inisialisasi integer idx sebagai threadIdx.x dengan vektor sumbu x. 6. Pembatasan eksekusi thread hanya 4 thread, thread yang berjalan hanya thread ID 0 sampai 3. 84
7. __syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya. 8. Kurung tutup tanda akhir program kernel. 9. Deklarasi program utama main, program akan dimulai dari main program. 10. Kurung buka program. 11. Deklarasi pointer a_h dan bb. 12. Deklarasi constant integer N dengan nilai 10. 13. Mengisi float x dengan data [10,12,14,16,18,20,22,24,26,28]. 14. Mengisi size dengan tipe data float sejumlah 10. 15. Alokasi a_h pada memory CPU dengan tipe data float sejumlah 10. 16. Alokasi bb pada GPU memory. 17. Mengisi array di a_h sejumlah 10 dengan tipe data float. 18. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari x disalin ke bb dengan arah pengiriman host to device. 19. Memanggil kernel dengan nama kernelku dengan jumlah block 2 dan jumlah thread 100 . 20. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari bb disalin ke a_h dengan arah pengiriman device to host. 21. Mencetak isi a_h. 22. Membebaskan alokasi memory a_h. 23. Membebaskan alokasi memory bb. 24. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 25. Kurung tutup tanda berakhirnya program.
Soal 4.2. Kode Sumber: 1. #include <stdio.h> 2. #include <cuda.h>
85
3. __global__ void square_array(float *bb) 4. { 5. int boss= threadIdx.x; 6. bb[boss] = bb[boss%4] + bb[boss+4]; 7. __syncthreads(); 8. bb[boss] = bb[boss%2] + bb[boss+2]; 9. __syncthreads(); 10. bb[boss] = bb[boss%1] + bb[boss+1]; 11. __syncthreads(); 12. } 13. int main(void) 14. { 15. float *a_h,*bb; 16. const int N = 8; 17. float x[8]={10,12,14,16,18,20,22,24}; 18. size_t size = N * sizeof(float); 19. a_h = (float *)malloc(size); 20. cudaMalloc((void **) &bb, size); 21. for (int i=0; i>> (bb); 24. cudaMemcpy(a_h,bb,sizeof(float)*N, cudaMemcpyDeviceToHost); 25. printf("10+12+14+16+18+20+22+24= %f\n",a_h[0]); 26. free(a_h); 27. cudaFree(bb); 28. system("pause"); 29. }
Penjelasan program per baris : 1.
inisialisasi pustaka <stdio.h>.
2.
inisialisasi pustaka <cuda.h>.
3.
Inisialisasi program kernel dengan nama square_array dengan passing parameter float bb .
4.
Kurung buka tanda awal program kernel.
5.
Inisialisasi integer boss sebagai threadIdx.x dengan vektor sumbu x.
6.
Memecah array menjadi dua bagian kemudian dijumlahkan, hasil penjumlahan disimpan pada array bb alamat 0 sampai 3.
7.
__syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya.
8.
Memecah array menjadi 4 bagian , penjumlahan dilakukan pada pecahan pertama dan kedua. Hasil penjumlahan disimpan pada array bb alamat 0 sampai 1. 86
9.
__syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya.
10. Melakukan penjumlahan array pada alamat 0 dan 1 dan hasilnya disimpan pada array 0. 11. __syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya. 12. Kurung tutup tanda akhir program kernel. 13. Deklarasi program utama main, program akan dimulai dari main program. 14. Kurung buka program. 15. Deklarasi pointer a_h dan bb. 16. Deklarasi constant integer N dengan nilai 8. 17. Mengisi float x dengan data [10,12,14,16,18,20,22,24]. 18. Mengisi size dengan tipe data float sejumlah 8. 19. Alokasi a_h pada memory CPU dengan tipe data float sejumlah 8. 20. Alokasi bb pada GPU memory. 21. Mengisi array di a_h sejumlah 8 dengan tipe data float. 22. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari x disalin ke bb dengan arah pengiriman host to device. 23. Memanggil kernel dengan nama kernelku dengan jumlah block 1 dan jumlah thread 10 . 24. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari bb disalin ke a_h dengan arah pengiriman device to host. 25. Mencetak isi a_h. 26. Membebaskan alokasi memory a_h. 27. Membebaskan alokasi memory bb. 28. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 29. Kurung tutup tanda berakhirnya program.
87
6.2. Set Dosen 2
Soal 5.1. Kode Sumber: 1. #include <stdio.h> 2. #include <cuda.h> 3. __global__ void square_array(float *bb, int N) 4. { 5. int idx = threadIdx.x; 6. bb [idx]= bb [ 9 - idx % 4]; 7. __syncthreads(); 8. } 9. int main(void) 10. { 11. float *a_h, *bb; 12. const int N = 10; 13. float x[10]={10,12,14,16,18,20,22,24,26,28}; 14. size_t size = N * sizeof(float); 15. a_h = (float *)malloc(size); 16. cudaMalloc((void **) &bb, size); 17. for (int i=0; i>> (bb, N); 20. cudaMemcpy(a_h,bb,sizeof(float)*N,cudaMemcpyDevi ceToHost); 21. for (int i=0; i<10; i++) printf("%f\n",a_h[i]); 22. free(a_h); 23. cudaFree(bb); 24. system("pause"); 25. }
Penjelasan program per baris : 1. inisialisasi pustaka <stdio.h>. 2. inisialisasi pustaka <cuda.h>. 3. Inisialisasi program kernel dengan nama square_array dengan passing parameter float bb . 4. Kurung buka tanda awal program kernel. 5. Inisialisasi integer idx sebagai threadIdx.x dengan vektor sumbu x. 6. Pembatasan eksekusi thread hanya 4 thread, thread yang berjalan hanya thread ID 4 sampai 7. 88
7. __syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya. 8. Kurung tutup tanda akhir program kernel. 9. Deklarasi program utama main, program akan dimulai dari main program. 10. Kurung buka program. 11. Deklarasi pointer a_h dan bb. 12. Deklarasi constant integer N dengan nilai 10. 13. Mengisi float x dengan data [10,12,14,16,18,20,22,24,26,28]. 14. Mengisi size dengan tipe data float sejumlah 10. 15. Alokasi a_h pada memory CPU dengan tipe data float sejumlah 10. 16. Alokasi bb pada GPU memory. 17. Mengisi array di a_h sejumlah 10 dengan tipe data float. 18. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari x disalin ke bb dengan arah pengiriman host to device. 19. Memanggil kernel dengan nama kernelku dengan jumlah block 2 dan jumlah thread 100 . 20. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari bb disalin ke a_h dengan arah pengiriman device to host. 21. Mencetak isi a_h. 22. Membebaskan alokasi memory a_h. 23. Membebaskan alokasi memory bb. 24. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 25. Kurung tutup tanda berakhirnya program.
Soal 5.2. Kode Sumber: 1. #include <stdio.h> 2. #include <cuda.h> 3. __global__ void square_array(float *bb, float *a, int N) 89
4. { 5. 6. 7. 8. 9. 10. 11. 12. }
int boss= threadIdx.x; bb[boss] = bb[boss%4] * bb[boss+4]; __syncthreads(); bb[boss] = bb[boss%2] * bb[boss+2]; __syncthreads(); bb[boss] = bb[boss%1] * bb[boss+1]; __syncthreads();
13. int main(void) 14. { 15. float *a_h,*bb; 16. const int N = 10; 17. float x[8]={10,12,14,16,18,20,22,24}; 18. size_t size = N * sizeof(float); 19. a_h = (float *)malloc(size); 20. cudaMalloc((void **) &bb, size); 21. for (int i=0; i>> (bb); 24. cudaMemcpy(a_h,bb,sizeof(float)*N,cudaMemcpyDevi ceToHost); 25. printf("10*12*14*16*18*20*22*24 = %f\n",a_h[0]); 26. free(a_h); 27. cudaFree(a_d); 28. system("pause"); 29. }
Penjelasan program per baris: 1.
inisialisasi pustaka <stdio.h>.
2.
inisialisasi pustaka <cuda.h>.
3.
Inisialisasi program kernel dengan nama square_array dengan passing parameter float bb .
4.
Kurung buka tanda awal program kernel.
5.
Inisialisasi integer boss sebagai threadIdx.x dengan vektor sumbu x.
6.
Memecah array menjadi dua bagian kemudian dikalikan, hasil perkalian disimpan pada array bb alamat 0 sampai 3.
7.
__syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya.
8.
Memecah array menjadi 4 bagian , perkalian dilakukan pada pecahan pertama dan kedua. Hasil perkalian disimpan pada array bb alamat 0 sampai 1.
90
9.
__syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya.
10. Melakukan perkalian array pada alamat 0 dan 1 dan hasilnya disimpan pada array 0. 11. __syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya. 12. Kurung tutup tanda akhir program kernel. 13. Deklarasi program utama main, program akan dimulai dari main program. 14. Kurung buka program. 15. Deklarasi pointer a_h dan bb. 16. Deklarasi constant integer N dengan nilai 8. 17. Mengisi float x dengan data [10,12,14,16,18,20,22,24]. 18. Mengisi size dengan tipe data float sejumlah 8. 19. Alokasi a_h pada memory CPU dengan tipe data float sejumlah 8. 20. Alokasi bb pada GPU memory. 21. Mengisi array di a_h sejumlah 8 dengan tipe data float. 22. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari x disalin ke bb dengan arah pengiriman host to device. 23. Memanggil kernel dengan nama kernelku dengan jumlah block 1 dan jumlah thread 10 . 24. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari bb disalin ke a_h dengan arah pengiriman device to host. 25. Mencetak isi a_h. 26. Membebaskan alokasi memory a_h. 27. Membebaskan alokasi memory bb. 28. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 29. Kurung tutup tanda berakhirnya program.
91
7. Daftar Pustaka 1. Farber.R, ” CUDA Application Design And Development”, MK, 2011. 2. Kirk David B., Hwu Wen-mei W., ”Programming Massively Parallel Processors”, MK, 2010. 3. Nvidia, “Getting Started With CUDA”, Nvidia, 2008. 4. Parallel Panorama, “My first CUDA program”, diakses dalam http://llpanorama.wordpress.com/2008/05/21/my-first-cuda-program/ pada 3 Agustus 2012. 5. Wikipedia, “CUDA”, diakses dalam http://en.wikipedia.org/wiki/CUDA pada 8 Februari 2012.
92
LAMPIRAN D PEDOMAN PRAKTIKUM TOPIK 2 NVIDIA CUDA THREADING 1. Tujuan Melalui pedoman praktikum topik 2 diharapkan mahasiswa menguasai konsep: • Menggunakan thread dalam 3 sumbu vektor x,y dan z. • Menggabungkan semua thread di dalam block. • Menggunakan timer untuk mengetahui waktu eksekusi GPU.
2. Dasar Teori 2.1 Konfigurasi thread [3, h.8] Setiap thread mempunyai ID yang disebut threadIdx, begitu juga dengan block mempunyai ID yang disebut blockIdx. ThreadIdx dan BlockIdx dapat dinyatakan dalam 3 dimensi vektor yaitu x, y dan z.
Thread pada sebuah block dapat juga digabungkan dengan thread pada block lain contohnya ada pada Gambar 1 .
Gambar 1. Grid [4, h.39]
Terdapat sebuah grid yang memiliki 3 block di dalamnya masing-masing block mempunyai 5 buah thread Yang mempunyai thread ID 0, 1, 2, 3 dan 4 apabila
93
semua thread ingin digabungkan untuk mengeksekusi kernel maka caranya adalah dengan konfigurasi sebagai berikut:
blockIdx.x *blockDim.x + threadIdx.x
misal: blockIdx.x (0) * blockDim.x (5) + threadIdx.x (0) = 0 blockIdx.x (0) * blockDim.x (5) + threadIdx.x (1) = 1 blockIdx.x (1) * blockDim.x (5) + threadIdx.x (0) = 5 blockIdx.x (2) * blockDim.x (5) + threadIdx.x (4) = 14
blockidx.x adalah ID dari block. blockDim.x adalah banyaknya thread di dalam block. threadIdx.x adalah ID dari thread. Maka thread ID nya menjadi adalah 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 Maka terdapat 15 thread dengan thread ID yang berbeda beda.
3. Contoh Program 3.1. Contoh Program 1 Contoh program 1 adalah contoh program eksekusi kernel dengan thread vektor x, y dan z [5] . Kode Sumber: 1. #include <stdio.h> 2. #include <cuda.h> 3. #include <cutil.h> 4. __global__ void kernelku(float *a) 5. { 6. int x=threadIdx.x; 7. int y=threadIdx.y; 8. int z=threadIdx.z; 9. a[x]=0; 10. __syncthreads(); 11. } 12. int main(void) 94
13. { 14. 15. 16. 17. 18. 19. 20. 21. 22. 23. 24. 25. 26. 27. 28. }
float *a_h, *a_d; const int N = 10; size_t size = N * sizeof(float); a_h = (float *)malloc(size); cudaMalloc((void **) &a_d, size); for (int i=0; i>> (a_d); cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); for (int i=0; i<10; i++) printf("%f\n",a_h[i]); free(a_h); cudaFree(a_d); system("pause");
Gambar 2. Printscreen sumbu x contoh program 1
Gambar 3. Printscreen sumbu y contoh program 1
95
Gambar 4. Printscreen sumbu z contoh program 1
Start
Inisialisasi threadIdx.x, threadIdx.y dan threadIdx.z
Thread mengeksekusi data
Sinkronisasi thread
End
gambar 4. Flowchart program kernel contoh program 1 (baris 4 sampai baris 11)
96
Gambar 5. Flowchart main program contoh program 1 (baris 12 sampai baris 28)
Penjelasan program per baris: 1. inisialisasi pustaka <stdio.h>. 2. inisialisasi pustaka <cuda.h>. 3. inisialisasi pustaka <cutil.h>. 4. Inisialisasi program kernel dengan nama kernelku dengan passing parameter float a . 5. Kurung buka tanda awal program kernel. 6. Inisialisasi integer x sebagai threadIdx.x dengan vektor sumbu x. 97
7. Inisialisasi integer y sebagai threadIdx.y dengan vektor sumbu y. 8. Inisialisasi integer z sebagai threadIdx.z dengan vektor sumbu z. 9. Array a yang dihandle oleh threadIdx.x akan memberikan nilai 0 sebanyak jumlah threadIdx.x. 10. __syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya. 11. Kurung tutup tanda akhir program kernel. 12. Deklarasi program utama main, program akan dimulai dari main program. 13. Kurung buka program. 14. Deklarasi pointer a_h dan a_d. 15. Deklarasi constant integer N dengan nilai 10. 16. Mengisi size dengan tipe data float sejumlah 10. 17. Alokasi a_h pada memory CPU dengan tipe data float sejumlah 10. 18. Inisialisasi a_d pada GPU memory. 19. Mengisi array di a_h sejumlah 10 dengan tipe data float. 20. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a_h disalin ke a_d dengan arah pengiriman host to device. 21. Inisialisasi thread dengan sumbu vektor x,y dan z. 22. Memanggil kernel dengan nama kernelku dengan jumlah block 2 dan konfigurasi thread (5,4,2) yang artinya jumlah thread 5 pada sumbu x, jumlah thread 4 pada sumbu y dan jumlah thread 2 pada sumbu z. 23. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a_d disalin ke a_h dengan arah pengiriman device to host. 24. Mencetak isi a_h. 25. Membebaskan alokasi memory a_h. 26. Membebaskan alokasi memory a_d. 27. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 28. Kurung tutup tanda berakhirnya program.
98
3.2. Contoh program 2: Contoh program 2 adalah contoh program penggabungan thread antar block [5].
Kode Sumber: 1. #include <stdio.h> 2. #include <cuda.h> 3. #include <cutil.h> 4. __global__ void kernelku(float *a) 5. { 6. int gap=blockIdx.x *blockDim.x + threadIdx.x; 7. a[gap]=0; 8. __syncthreads(); 9. } 10. int main(void) 11. { 12. float *a_h, *a_d; 13. const int N = 20; 14. size_t size = N * sizeof(float); 15. a_h = (float *)malloc(size); 16. cudaMalloc((void **) &a_d, size); 17. for (int i=0; i>> (a_d); 21. cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); 22. for (int i=0; i<20; i++) printf("%f\n",a_h[i]); 23. free(a_h); 24. cudaFree(a_d); 25. system("pause"); 26. }
99
Gambar 6. Printscreen contoh program 2 Apabila jumlah block diganti dengan dua maka printscreen program menjadi seperti berikut
Gambar 7. Printscreen contoh program 2
100
Gambar 8. Flowchart program kernel contoh program 2 (baris 4 sampai baris 9)
101
Gambar 9. Flowchart main program contoh program 2 (baris 10 sampai baris 26)
Penjelasan program per baris : 1.
Baris pertama merupakan inisialisasi pustaka <stdio.h>.
2.
Baris kedua merupakan inisialisasi pustaka <cuda.h>.
3.
Baris ketiga merupakan inisialisasi pustaka <cutil.h>.
4.
Inisialisasi program kernel dengan nama kernelku dengan passing parameter float a .
5.
Kurung buka tanda awal program kernel.
6.
Inisialisasi integer gap dengan gabungan thread antar block.
7.
Array a yang dihandle oleh gap akan memberikan nilai 0 sebanyak jumlah gabungan thread antar block.
102
8.
__syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya.
9.
Kurung tutup tanda akhir program kernel.
10. Deklarasi program utama main, program akan dimulai dari main program. 11. Kurung buka program. 12. Deklarasi pointer a_h dan a_d. 13. Deklarasi constant integer N dengan nilai 20. 14. Mengisi size dengan tipe data float sejumlah 20. 15. Alokasi a_h pada memory CPU dengan tipe data float sejumlah 20. 16. Inisialisasi a_d pada GPU memory. 17. Mengisi array di a_h sejumlah 20 dengan tipe data float. 18. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a_h disalin ke a_d dengan arah pengiriman host to device. 19. Inisialisasi thread dengan sumbu vektor x,y dan z. 20. Memanggil kernel dengan nama kernelku dengan jumlah block 1 dan konfigurasi thread (5,4,2) yang artinya jumlah thread 5 pada sumbu x, jumlah thread 4 pada sumbu y dan jumlah thread 2 pada sumbu z. 21. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a_d disalin ke a_h dengan arah pengiriman device to host. 22. Mencetak isi a_h. 23. Membebaskan alokasi memory a_h. 24. Membebaskan alokasi memory a_d. 25. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 26. Kurung tutup tanda berakhirnya program.
3.3. Contoh program 3: Timer pada Nvidia CUDA digunakan untuk menghitung waktu yang dibutuhkan GPU untuk mengeksekusi kernel. Pada saat memanggil timer, library cutil.h harus dipanggil. Contoh program 3 adalah contoh program penggunaan timer pada Nvidia CUDA [5]. 103
Kode Sumber: 1. #include <stdio.h> 2. #include <cuda.h> 3. #include <cutil.h> 4. __global__ void kernelku(float *a) 5. { 6. int gap=blockIdx.x *blockDim.x + threadIdx.x; 7. a[gap]=gap*gap; 8. __syncthreads(); 9. }
10. int main(void) 11. { 12. unsigned int timer=0; 13. cutCreateTimer( &timer ); 14. float *a_h, *a_d; 15. const int N = 1<<25; 16. size_t size = N * sizeof(float); 17. a_h = (float *)malloc(size); 18. cudaMalloc((void **) &a_d, size); 19. for (int i=0; i>> (a_d); 24. cudaThreadSynchronize(); 25. cudaThreadSynchronize(); 26. cutStopTimer( timer ); 27. cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); 28. printf("CUDA execution time = %f ms\n",cutGetTimerValue( timer )); 29. for (int i=0; i<20; i++) printf("%f\n",a_h[i]); 30. free(a_h); 31. cudaFree(a_d); 32. system("pause"); 33. }
104
Gambar 10. Printscreen contoh program 3
Gambar 11. Flowchart program kernel contoh program 3 (baris 4 sampai baris 9)
105
Gambar 12. Flowchart main program contoh program 3 (baris 10 sampai baris 33)
Penjelasan program per baris : 1. Baris pertama merupakan inisialisasi pustaka <stdio.h>. 2. Baris kedua merupakan inisialisasi pustaka <cuda.h>. 3. Baris ketiga merupakan inisialisasi pustaka <cutil.h>. 4. Inisialisasi program kernel dengan nama kernelku dengan passing parameter float a . 5. Kurung buka tanda awal program kernel. 106
6. Inisialisasi integer gap dengan gabungan thread antar block. 7. Array a yang dihandle oleh gap akan memberikan nilai 0 sebanyak jumlah gabungan thread antar block. 8. __syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya. 9. Kurung tutup tanda akhir program kernel. 10. Deklarasi program utama main, program akan dimulai dari main program. 11. Kurung buka program. 12. Inisialisasi timer dengan nilai awal 0. 13. Pemanggilan fungsi cutCreateTimer pada library <cutil.h> dengan menunjuk pointer timer. 14. Deklarasi pointer a_h dan a_d. 15. Deklarasi constant integer N dengan nilai 2 pangkat 25. 16. Mengisi size dengan tipe data float sejumlah N. 17. Alokasi a_h pada memory CPU dengan tipe data float sejumlah N. 18. Inisialisasi a_d pada GPU memory. 19. Mengisi array di a_h sejumlah 20 dengan tipe data float. 20. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a_h disalin ke a_d dengan arah pengiriman host to device. 21. Inisialisasi thread dengan sumbu vektor x,y dan z. 22. Timer pada GPU mulai dinyalakan. 23. Memanggil kernel dengan nama kernelku dengan jumlah block 10 dan konfigurasi thread (64,4,2) yang artinya jumlah thread 64 pada sumbu x, jumlah thread 4 pada sumbu y dan jumlah thread 2 pada sumbu z. 24. Fungsi cudaThreadSynchronize() bertujuan untuk memastikan semua thread selesai dan hasilnya siap untuk dikirimkan ke CPU. 25. Fungsi cudaThreadSynchronize() bertujuan untuk memastikan semua thread selesai dan hasilnya siap untuk dikirimkan ke CPU. 26. Timer dimatikan. 27. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a_d disalin ke a_h dengan arah pengiriman device to host. 107
28. Mencetak waktu eksekusi GPU. 29. Mencetak isi a_h. 30. Membebaskan alokasi memory a_h. 31. Membebaskan alokasi memory a_d. 32. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 33. Kurung tutup tanda berakhirnya program.
108
4. Soal Praktikum Set 1 4.1. Buatlah program penjumlahan koordinat seperti pada gambar 13. Gunakanlah penjumlahan vektor x dan y . Pada gambar 2 konfigurasi thread adalah (6,6).
Gambar 13.
Pada gambar 13 jumlah thread yang digunakan adalah 36 sehingga semua koordinat akan tereksekusi apabila konfigurasi thread yang digunakan adalah (6,2) maka hasilnya akan terlihat seperti pada gambar 14, perlihatkan juga waktu eksekusi GPUnya.
Gambar 14.
4.2. Perkalian matrik Ubahlah kode sequential perkalian matrik berikut ini menjadi kode paralel yang dapat dikerjakan oleh GPU seperti printscreen gambar 15.
109
Kode sequential perkalian matriks [1]: 1. Void perkalianmatrik(float *A,float *B,float *C) 2. for (int i=0; i< lebarmatrik; ++i) 3. for (int j=0; j< lebarmatrik; ++j) 4. { 5. float sum=0; 6. for( int k =0; k < lebarmatrik; ++k) 7. { 8. float matriksA= A[i * lebarmatrik +k]; 9. float matriksB= B[k * lebarmatrik +j]; 10. sum +=matriksA * matriksB; 11. } 12. C[j * lebarmatrik + i] = sum; 13. } 14. }
Gambar 15.
110
5. Soal Praktikum Set 2 5.1. Buatlah program perkalian koordinat seperti pada gambar 13. Gunakanlah perkalian vektor x dan y . Pada gambar 13 konfigurasi thread adalah (6,6).
Gambar 13.
Pada gambar 13 jumlah thread yang digunakan adalah 36 sehingga semua koordinat akan tereksekusi apabila konfigurasi thread yang digunakan adalah (6,4) maka hasilnya akan terlihat seperti pada gambar 14, perlihatkan juga waktu eksekusi GPUnya.
Gambar 14.
111
5.2. Perkalian matrik Ubahlah kode sequential perkalian matrik berikut ini menjadi kode paralel yang dapat dikerjakan oleh GPU seperti printscreen gambar 15.
Kode sequential perkalian matriks[1]: 1. Void perkalianmatrik(float *A,float *B,float *C) 2. for (int i=0; i< lebarmatrik; ++i) 3. for (int j=0; j< lebarmatrik; ++j) 4. { 5. float sum=0; 6. for( int k =0; k < lebarmatrik; ++k) 7. { 8. float matriksA= A[i * lebarmatrik +k]; 9. float matriksB= B[k * lebarmatrik +j]; 10. sum +=matriksA * matriksB; 11. } 12. C[j * lebarmatrik + i] = sum; 13. } 14. }
Gambar 15.
112
6. Set Dosen 6.1.Set Dosen 1 Soal 4.1. Kode Sumber [6]: 1. #include <stdio.h> 2. #include <cuda.h> 3. #include <cutil.h> 4. #define COLUMNS 6 5. #define ROWS 6 6.__global__ void add(int *a, int *b, int *c) 7. { 8. int x = threadIdx.x; 9. int y = threadIdx.y; 10. int i = (COLUMNS*y) + x; 11. c[i] = a[i] + b[i]; 12. __syncthreads(); 13. } 14. int main(void) 15. { 16. unsigned int timer=0; 17. cutCreateTimer( &timer ); 18. int a[ROWS][COLUMNS], b[ROWS][COLUMNS], c[ROWS][COLUMNS]; 19. int *dev_a, *dev_b, *dev_c; 20. cudaMalloc((void **) &dev_a, ROWS*COLUMNS*sizeof(int)); 21. cudaMalloc((void **) &dev_b, ROWS*COLUMNS*sizeof(int)); 22. cudaMalloc((void **) &dev_c, ROWS*COLUMNS*sizeof(int)); 23. for (int y = 0; y < ROWS; y++) 24. for (int x = 0; x < COLUMNS; x++) 25. { 26. a[y][x] = x; 27. b[y][x] = y; 28. } 29. cudaMemcpy(dev_a, a, ROWS*COLUMNS*sizeof(int), cudaMemcpyHostToDevice); 30. cudaMemcpy(dev_b, b, ROWS*COLUMNS*sizeof(int), cudaMemcpyHostToDevice); 31. dim3 grid(COLUMNS,ROWS); 32. dim3 threadperblock(6,4); 33. cutStartTimer( timer ); 34. add<<>>(dev_a, dev_b, dev_c); 35. cudaThreadSynchronize(); 36. cudaThreadSynchronize(); 113
37. cutStopTimer( timer ); 38. cudaMemcpy(c, dev_c, ROWS*COLUMNS*sizeof(int), cudaMemcpyDeviceToHost); 39. for (int y = 0; y < ROWS; y++) 40. { 41. for (int x = 0; x < COLUMNS; x++) 42. { 43. printf("[%d][%d]=%d ",y,x,c[y][x]); 44. } 45. printf("\n"); 46. } 47. printf("CUDA execution time = %f ms\n",cutGetTimerValue( timer )); 48. cudaFree(dev_a); 49. cudaFree(dev_b); 50. cudaFree(dev_c); 51. system("pause"); 52. }
Penjelasan program per baris: 1. inisialisasi pustaka <stdio.h>. 2. inisialisasi pustaka <cuda.h>. 3. inisialisasi pustaka <cutil.h>. 4. Mendefinisikan kolom dengan lebar 6. 5. Mendefinisikan baris dengan tinggi 6. 6. Inisialisasi program kernel dengan nama kernelku dengan passing parameter int a, b dan c. 7. Kurung buka tanda awal program kernel. 8. Inisialisasi integer x sebagai threadIdx.x dengan vektor sumbu x. 9. Inisialisasi integer y sebagai threadIdx.x dengan vektor sumbu y. 10. Inisialisasi integer i sebagai perkalian kolom dengan int y ditambah dengan int x, threadIdx.y akan mulai mengeksekusi setelah threadIdx.x selesai dieksekusi. 11. Merupakan penambahan float a dengan float b dengan menggunakan thread i sebagai eksekutornya setelah thread i dengan thread ID 0 sampai 5 selesai mengeksekusi maka otomatis dilanjutkan dengan thread ID 6 sampai 11 begitu seterusnya sampai batas threadIdx.y. 12. __syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya.
114
13. Kurung tutup tanda berakhirnya program kernel. 14. Deklarasi program utama main, program akan dimulai dari main program. 15. Kurung buka program. 16. Inisialisasi timer dengan nilai awal 0. 17. Pemanggilan fungsi cutCreateTimer pada library <cutil.h> dengan menunjuk pointer timer. 18. Inisialisasi int a,b dan c dengan array 2 dimensi. 19. Inisialiasasi pointer dev_a,dev_b dan dev_c. 20. Inisialisasi dev_a pada GPU memory. 21. Inisialisasi dev_b pada GPU memory. 22. Inisialisasi dev_c pada GPU memory. 23. Perulangan baris. 24. Perulangan kolom. 25. Pengisian kolom dan baris pada alamat array int a dan b beserta nilainya, nilai int a akan diisi dengan nilai kolom sedangkan int b akan diisi dengan nilai baris. 29. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a disalin ke dev_a dengan arah pengiriman host to device. 30. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari b disalin ke dev_b dengan arah pengiriman host to device. 31. Inisialisasi block dengan sumbu vektor x dan y. 32. Inisialisasi thread dengan sumbu vektor x dan y. 33. Timer pada GPU dinyalakan. 34. Memanggil fungsi kernel dengan nama add dengan konfigurasi block dan thread 2 dimensi dengan parameter yang dipassingkan dev_a, dev_b dan dev_c. 35. Fungsi cudaThreadSynchronize() bertujuan untuk memastikan semua thread selesai dan hasilnya siap untuk dikirimkan ke CPU. 37. Mematikan timer.
115
38. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari dev_c disalin ke c dengan arah pengiriman device to host. 39. Perulangan untuk mencetak baris dan kolom beserta isi dari array c yang merupakan hasil penjumlahan dari int a dan b. 47. Mencetak waktu eksekusi GPU. 48. Membebaskan alokasi memory dev_a. 49. Membebaskan alokasi memory dev_b. 50. Membebaskan alokasi memory dev_c. 51. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 52. Kurung tutup tanda berakhirnya program
Soal 4.2. Kode Sumber [1]: 1. #include <cuda.h> 2. #include <stdio.h> 3. #define lebar 4 4. __global__ void kalimatrix( float * A, float *B, float *C) 5. { 6. int kolom = threadIdx.x; 7. int baris = threadIdx.y; 8. float p=0; 9. for( int k =0; k < lebar; ++k) 10. { 11. float A_elem= A[baris *lebar +k]; 12. float B_elem= B[k *lebar +kolom]; 13. p +=A_elem * B_elem; 14. } 15. C[baris * lebar + kolom] = p; 16. __syncthreads(); 17. } 18. int main(void) 19. { 20. Float *M_C, *dMA, *dMB, *dMC; 21. int N = 16; 22. float M_A[16]={1,4,5,6,7,9,10,2,4,1,1,4,6,10,4,5}; 23. float M_B[16]={1,1,1,3,1,1,4,5,2,2,2,2,2,2,2,2}; 24. size_t size = N * sizeof(float); 116
25. 26. 27. 28. 29. 30. 31. 32. 33. 34. 35. 36. 37. 38. 39. 40. 41. 42. 43. 44. 45. 46. 47. 48. 49. 50. 51. 52. 53. 54. 55. 56. 57. 58. 59. 60. 61. 62. 63. 64. }
M_C = (float *)malloc(size); cudaMalloc((void **) &dMC, size); cudaMalloc((void **) &dMA, size); cudaMalloc((void **) &dMB, size); cudaMemcpy(dMC, M_C , size, cudaMemcpyHostToDevice); cudaMemcpy(dMA, M_A, size, cudaMemcpyHostToDevice); cudaMemcpy(dMB, M_B, size, cudaMemcpyHostToDevice); printf("\n"); printf("matrik A: \n"); for (int a=0;a >> (dMA,dMB,dMC); cudaMemcpy(M_C, dMC, sizeof(float)*N, cudaMemcpyDeviceToHost); printf("matrik C: \n"); for (int a=0;a
Penjelasan program per baris: 1. Inisialisasi pustaka <cuda.h>. 2. Inisialisasi pustaka <studio.h>. 3. Mendefinisikan lebar dengan nilai 4. 4. Inisialisasi program kernel dengan nama kalimatrix dengan passing parameter float a, b dan c. 5. Kurung buka program kernel. 117
6. Inisialisasi thread vektor x dengan nama kolom. 7. Inisialisasi thread vektor y dengan nama baris. 8. Inisialisasi float p dengan nilai awal 0. 9. Perulangan untuk mengalikan matrik. 11. Inisialisasi elemen kolom pada matrik A yang disimpan pada float A_elem. 12. Inisialisasi elemen baris pada matrik B yang disimpan pada float B_elem. 13. Perkalian matrik A dan B dan disimpan pada float p. 15. Nilai p disimpan pada matrik C dengan urutan mulai alamat 0 sampai selesai. 17. Kurung tutup tanda akhir program kernel. 18. Deklarasi program utama main, program akan dimulai dari main program. 20. Inisialisasi pointer M_C, dMA, dMB, dMC. 21. Inisialisasi int N dengan jumlah 16. 22. Inisialisasi float M_A dengan data 1,4,5,6,7,9,10,2,4,1,1,4,6,10,4,5. 23. Inisialisasi float M_B dengan data 1,1,1,3,1,1,4,5,2,2,2,2,2,2,2,2. 24. Deklarasi size dengan tipe data float dengan jumlah N. 25. Alokasi M_C pada CPU memory. 26. Alokasi dMC pada GPU memory. 27. Alokasi dMB pada GPU memory. 28. Alokasi dMA pada GPU memory. 29. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari M_C disalin ke dMC dengan arah pengiriman host to device. 30. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari M_A disalin ke dMA dengan arah pengiriman host to device. 31. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari M_B disalin ke dMB dengan arah pengiriman host to device. 33. Mencetak isi matrik A. 42. Mencetak isi matrik B. 51. Inisialisasi jumlah threadIdx.x dan threadIdx.y. 52. Pemanggilan program kernel kalimatrix. 118
53. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari dMC disalin ke M_C dengan arah pengiriman device to host. 54. Mencetak isi matrik C. 64. kurung tutup tanda akhir program.
6.2. Set Dosen 2 Soal 5.1. Kode Sumber [6]: 1.#include <stdio.h> 2.#include <cuda.h> 3.#include <cutil.h> 4.#define COLUMNS 6 5.#define ROWS 6 6.__global__ void add(int *a, int *b, int *c) 7.{ 8.int x = threadIdx.x; 9.int y = threadIdx.y; 10. int i = (COLUMNS*y) + x; 11. c[i] = a[i] * b[i]; 12. __syncthreads(); 13. } 14. int main(void) 15. { 16. unsigned int timer=0; 17. cutCreateTimer( &timer ); 18. int a[ROWS][COLUMNS], b[ROWS][COLUMNS], c[ROWS][COLUMNS]; 19. int *dev_a, *dev_b, *dev_c; 20. cudaMalloc((void **) &dev_a, ROWS*COLUMNS*sizeof(int)); 21. cudaMalloc((void **) &dev_b, ROWS*COLUMNS*sizeof(int)); 22. cudaMalloc((void **) &dev_c, ROWS*COLUMNS*sizeof(int)); 23. for (int y = 0; y < ROWS; y++) 24. for (int x = 0; x < COLUMNS; x++) 25. { 26. a[y][x] = x; 27. b[y][x] = y; 28. }
119
29. cudaMemcpy(dev_a, a, ROWS*COLUMNS*sizeof(int), cudaMemcpyHostToDevice); 30. cudaMemcpy(dev_b, b, ROWS*COLUMNS*sizeof(int), cudaMemcpyHostToDevice); 31. dim3 grid(COLUMNS,ROWS); 32. dim3 threadperblock(6,4); 33. cutStartTimer( timer ); 34. add<<>>(dev_a, dev_b, dev_c); 35. cudaThreadSynchronize(); 36. cudaThreadSynchronize(); 37. cutStopTimer( timer ); 38. cudaMemcpy(c, dev_c, ROWS*COLUMNS*sizeof(int), cudaMemcpyDeviceToHost); 39. for (int y = 0; y < ROWS; y++) 40. { 41. for (int x = 0; x < COLUMNS; x++) 42. { 43. printf("[%d][%d]=%d ",y,x,c[y][x]); 44. } 45. printf("\n"); 46. } 47. printf("CUDA execution time = %f ms\n",cutGetTimerValue( timer )); 48. cudaFree(dev_a); 49. cudaFree(dev_b); 50. cudaFree(dev_c); 51. system("pause"); 52. }
Penjelasan program per baris: 1. inisialisasi pustaka <stdio.h>. 2. inisialisasi pustaka <cuda.h>. 3. inisialisasi pustaka <cutil.h>. 4. Mendefinisikan kolom dengan lebar 6. 5. Mendefinisikan baris dengan tinggi 6. 6. Inisialisasi program kernel dengan nama add dengan passing parameter int a, b dan c. 7. Kurung buka tanda awal program kernel. 8. Inisialisasi integer x sebagai threadIdx.x dengan vektor sumbu x. 9. Inisialisasi integer y sebagai threadIdx.x dengan vektor sumbu y. 10. Inisialisasi integer i sebagai perkalian kolom dengan int y ditambah dengan int x, threadIdx.y akan mulai mengeksekusi setelah threadIdx.x selesai dieksekusi.
120
11. Merupakan perkalian float a dengan float b dengan menggunakan thread i sebagai eksekutornya setelah thread i dengan thread ID 0 sampai 5 selesai mengeksekusi maka otomatis dilanjutkan dengan thread ID 6 sampai 11 begitu seterusnya sampai batas threadIdx.y. 12. __syncthreads(); merupakan fungsi memastikan semua thread telah selesai melakukan instruksi sebelum dilanjutkan ke baris berikutnya. 13. Kurung tutup tanda berakhirnya program kernel. 14. Deklarasi program utama main, program akan dimulai dari main program. 15. Kurung buka program. 16. Inisialisasi timer dengan nilai awal 0. 17. Pemanggilan fungsi cutCreateTimer pada library <cutil.h> dengan menunjuk pointer timer. 18. Inisialisasi int a,b dan c dengan array 2 dimensi. 19. Inisialiasasi pointer dev_a,dev_b dan dev_c. 20. Inisialisasi dev_a pada GPU memory. 21. Inisialisasi dev_b pada GPU memory. 22. Inisialisasi dev_c pada GPU memory. 23. Perulangan baris. 24. Perulangan kolom. 25. Pengisian kolom dan baris pada alamat array int a dan b beserta nilainya, nilai int a akan diisi dengan nilai kolom sedangkan int b akan diisi dengan nilai baris. 29. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari a disalin ke dev_a dengan arah pengiriman host to device. 30. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari b disalin ke dev_b dengan arah pengiriman host to device. 31. Inisialisasi block dengan sumbu vektor x dan y. 32. Inisialisasi thread dengan sumbu vektor x dan y. 33. Timer pada GPU dinyalakan.
121
34. Memanggil fungsi kernel dengan nama add dengan konfigurasi block dan thread 2 dimensi dengan parameter yang dipassingkan dev_a, dev_b dan dev_c. 35. Fungsi cudaThreadSynchronize() bertujuan untuk memastikan semua thread selesai dan hasilnya siap untuk dikirimkan ke CPU. 37. Mematikan timer. 38. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari dev_c disalin ke c dengan arah pengiriman device to host. 39. Perulangan untuk mencetak baris dan kolom beserta isi dari array c yang merupakan hasil perkalian dari array int a dan b. 47. Mencetak waktu eksekusi GPU. 48. Membebaskan alokasi memory dev_a. 49. Membebaskan alokasi memory dev_b. 50. Membebaskan alokasi memory dev_c. 51. Menahan tampilan agar menunggu input yang masuk baru program akan keluar. 52. Kurung tutup tanda berakhirnya program. Soal 5.2. Kode Sumber [1]: 1.#include <cuda.h> 2.#include <stdio.h> 3.#define lebar 4 4.__global__ void kalimatrix( float * A, float *B, float *C) 5.{ 6.int kolom = threadIdx.x; 7.int baris = threadIdx.y; 8.float p=0; 9.for( int k =0; k < lebar; ++k) 10. { 11. float A_elem= A[baris *lebar +k]; 12. float B_elem= B[k *lebar +kolom]; 13. p +=A_elem * B_elem; 14. } 15. C[baris * lebar + kolom] = p; 16. __syncthreads(); 17. } 122
18. int main(void) 19. { 20. Float *M_C, *dMA, *dMB, *dMC; 21. int N = 16; 22. float M_A[16]={1,4,5,6,7,9,10,2,4,1,1,4,6,10,4,5}; 23. float M_B[16]={1,1,1,3,1,1,4,5,2,2,2,2,2,2,2,2}; 24. size_t size = N * sizeof(float); 25. M_C = (float *)malloc(size); 26. cudaMalloc((void **) &dMC, size); 27. cudaMalloc((void **) &dMA, size); 28. cudaMalloc((void **) &dMB, size); 29. cudaMemcpy(dMC, M_C , size, cudaMemcpyHostToDevice); 30. cudaMemcpy(dMA, M_A, size, cudaMemcpyHostToDevice); 31. cudaMemcpy(dMB, M_B, size, cudaMemcpyHostToDevice); 32. printf("\n"); 33. printf("matrik A: \n"); 34. for (int a=0;a >> (dMA,dMB,dMC); 53. cudaMemcpy(M_C, dMC, sizeof(float)*N, cudaMemcpyDeviceToHost); 54. printf("matrik C: \n"); 55. for (int a=0;a
Penjelasan program per baris: 1. Inisialisasi pustaka <cuda.h>. 2. Inisialisasi pustaka <studio.h>. 123
3. Mendefinisikan lebar dengan nilai 4. 4. Inisialisasi program kernel dengan nama kalimatrix dengan passing parameter float a, b dan c. 5. Kurung buka program kernel. 6. Inisialisasi thread vektor x dengan nama kolom. 7. Inisialisasi thread vektor y dengan nama baris. 8. Inisialisasi float p dengan nilai awal 0. 9. Perulangan untuk mengalikan matrik. 14. Inisialisasi elemen kolom pada matrik A yang disimpan pada float A_elem. 15. Inisialisasi elemen baris pada matrik B yang disimpan pada float B_elem. 16. Perkalian matrik A dan B dan disimpan pada float p. 16. Nilai p disimpan pada matrik C dengan urutan mulai alamat 0 sampai selesai. 19. Kurung tutup tanda akhir program kernel. 20. Deklarasi program utama main, program akan dimulai dari main program. 32. Inisialisasi pointer M_C, dMA, dMB, dMC. 33. Inisialisasi int N dengan jumlah 16. 34. Inisialisasi float M_A dengan data 1,4,5,6,7,9,10,2,4,1,1,4,6,10,4,5. 35. Inisialisasi float M_B dengan data 1,1,1,3,1,1,4,5,2,2,2,2,2,2,2,2. 36. Deklarasi size dengan tipe data float dengan jumlah N. 37. Alokasi M_C pada CPU memory. 38. Alokasi dMC pada GPU memory. 39. Alokasi dMB pada GPU memory. 40. Alokasi dMA pada GPU memory. 41. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari M_C disalin ke dMC dengan arah pengiriman host to device. 42. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari M_A disalin ke dMA dengan arah pengiriman host to device. 43. Menyalin memory dari CPU ke GPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari M_B disalin ke dMB dengan arah pengiriman host to device. 124
34. Mencetak isi matrik A. 43. Mencetak isi matrik B. 55. Inisialisasi jumlah threadIdx.x dan threadIdx.y. 56. Pemanggilan program kernel kalimatrix. 57. Menyalin memory dari GPU ke CPU dengan format (tujuan, sumber , ukuran , arah pengiriman). Isi dari dMC disalin ke M_C dengan arah pengiriman device to host. 58. Mencetak isi matrik C. 64. kurung tutup tanda akhir program.
7.Daftar Pustaka 1. cse.shirazu.ac.ir, ”MatmultBasic”, diakses dalam http://www.cse.shirazu.ac.ir/~azimi/gpu89/lectures/07-MatMult-Basic.pdf diakses pada 14 September 2012. 2. Kirk David B., Hwu Wen-mei W., ”Programming Massively Parallel Processors”, MK, 2010. 3. Nvidia, ”Nvidia CUDA C Programming Guide version 4”, Nvidia,2011. 4. Nvidia, “Getting Started With CUDA”, Nvidia, 2008. 5. Parallel Panorama, ”Thread and blocks and grids”, diakses dalam http://llpanorama.wordpress.com/2008/06/11/threads-and-blocks-and-grids-ohmy/ diakses pada 3 Agustus 2012. 6. Stack Overflow, “Indexing scheme vs array of pointers in CUDA”, diakses dalam http://stackoverflow.com/questions/11817471/indexing-scheme-vs-array-ofpointers-in-cuda diakses pada 3 Agustus 2012.
125
126
