three working kernels

Author: divyegala

cpp/include/learning_cuda/count/count.cuh

@@ -52,4 +52,40 @@ int count_if(T* arr, int size,
     return count_h;
 }
 
-} // namespace manual_reduction
+} // namespace manual_reduction
+
+namespace syncthreads_count_reduction {
+
+template <typename T, class F>
+int count_if(T* arr, int size,
+                F count_if_op) {
+    int *count_d, count_h, N_BLK;
+
+    common::_prep_count_if(&count_d, N_BLK, size);
+
+    count_kernel<<<N_BLK, TPB>>> (arr, size, count_d, count_if_op);
+
+    common::_finish_count_if(&count_d, count_h);
+
+    return count_h;
+}
+
+} // namespace syncthreads_count_reduction
+
+namespace ballot_sync_reduction {
+
+template <typename T, class F>
+int count_if(T* arr, int size,
+                F count_if_op) {
+    int *count_d, count_h, N_BLK;
+
+    common::_prep_count_if(&count_d, N_BLK, size);
+
+    count_kernel<<<N_BLK, TPB>>> (arr, size, count_d, count_if_op);
+
+    common::_finish_count_if(&count_d, count_h);
+
+    return count_h;
+}
+
+} // namespace ballot_count_reduction

cpp/include/learning_cuda/count/count_kernels.cuh

@@ -1,4 +1,4 @@
-#include <cuda_runtime.h>
+// #include <cuda_runtime.h>
 
 namespace naive {
 
@@ -38,7 +38,7 @@ void count_kernel(T* arr, int size, int *count, F count_if_op) {
 
     if (tid < size) {
         for (int offset = blockDim.x / 2; offset > 0; offset /=2 ) {
-            local_count_array[threadIdx.x] = local_count_array[threadIdx.x + offset];
+            local_count_array[threadIdx.x] += local_count_array[threadIdx.x + offset];
             __syncthreads();
         }
     }
@@ -48,4 +48,45 @@ void count_kernel(T* arr, int size, int *count, F count_if_op) {
     }
 }
 
-} // namespace manual_reduction
+} // namespace manual_reduction
+
+namespace syncthreads_count_reduction {
+
+template <typename T, class F>
+__global__ 
+void count_kernel(T* arr, int size, int *count, F count_if_op) {
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
+
+    bool predicate = tid < size && count_if_op(arr[tid]);
+
+    int block_count = __syncthreads_count(predicate);
+
+    if(threadIdx.x == 0) {
+        atomicAdd(count, block_count);
+    }
+    
+}
+
+} // namespace syncthreads_count
+
+#define FULL_MASK 0xffffffff
+
+namespace ballot_sync_reduction {
+
+template <typename T, class F>
+__global__ 
+void count_kernel(T* arr, int size, int *count, F count_if_op) {
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
+
+    bool predicate = tid < size && count_if_op(arr[tid]);
+
+    unsigned ballot_mask = __ballot_sync(FULL_MASK, predicate);
+    int warp_count = __popc(ballot_mask);
+
+    if(threadIdx.x % 32 == 0) {
+        atomicAdd(count, warp_count);
+    }
+    
+}
+
+} // ballot syncthreads_count

cpp/src/count/count.cu

@@ -12,7 +12,7 @@
 
 
 int main() {
-    int n_elems = 100000;
+    int n_elems = 1000;
 
     thrust::device_vector<int> rand_d(n_elems, 0);
 
@@ -38,10 +38,20 @@ int main() {
 
     std::cout << "\nNaive Count: " << naive_count << std::endl;
 
-    int man_red_count = naive::count_if(thrust::raw_pointer_cast(rand_d.data()),
+    int man_count = manual_reduction::count_if(thrust::raw_pointer_cast(rand_d.data()),
     n_elems, is_greater_than_10);
 
-    std::cout << "\nManual Reduction Count: " << man_red_count << std::endl;
+    std::cout << "\nManual Reduction Count: " << man_count << std::endl;
+
+    int syn_count = syncthreads_count_reduction::count_if(thrust::raw_pointer_cast(rand_d.data()),
+    n_elems, is_greater_than_10);
+
+    std::cout << "\nSyncthread Reduction Count: " << syn_count << std::endl;
+
+    int bal_count = ballot_sync_reduction::count_if(thrust::raw_pointer_cast(rand_d.data()),
+    n_elems, is_greater_than_10);
+
+    std::cout << "\nBallot Reduction Count: " << bal_count << std::endl;
 
     return 0;
 }