[HIPIFY][CUB][#1460] Add "using namespace cub" translation support

+ Add cub_03.cu
2019-10-18 18:51:40 +03:00
parent f0ed210b19
commit 98874c0e7f
3 changed files with 57 additions and 0 deletions
@@ -0,0 +1,37 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+// CHECK: #include <hip/hip_runtime.h>
+#include <iostream>
+// CHECK: #include <hiprand.h>
+#include <curand.h>
+// CHECK: #include <hipcub/hipcub.hpp>
+#include <cub/cub.cuh>
+
+#include <iostream>
+
+// using namespace hipcub;
+using namespace cub;
+
+/**
+* Simple CUDA kernel for computing tiled partial sums
+*/
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD,
+          BlockLoadAlgorithm LOAD_ALGO,
+          BlockScanAlgorithm SCAN_ALGO>
+__global__ void ScanTilesKernel(int *d_in, int *d_out) {
+  // Specialize collective types for problem context
+  // TODO: typedef cub::BlockLoad<int*, BLOCK_THREADS, ITEMS_PER_THREAD, LOAD_ALGO> BlockLoadT;
+  typedef BlockLoad<int*, BLOCK_THREADS, ITEMS_PER_THREAD, LOAD_ALGO> BlockLoadT;
+  typedef BlockScan<int, BLOCK_THREADS, SCAN_ALGO> BlockScanT;
+  // Allocate on-chip temporary storage
+  __shared__ union {
+    typename BlockLoadT::TempStorage load;
+    typename BlockScanT::TempStorage reduce;
+  } temp_storage;
+  // Load data per thread
+  int thread_data[ITEMS_PER_THREAD];
+  int offset = blockIdx.x * (BLOCK_THREADS * ITEMS_PER_THREAD);
+  BlockLoadT(temp_storage.load).Load(d_in + offset, offset);
+  __syncthreads();
+  // Compute the block-wide prefix sum
+  BlockScanT(temp_storage).Sum(thread_data);
+}