From 125f5d055d251ce3f8a13d079dfee1bfd8916675 Mon Sep 17 00:00:00 2001
From: Habibi <muhhabibih@gmail.com>
Date: Thu, 11 Apr 2019 17:23:59 +0700
Subject: [PATCH 01/10] add radix sort cu

---
 radix_sort.cu | 95 +++++++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 95 insertions(+)
 create mode 100644 radix_sort.cu

diff --git a/radix_sort.cu b/radix_sort.cu
new file mode 100644
index 0000000..967c3d1
--- /dev/null
+++ b/radix_sort.cu
@@ -0,0 +1,95 @@
+#include "cuda_runtime.h"
+#include <stdio.h>
+#include <math.h>
+#include <time.h>
+#include <iostream>
+
+#define SEED 13516085
+#define NUMBER_OF_BLOCKS 32
+#define NUMBER_OF_THREADS 256
+
+using namespace std;
+
+
+int* harr;
+
+__host__ void create_random_array(int* arr, int n) {
+	srand(SEED);
+	for (int i = 0; i < n; i++) {
+		arr[i] = (int)rand();
+	}
+}
+
+__global__ void fill_count(int* arr, int n, int* counts) {
+  int idx = threadIdx.x + blockIdx.x * blockDim.x;
+  int bitmask = 1 << blockIdx.x;
+
+  counts[idx] = 0;
+
+  for(int i = threadIdx.x; i < n; i += NUMBER_OF_THREADS) {
+    if (!(arr[i] & bitmask)) {
+      counts[idx]++;
+    }
+  }
+}
+
+void radix_sort_parallel(int* arr, int n) {
+  int *d_arr;
+  int *d_counts;
+  int *h_counts = (int*) malloc(NUMBER_OF_BLOCKS*NUMBER_OF_THREADS*sizeof(int));
+  int *temp_arr = (int*) malloc(n*sizeof(int)); 
+  int counts_per_bit[32] = {0};
+
+  cudaMalloc((void**)&d_counts, NUMBER_OF_BLOCKS*NUMBER_OF_THREADS*sizeof(int));
+  cudaMalloc((void**)&d_arr, n*sizeof(int));
+
+  cudaMemcpy(d_arr, arr, n*sizeof(int), cudaMemcpyHostToDevice);
+
+  fill_count<<<NUMBER_OF_BLOCKS, NUMBER_OF_THREADS>>>(d_arr, n, d_counts);
+  cudaDeviceSynchronize();
+
+  cudaMemcpy(h_counts, d_counts, NUMBER_OF_BLOCKS*NUMBER_OF_THREADS*sizeof(int), cudaMemcpyDeviceToHost);
+
+  for(int i = 0; i < 32; i++) {
+    for(int j = 0; j < NUMBER_OF_THREADS; j++) {
+      counts_per_bit[i] += h_counts[i*NUMBER_OF_THREADS + j];
+    }
+  }
+
+  for(int i = 0; i < 32; i++) {
+    int banyak1 = 0;
+    int banyak0 = 0;
+    int bitmask = 1 << i;
+
+    for(int j = 0; j < n; j++) {
+      if (arr[j] & bitmask) {
+        temp_arr[counts_per_bit[i] + banyak1] = arr[j]; 
+        banyak1++;
+      } else {
+        temp_arr[banyak0] = arr[j];
+        banyak0++;
+      }
+    }
+
+    memcpy(arr, temp_arr, n*sizeof(int));
+  }
+
+  cudaFree(d_arr);
+  cudaFree(d_counts);
+}
+
+int main(int argc, char *argv[]) {
+  int number_of_int = strtol(argv[1], NULL, 10);
+
+  harr = (int*) malloc(number_of_int*sizeof(int));
+
+  create_random_array(harr, number_of_int);
+  
+  clock_t start = clock();
+  radix_sort_parallel(harr, number_of_int);
+  clock_t end = clock();
+
+  printf("Waktu yang dibutuhkan untuk Radix Sort Parallel: %lfms\n", (double)(end - start)/CLOCKS_PER_SEC * 1000 * 1000);
+
+  return 0;
+}
\ No newline at end of file
-- 
GitLab


From d0bda9a86267748336fd0211ce5e7f2e2b60fc50 Mon Sep 17 00:00:00 2001
From: devipramita <devipram99@gmail.com>
Date: Thu, 11 Apr 2019 20:57:25 +0700
Subject: [PATCH 02/10] add readme

---
 README.md | 1 +
 1 file changed, 1 insertion(+)
 create mode 100644 README.md

diff --git a/README.md b/README.md
new file mode 100644
index 0000000..2179e8b
--- /dev/null
+++ b/README.md
@@ -0,0 +1 @@
+# CUDA
\ No newline at end of file
-- 
GitLab


From 1e73e105ba48da89af0ba309b658e791f2eaf203 Mon Sep 17 00:00:00 2001
From: devipramita <devipram99@gmail.com>
Date: Thu, 11 Apr 2019 20:58:45 +0700
Subject: [PATCH 03/10] add readme

---
 README.md | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 2179e8b..76f10bf 100644
--- a/README.md
+++ b/README.md
@@ -1 +1,3 @@
-# CUDA
\ No newline at end of file
+# CUDA
+
+### 1. 
\ No newline at end of file
-- 
GitLab


From 7a336d5f0853944f9bb18dee9e72957e35f32aa5 Mon Sep 17 00:00:00 2001
From: devipramita <devipram99@gmail.com>
Date: Thu, 11 Apr 2019 21:01:30 +0700
Subject: [PATCH 04/10] update readme

---
 README.md | 10 +++++++++-
 1 file changed, 9 insertions(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 76f10bf..6ba7f00 100644
--- a/README.md
+++ b/README.md
@@ -1,3 +1,11 @@
 # CUDA
 
-### 1. 
\ No newline at end of file
+### 1. Deskripsi Solusi Paralel
+
+### 2. Analisis Solusi yang diberikan
+
+### 3. Jelaskan pemetaan ke thread komputasi dan jumlah thread per block yang digunakan
+
+### 4. Pengukuran kinerja untuk tiap kasus uji
+
+### 5. Analisis perbandingan kinerja serial dan paralel
\ No newline at end of file
-- 
GitLab


From ffb72a9cd6d785c2d17b412103c5ee0d1ff984f4 Mon Sep 17 00:00:00 2001
From: devipramita <devipram99@gmail.com>
Date: Thu, 11 Apr 2019 21:27:28 +0700
Subject: [PATCH 05/10] update readme 1,3,5

---
 README.md | 11 +++++++++--
 1 file changed, 9 insertions(+), 2 deletions(-)

diff --git a/README.md b/README.md
index 6ba7f00..1f4d557 100644
--- a/README.md
+++ b/README.md
@@ -1,11 +1,18 @@
 # CUDA
 
 ### 1. Deskripsi Solusi Paralel
+Memparalelkan proses pengurutan Radix Sort pada bagian menghitung jumlah binary 0 dan 1 di fungsi Fill Count. <br />
+Penghitungan jumlah bit 0 dilakukan pada setiap thread untuk tiap block, menggunakan threadIdx, kemudian hasilnya dimasukkan ke dalam sebuah array baru berisikan hasil penghitungan tersebut.
 
 ### 2. Analisis Solusi yang diberikan
+##### apakah ada solusi yang memberikan kinerja lebih baik?
 
-### 3. Jelaskan pemetaan ke thread komputasi dan jumlah thread per block yang digunakan
+
+### 3. Jelaskan pemetaan thread ke komputasi dan jumlah thread per block yang digunakan
+Jumlah thread yang digunakan per block adalah 256 threads.
 
 ### 4. Pengukuran kinerja untuk tiap kasus uji
 
-### 5. Analisis perbandingan kinerja serial dan paralel
\ No newline at end of file
+
+### 5. Analisis perbandingan kinerja serial dan paralel
+Kinerja paralel memberikan kinerja yang lebih lambat dibandingkan kinerja serial.
\ No newline at end of file
-- 
GitLab


From 4295bf251e84b35f48e56321440b3e0d5fde3dbd Mon Sep 17 00:00:00 2001
From: devipramita <devipram99@gmail.com>
Date: Thu, 11 Apr 2019 21:49:47 +0700
Subject: [PATCH 06/10] modified file structure

---
 README.md         | 22 +++++++----
 src/radix_sort.cu | 95 +++++++++++++++++++++++++++++++++++++++++++++++
 2 files changed, 109 insertions(+), 8 deletions(-)
 create mode 100644 src/radix_sort.cu

diff --git a/README.md b/README.md
index 1f4d557..8da379b 100644
--- a/README.md
+++ b/README.md
@@ -1,18 +1,24 @@
 # CUDA
 
-### 1. Deskripsi Solusi Paralel
-Memparalelkan proses pengurutan Radix Sort pada bagian menghitung jumlah binary 0 dan 1 di fungsi Fill Count. <br />
-Penghitungan jumlah bit 0 dilakukan pada setiap thread untuk tiap block, menggunakan threadIdx, kemudian hasilnya dimasukkan ke dalam sebuah array baru berisikan hasil penghitungan tersebut.
+### Petunjuk Penggunaan Program
 
-### 2. Analisis Solusi yang diberikan
-##### apakah ada solusi yang memberikan kinerja lebih baik?
 
+### Pembagian Tugas
+13516085 M. Habibi Haidir
+13516142 Maharani Devira
 
-### 3. Jelaskan pemetaan thread ke komputasi dan jumlah thread per block yang digunakan
+### Laporan Pengerjaan
+##### 1. Deskripsi Solusi Paralel
+Memparalelkan proses pengurutan Radix Sort pada bagian menghitung jumlah binary 0 dan 1 di fungsi Fill Count. Penghitungan jumlah bit 0 dilakukan pada setiap thread untuk tiap block, menggunakan threadIdx, kemudian hasilnya dimasukkan ke dalam sebuah array baru berisikan hasil penghitungan pada seluruh thread yang ada. Dari array baru inilah proses radix sort akan lebih mudah dijalankan dan setiap elemen angka input terurut akan dituliskan secara serial ke dalma sebuah array final.
+
+##### 2. Analisis Solusi yang diberikan, apakah ada solusi yang memberikan kinerja lebih baik?
+Kemungkinan kinerja akan lebih baik lagi jika penulisan array final dilakukan secara paralel juga sehingga proses radix sort dapat berjalan lebih cepat lagi.
+
+##### 3. Jelaskan pemetaan thread ke komputasi dan jumlah thread per block yang digunakan
 Jumlah thread yang digunakan per block adalah 256 threads.
 
-### 4. Pengukuran kinerja untuk tiap kasus uji
+##### 4. Pengukuran kinerja untuk tiap kasus uji
 
 
-### 5. Analisis perbandingan kinerja serial dan paralel
+##### 5. Analisis perbandingan kinerja serial dan paralel
 Kinerja paralel memberikan kinerja yang lebih lambat dibandingkan kinerja serial.
\ No newline at end of file
diff --git a/src/radix_sort.cu b/src/radix_sort.cu
new file mode 100644
index 0000000..967c3d1
--- /dev/null
+++ b/src/radix_sort.cu
@@ -0,0 +1,95 @@
+#include "cuda_runtime.h"
+#include <stdio.h>
+#include <math.h>
+#include <time.h>
+#include <iostream>
+
+#define SEED 13516085
+#define NUMBER_OF_BLOCKS 32
+#define NUMBER_OF_THREADS 256
+
+using namespace std;
+
+
+int* harr;
+
+__host__ void create_random_array(int* arr, int n) {
+	srand(SEED);
+	for (int i = 0; i < n; i++) {
+		arr[i] = (int)rand();
+	}
+}
+
+__global__ void fill_count(int* arr, int n, int* counts) {
+  int idx = threadIdx.x + blockIdx.x * blockDim.x;
+  int bitmask = 1 << blockIdx.x;
+
+  counts[idx] = 0;
+
+  for(int i = threadIdx.x; i < n; i += NUMBER_OF_THREADS) {
+    if (!(arr[i] & bitmask)) {
+      counts[idx]++;
+    }
+  }
+}
+
+void radix_sort_parallel(int* arr, int n) {
+  int *d_arr;
+  int *d_counts;
+  int *h_counts = (int*) malloc(NUMBER_OF_BLOCKS*NUMBER_OF_THREADS*sizeof(int));
+  int *temp_arr = (int*) malloc(n*sizeof(int)); 
+  int counts_per_bit[32] = {0};
+
+  cudaMalloc((void**)&d_counts, NUMBER_OF_BLOCKS*NUMBER_OF_THREADS*sizeof(int));
+  cudaMalloc((void**)&d_arr, n*sizeof(int));
+
+  cudaMemcpy(d_arr, arr, n*sizeof(int), cudaMemcpyHostToDevice);
+
+  fill_count<<<NUMBER_OF_BLOCKS, NUMBER_OF_THREADS>>>(d_arr, n, d_counts);
+  cudaDeviceSynchronize();
+
+  cudaMemcpy(h_counts, d_counts, NUMBER_OF_BLOCKS*NUMBER_OF_THREADS*sizeof(int), cudaMemcpyDeviceToHost);
+
+  for(int i = 0; i < 32; i++) {
+    for(int j = 0; j < NUMBER_OF_THREADS; j++) {
+      counts_per_bit[i] += h_counts[i*NUMBER_OF_THREADS + j];
+    }
+  }
+
+  for(int i = 0; i < 32; i++) {
+    int banyak1 = 0;
+    int banyak0 = 0;
+    int bitmask = 1 << i;
+
+    for(int j = 0; j < n; j++) {
+      if (arr[j] & bitmask) {
+        temp_arr[counts_per_bit[i] + banyak1] = arr[j]; 
+        banyak1++;
+      } else {
+        temp_arr[banyak0] = arr[j];
+        banyak0++;
+      }
+    }
+
+    memcpy(arr, temp_arr, n*sizeof(int));
+  }
+
+  cudaFree(d_arr);
+  cudaFree(d_counts);
+}
+
+int main(int argc, char *argv[]) {
+  int number_of_int = strtol(argv[1], NULL, 10);
+
+  harr = (int*) malloc(number_of_int*sizeof(int));
+
+  create_random_array(harr, number_of_int);
+  
+  clock_t start = clock();
+  radix_sort_parallel(harr, number_of_int);
+  clock_t end = clock();
+
+  printf("Waktu yang dibutuhkan untuk Radix Sort Parallel: %lfms\n", (double)(end - start)/CLOCKS_PER_SEC * 1000 * 1000);
+
+  return 0;
+}
\ No newline at end of file
-- 
GitLab


From 2e37aecdc151c7f93e682fd2379573fb48386709 Mon Sep 17 00:00:00 2001
From: devipramita <devipram99@gmail.com>
Date: Thu, 11 Apr 2019 22:56:29 +0700
Subject: [PATCH 07/10] update readme 3,4

---
 README.md | 16 +++++++++++-----
 1 file changed, 11 insertions(+), 5 deletions(-)

diff --git a/README.md b/README.md
index 8da379b..c4b89d8 100644
--- a/README.md
+++ b/README.md
@@ -1,11 +1,11 @@
 # CUDA
 
 ### Petunjuk Penggunaan Program
-
+Menjalankan perintah "make"
 
 ### Pembagian Tugas
-13516085 M. Habibi Haidir
-13516142 Maharani Devira
+13516085 M. Habibi Haidir : fungsi Fill Count, Radix Sort Paralel, Main
+13516142 Maharani Devira : Main, Readme
 
 ### Laporan Pengerjaan
 ##### 1. Deskripsi Solusi Paralel
@@ -15,10 +15,16 @@ Memparalelkan proses pengurutan Radix Sort pada bagian menghitung jumlah binary
 Kemungkinan kinerja akan lebih baik lagi jika penulisan array final dilakukan secara paralel juga sehingga proses radix sort dapat berjalan lebih cepat lagi.
 
 ##### 3. Jelaskan pemetaan thread ke komputasi dan jumlah thread per block yang digunakan
-Jumlah thread yang digunakan per block adalah 256 threads.
+Jumlah block yang digunakan adalah 32 block dengan jumlah thread yang digunakan per block adalah 256 threads. Kedua angka tersebut dipilih karena disesuaikan dengan jumlah bit dari bilangan integer dan berpengaruh pada logika komputasi program ini yang bekerja paralel antar bit integer.
 
 ##### 4. Pengukuran kinerja untuk tiap kasus uji
-
+| N | Serial | Paralel |
+| -------------------- |
+| 5000 | 1787 ms | 229729 ms |
+| 50000 | 17628 ms | 249938 ms |
+| 100000 | 25556 ms | 263562 ms |
+| 200000 | 42970 ms | 298985 ms |
+| 400000 | 107584 ms | 369119 ms |
 
 ##### 5. Analisis perbandingan kinerja serial dan paralel
 Kinerja paralel memberikan kinerja yang lebih lambat dibandingkan kinerja serial.
\ No newline at end of file
-- 
GitLab


From 825653a70eeca08bd02c0c9e609de2f73bc4d81a Mon Sep 17 00:00:00 2001
From: devipramita <devipram99@gmail.com>
Date: Thu, 11 Apr 2019 22:59:40 +0700
Subject: [PATCH 08/10] update readme 3,4

---
 README.md | 17 +++++++++++------
 1 file changed, 11 insertions(+), 6 deletions(-)

diff --git a/README.md b/README.md
index c4b89d8..b50ddb4 100644
--- a/README.md
+++ b/README.md
@@ -18,13 +18,18 @@ Kemungkinan kinerja akan lebih baik lagi jika penulisan array final dilakukan se
 Jumlah block yang digunakan adalah 32 block dengan jumlah thread yang digunakan per block adalah 256 threads. Kedua angka tersebut dipilih karena disesuaikan dengan jumlah bit dari bilangan integer dan berpengaruh pada logika komputasi program ini yang bekerja paralel antar bit integer.
 
 ##### 4. Pengukuran kinerja untuk tiap kasus uji
-| N | Serial | Paralel |
-| -------------------- |
-| 5000 | 1787 ms | 229729 ms |
-| 50000 | 17628 ms | 249938 ms |
-| 100000 | 25556 ms | 263562 ms |
-| 200000 | 42970 ms | 298985 ms |
+| N      | Serial    | Paralel   |
+| :----: | :-------: | :-------: |
+| 5000   | 1787 ms   | 229729 ms |
+| 50000  | 17628 ms  | 249938 ms |
+| 100000 | 25556 ms  | 263562 ms |
+| 200000 | 42970 ms  | 298985 ms |
 | 400000 | 107584 ms | 369119 ms |
 
+| Left Aligned | Centered | Right Aligned | Left Aligned | Centered | Right Aligned |
+| :----------- | :------: | ------------: | :----------- | :------: | ------------: |
+| Cell 1       | Cell 2   | Cell 3        | Cell 4       | Cell 5   | Cell 6        |
+| Cell 7       | Cell 8   | Cell 9        | Cell 10      | Cell 11  | Cell 12       |
+
 ##### 5. Analisis perbandingan kinerja serial dan paralel
 Kinerja paralel memberikan kinerja yang lebih lambat dibandingkan kinerja serial.
\ No newline at end of file
-- 
GitLab


From 426a25df3935fbf91843863261934864123a47e6 Mon Sep 17 00:00:00 2001
From: devipramita <devipram99@gmail.com>
Date: Thu, 11 Apr 2019 23:00:35 +0700
Subject: [PATCH 09/10] update readme 3,4

---
 README.md | 5 -----
 1 file changed, 5 deletions(-)

diff --git a/README.md b/README.md
index b50ddb4..706359f 100644
--- a/README.md
+++ b/README.md
@@ -26,10 +26,5 @@ Jumlah block yang digunakan adalah 32 block dengan jumlah thread yang digunakan
 | 200000 | 42970 ms  | 298985 ms |
 | 400000 | 107584 ms | 369119 ms |
 
-| Left Aligned | Centered | Right Aligned | Left Aligned | Centered | Right Aligned |
-| :----------- | :------: | ------------: | :----------- | :------: | ------------: |
-| Cell 1       | Cell 2   | Cell 3        | Cell 4       | Cell 5   | Cell 6        |
-| Cell 7       | Cell 8   | Cell 9        | Cell 10      | Cell 11  | Cell 12       |
-
 ##### 5. Analisis perbandingan kinerja serial dan paralel
 Kinerja paralel memberikan kinerja yang lebih lambat dibandingkan kinerja serial.
\ No newline at end of file
-- 
GitLab


From f62c38c87d32177d383fe9ffc03bf10a967e8b7f Mon Sep 17 00:00:00 2001
From: devipramita <devipram99@gmail.com>
Date: Thu, 11 Apr 2019 23:01:39 +0700
Subject: [PATCH 10/10] update readme 3,4

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 706359f..27c281d 100644
--- a/README.md
+++ b/README.md
@@ -1,7 +1,7 @@
 # CUDA
 
 ### Petunjuk Penggunaan Program
-Menjalankan perintah "make"
+Menjalankan perintah ./radix_sort N (dengan N = 5000, 50000, 100000, 200000, 400000)
 
 ### Pembagian Tugas
 13516085 M. Habibi Haidir : fungsi Fill Count, Radix Sort Paralel, Main
-- 
GitLab