restructuring to accept TBP in template

Author: divyegala

conda/learn_dev_cuda10.2.yml

@@ -1,4 +1,4 @@
 name: learn_dev
 dependencies:
-- cmake>=3.14
+- cmake>=3.18
 - cudatoolkit=10.2

cpp/include/learning_cuda/count/count.cuh

@@ -1,9 +1,10 @@
-#include "count_kernels.cuh"
+#pragma once
 
-#define TPB 32
+#include "count_kernels.cuh"
 
 namespace common {
 
+template <int TPB>
 void _prep_count_if(int **count_d, int &N_BLK, int size) {
     cudaMalloc(count_d, sizeof(int));
     cudaMemset(*count_d, 0, sizeof(int));
@@ -20,14 +21,16 @@ void _finish_count_if(int **count_d, int &count_h) {
 
 namespace naive {
 
-template <typename T, class F>
-int count_if(T* arr, int size,
-             F count_if_op) {
+template <typename T, int TPB, typename CountIfOp>
+inline int count_if(T* arr, int size, CountIfOp count_if_op) {
     int *count_d, count_h, N_BLK;
 
-    common::_prep_count_if(&count_d, N_BLK, size);
+    common::_prep_count_if<TPB> (&count_d, N_BLK, size);
 
-    count_kernel<<<N_BLK, TPB>>> (arr, size, count_d, count_if_op);
+    std:: cout << "TPB: " << TPB << "\n";
+    std:: cout << "N_BLK: " << N_BLK << "\n";
+    std:: cout << "size: " << size << "\n";
+    detail::count_kernel<<<N_BLK, TPB>>> (arr, size, count_d, count_if_op);
 
     common::_finish_count_if(&count_d, count_h);
 
@@ -38,14 +41,15 @@ int count_if(T* arr, int size,
 
 namespace manual_reduction {
 
-template <typename T, class F>
-int count_if(T* arr, int size,
-             F count_if_op) {
+template <typename T, int TPB, typename CountIfOp>
+inline int count_if(T* arr, int size,
+                    CountIfOp count_if_op) {
     int *count_d, count_h, N_BLK;
 
-    common::_prep_count_if(&count_d, N_BLK, size);
+    common::_prep_count_if<TPB> (&count_d, N_BLK, size);
 
-    count_kernel<<<N_BLK, TPB>>> (arr, size, count_d, count_if_op);
+    detail::count_kernel<T, TPB> <<<N_BLK, TPB>>> (arr, size, count_d,
+                                                   count_if_op);
 
     common::_finish_count_if(&count_d, count_h);
 
@@ -56,14 +60,14 @@ int count_if(T* arr, int size,
 
 namespace syncthreads_count_reduction {
 
-template <typename T, class F>
-int count_if(T* arr, int size,
-                F count_if_op) {
+template <typename T, int TPB, typename CountIfOp>
+inline int count_if(T* arr, int size,
+                    CountIfOp count_if_op) {
     int *count_d, count_h, N_BLK;
 
-    common::_prep_count_if(&count_d, N_BLK, size);
+    common::_prep_count_if<TPB> (&count_d, N_BLK, size);
 
-    count_kernel<<<N_BLK, TPB>>> (arr, size, count_d, count_if_op);
+    detail::count_kernel<<<N_BLK, TPB>>> (arr, size, count_d, count_if_op);
 
     common::_finish_count_if(&count_d, count_h);
 
@@ -74,14 +78,15 @@ int count_if(T* arr, int size,
 
 namespace ballot_sync_reduction {
 
-template <typename T, class F>
-int count_if(T* arr, int size,
-                F count_if_op) {
+template <typename T, int TPB, typename CountIfOp>
+inline int count_if(T* arr, int size,
+                    CountIfOp count_if_op) {
     int *count_d, count_h, N_BLK;
 
-    common::_prep_count_if(&count_d, N_BLK, size);
+    common::_prep_count_if<TPB> (&count_d, N_BLK, size);
 
-    count_kernel<<<N_BLK, TPB>>> (arr, size, count_d, count_if_op);
+    detail::count_kernel<T, TPB> <<<N_BLK, TPB>>> (arr, size, count_d,
+                                                   count_if_op);
 
     common::_finish_count_if(&count_d, count_h);
 

cpp/include/learning_cuda/count/count_kernels.cuh

@@ -1,29 +1,35 @@
-// #include <cuda_runtime.h>
+#pragma once
+
+#include <cuda_runtime.h>
 
 namespace naive {
+namespace detail {
 
-template <typename T, class F>
+template <typename T, typename CountIfOp>
 __global__ 
-void count_kernel(T* arr, int size, int *count, F count_if_op) {
+void count_kernel(T* arr, int size, int *count, CountIfOp count_if_op) {
     int tid = threadIdx.x + blockIdx.x * blockDim.x;
 
     if (tid < size) {
+
         if (count_if_op(arr[tid])) {
             atomicAdd(count, 1);
         }
     }
 }
 
+} // namespace detail
 } // namespace naive
 
 namespace manual_reduction {
+namespace detail {
 
-template <typename T, class F>
+template <typename T, int TPB, typename CountIfOp>
 __global__ 
-void count_kernel(T* arr, int size, int *count, F count_if_op) {
+void count_kernel(T* arr, int size, int *count, CountIfOp count_if_op) {
     int tid = threadIdx.x + blockIdx.x * blockDim.x;
 
-    __shared__ int local_count_array[32];
+    __shared__ int local_count_array[TPB];
 
     if (tid < size) {
         if (count_if_op(arr[tid])) {
@@ -47,14 +53,16 @@ void count_kernel(T* arr, int size, int *count, F count_if_op) {
         atomicAdd(count, local_count_array[threadIdx.x]);
     }
 }
-    
+
+} // namespace detail
 } // namespace manual_reduction
 
 namespace syncthreads_count_reduction {
+namespace detail {
 
-template <typename T, class F>
+template <typename T, typename CountIfOp>
 __global__ 
-void count_kernel(T* arr, int size, int *count, F count_if_op) {
+void count_kernel(T* arr, int size, int *count, CountIfOp count_if_op) {
     int tid = threadIdx.x + blockIdx.x * blockDim.x;
 
     bool predicate = tid < size && count_if_op(arr[tid]);
@@ -67,15 +75,17 @@ void count_kernel(T* arr, int size, int *count, F count_if_op) {
 
 }
 
+} // namespace detail
 } // namespace syncthreads_count
 
 #define FULL_MASK 0xffffffff
 
 namespace ballot_sync_reduction {
+namespace detail {
 
-template <typename T, class F>
+template <typename T, int TPB, typename CountIfOp>
 __global__ 
-void count_kernel(T* arr, int size, int *count, F count_if_op) {
+void count_kernel(T* arr, int size, int *count, CountIfOp count_if_op) {
     int tid = threadIdx.x + blockIdx.x * blockDim.x;
 
     bool predicate = tid < size && count_if_op(arr[tid]);
@@ -89,4 +99,5 @@ void count_kernel(T* arr, int size, int *count, F count_if_op) {
 
 }
 
-} // ballot syncthreads_count
+} // namespace detail
+} // namespace ballot_sync_reduction

cpp/src/count/count.cu

@@ -12,7 +12,7 @@
 
 
 int main() {
-    int n_elems = 1000;
+    int n_elems = 100;
 
     thrust::device_vector<int> rand_d(n_elems, 0);
 
@@ -33,23 +33,27 @@ int main() {
 
     std::cout << "\nThrust Count: " << thrust_count << std::endl;
 
-    int naive_count = naive::count_if(thrust::raw_pointer_cast(rand_d.data()),
-                                      n_elems, is_greater_than_10);
+    int *rand_d_ptr = thrust::raw_pointer_cast(rand_d.data());
+
+    int naive_count = naive::count_if<int, 32>(rand_d_ptr, n_elems,
+                                               is_greater_than_10);
 
     std::cout << "\nNaive Count: " << naive_count << std::endl;
 
-    int man_count = manual_reduction::count_if(thrust::raw_pointer_cast(rand_d.data()),
-    n_elems, is_greater_than_10);
+    int man_count = manual_reduction::count_if<int, 32>(rand_d_ptr, n_elems,
+                                                        is_greater_than_10);
 
     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);
+    int syn_count = syncthreads_count_reduction::count_if<int, 32>(rand_d_ptr,
+                                                                   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);
+    int bal_count = ballot_sync_reduction::count_if<int, 32>(rand_d_ptr, 
+                                                             n_elems,
+                                                             is_greater_than_10);
 
     std::cout << "\nBallot Reduction Count: " << bal_count << std::endl;