edit radix par

src/radix_sort_par.cu

@@ -19,11 +19,14 @@ int getMax(int arr[], int n)
 __global__
 void countSort(int arr[], int n, int exp) 
 { 
+	int index = threadIdx.x;
+    int stride = blockDim.x;
+
     int *output= (int*)malloc(sizeof(int)*n); // output array 
     int i, count[10] = {0}; 
   
     // Store count of occurrences in count[] 
-    for (i = 0; i < n; i++) 
+    for (i = index; i < n; i+=stride) 
         count[ (arr[i]/exp)%10 ]++; 
   
     // Change count[i] so that count[i] now contains actual 
@@ -32,7 +35,7 @@ void countSort(int arr[], int n, int exp)
         count[i] += count[i - 1]; 
   
     // Build the output array 
-    for (i = n - 1; i >= 0; i--) 
+    for (i = n - 1; i >= index; i-=stride) 
     { 
         output[count[ (arr[i]/exp)%10 ] - 1] = arr[i]; 
         count[ (arr[i]/exp)%10 ]--; 
@@ -40,7 +43,7 @@ void countSort(int arr[], int n, int exp)
   
     // Copy the output array to arr[], so that arr[] now 
     // contains sorted numbers according to current digit 
-    for (i = 0; i < n; i++) 
+    for (i = index; i < n; i+=stride) 
         arr[i] = output[i]; 
 } 
   
@@ -63,11 +66,11 @@ void radixsort(int *arr, int n)
 
     cudaMemcpy(d_arr, arr, sizeof(int)*n,cudaMemcpyHostToDevice);
     for (int exp = 1; m/exp > 0; exp *= 10){ 
-        countSort<<<1,32>>>(d_arr, n, exp); 
+        countSort<<<1,1024>>>(d_arr, n, exp); 
 
-        //transfer data back to host memory
-        cudaMemcpy(arr, d_arr, sizeof(int)*n, cudaMemcpyDeviceToHost);
-    }
+    //transfer data back to host memory
+    cudaMemcpy(arr, d_arr, sizeof(int)*n, cudaMemcpyDeviceToHost);
+    
     //deallocate device memory
     cudaFree(d_arr);