From 4b30c2411d7949dec7ffe66a4c8c070ce4b2a59f Mon Sep 17 00:00:00 2001 From: Adylan Roaffa <adylanrff@gmail.com> Date: Thu, 11 Apr 2019 00:01:16 +0700 Subject: [PATCH] add get max cuda --- src/cudaUtil/cudaUtil.cu | 34 ++++++++++++++++++++++++++++++++++ src/cudaUtil/cudaUtil.cuh | 3 ++- src/main.cu | 34 ++++++++++++++++++++++++++-------- src/radixSort/radixSort.cu | 10 ++++++++++ src/radixSort/radixSort.cuh | 3 ++- src/util/util.hpp | 1 - 6 files changed, 74 insertions(+), 11 deletions(-) diff --git a/src/cudaUtil/cudaUtil.cu b/src/cudaUtil/cudaUtil.cu index e69de29..430b9e2 100644 --- a/src/cudaUtil/cudaUtil.cu +++ b/src/cudaUtil/cudaUtil.cu @@ -0,0 +1,34 @@ +#include "cudaUtil.cuh" +#include <stdio.h> + +__device__ void getMax(int* arr, int n, int* maxBitIdx){ + int myId = threadIdx.x + blockDim.x * blockIdx.x; + int tId = threadIdx.x; + + for (unsigned int s = blockDim.x/2; s>0; s>>=1){ + if (tId < s){ + if (arr[myId] < arr[myId + s]){ + arr[myId] = arr[myId+s]; + } + } + __syncthreads(); + } + + if (tId == 0){ + int currentNum = arr[myId]; + int bitIndex = 0; + while(currentNum > 0){ + currentNum >>=1; + bitIndex +=1; + } + *maxBitIdx = bitIndex; + } +} + +__global__ void printArr(int* arr, int n){ + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i<n; i+=stride){ + printf("%d ", arr[i]); + } +} diff --git a/src/cudaUtil/cudaUtil.cuh b/src/cudaUtil/cudaUtil.cuh index e7373de..69ce35c 100644 --- a/src/cudaUtil/cudaUtil.cuh +++ b/src/cudaUtil/cudaUtil.cuh @@ -1,6 +1,7 @@ #ifndef CUDA_UTIL_CUH #define CUDA_UTIL_CUH - +__device__ void getMax(int* arr, int n, int* out); +__global__ void printArr(int* arr, int n); #endif \ No newline at end of file diff --git a/src/main.cu b/src/main.cu index 6a0c7b6..f8181a4 100644 --- a/src/main.cu +++ b/src/main.cu @@ -1,23 +1,41 @@ #include <stdio.h> #include "util/util.hpp" +#include "cudaUtil/cudaUtil.cuh" +#include "radixSort/radixSort.cuh" __global__ void cuda_hello(){ printf("Hello World from GPU!\n"); } -int main() { - int* arr; - int* d_arr; +int main(int argc, char *argv[]) { + + if (argc != 2){ + fprintf(stderr, "Usage: <number of elements>\n"); + exit(1); + } - arr = (int *) malloc(sizeof(int) * 10); - cudaMalloc((void**)&d_arr, sizeof(int) * 10); + int n = atoi(argv[1]); - cudaMemcpy(d_arr, arr, sizeof(float) * 10, cudaMemcpyHostToDevice); + int* arr = (int *) malloc(sizeof(int) * n); + + int* d_arr; + cudaMalloc((void**)&d_arr, sizeof(int) * n); - generate_random(arr, 10); - print_arr(arr, 10); + int* d_out; + cudaMalloc((void**)&d_out, sizeof(int) * n); + generate_random(arr, n); + + cudaMemcpy(d_arr, arr, sizeof(float) * n, cudaMemcpyHostToDevice); + cudaDeviceSynchronize(); + + radixSort<<<1,256>>>(d_out, d_arr, n); + cudaDeviceSynchronize(); + cuda_hello<<<1,1>>>(); cudaDeviceSynchronize(); + + cudaFree(d_arr); + return 0; } diff --git a/src/radixSort/radixSort.cu b/src/radixSort/radixSort.cu index 8e8917a..b539d5d 100644 --- a/src/radixSort/radixSort.cu +++ b/src/radixSort/radixSort.cu @@ -1,5 +1,15 @@ #include "radixSort.cuh" +#include "../cudaUtil/cudaUtil.cuh" +#include <stdio.h> +__global__ void radixSort(int* out, int* arr, int n){ + int* maxBitIdx = (int*) malloc(sizeof(int)); + + if (threadIdx.x == 0){ + getMax(arr,n,maxBitIdx); + printf("%d\n", *maxBitIdx); + } + __syncthreads(); } diff --git a/src/radixSort/radixSort.cuh b/src/radixSort/radixSort.cuh index 8409449..2814cd8 100644 --- a/src/radixSort/radixSort.cuh +++ b/src/radixSort/radixSort.cuh @@ -1,6 +1,7 @@ #ifndef RADIX_SORT_CUH #define RADIX_SORT_CUH -void radixSort(int* out, int* arr, int n); +__global__ void radixSort(int* out, int* arr, int n); +__device__ void sortMSB(int* out, int* arr, int n); #endif \ No newline at end of file diff --git a/src/util/util.hpp b/src/util/util.hpp index 5e23e74..114c8ee 100644 --- a/src/util/util.hpp +++ b/src/util/util.hpp @@ -3,7 +3,6 @@ void generate_random(int* arr, int n); void print_arr(int * array, int num_elements); -int get_max(int arr[], int n); void assert_sorted(int arr[], int n, double time_elapsed); #endif \ No newline at end of file -- GitLab