diff --git a/README.md b/README.md
new file mode 100644
index 0000000000000000000000000000000000000000..27c281dfd12e1546562142f7ddd86439bdd0ce41
--- /dev/null
+++ b/README.md
@@ -0,0 +1,30 @@
+# CUDA
+
+### Petunjuk Penggunaan Program
+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
+13516142 Maharani Devira : Main, Readme
+
+### 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 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
diff --git a/radix_sort.cu b/radix_sort.cu
new file mode 100644
index 0000000000000000000000000000000000000000..967c3d1a1a5fffa63f2604f5c3d54506119ad9b4
--- /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
diff --git a/src/radix_sort.cu b/src/radix_sort.cu
new file mode 100644
index 0000000000000000000000000000000000000000..967c3d1a1a5fffa63f2604f5c3d54506119ad9b4
--- /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