From d39793f0f740e7f25b618687bd508aedb260cfcb Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Tue, 22 Oct 2019 19:07:37 +0300
Subject: [PATCH] [HIPIFY] Disable delayed template parsing

By implicit unconditional passing -fno-delayed-template-parsing option (which appeared in LLVM 3.8.0, thus doesn't need compatibility wrapping) to hipify-clang.

[Reason] To parse uncalled template functions otherwise they are not parsed without calling, thus not hipified.

Affects cub_03.cu test, which has uncalled global template function.


[ROCm/hip commit: b6e6f12b546a0ebca0c35ae7efeafa0ca4524cdd]
---
 projects/hip/hipify-clang/src/main.cpp               |  1 +
 .../hipify-clang/unit_tests/libraries/CUB/cub_03.cu  | 12 ++++++------
 2 files changed, 7 insertions(+), 6 deletions(-)

diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp
index 2214567df3..64037c43dd 100644
--- a/projects/hip/hipify-clang/src/main.cpp
+++ b/projects/hip/hipify-clang/src/main.cpp
@@ -199,6 +199,7 @@ int main(int argc, const char **argv) {
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN));
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN));
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
+    Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-fno-delayed-template-parsing", ct::ArgumentInsertPosition::BEGIN));
     if (!CudaPath.empty()) {
       std::string sCudaPath = "--cuda-path=" + CudaPath;
       Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN));
diff --git a/projects/hip/tests/hipify-clang/unit_tests/libraries/CUB/cub_03.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/CUB/cub_03.cu
index 9fdbc17515..bc914d419d 100644
--- a/projects/hip/tests/hipify-clang/unit_tests/libraries/CUB/cub_03.cu
+++ b/projects/hip/tests/hipify-clang/unit_tests/libraries/CUB/cub_03.cu
@@ -1,8 +1,6 @@
 // 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>
 
@@ -11,12 +9,14 @@ 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>
+          // CHECK: hipcub::BlockLoadAlgorithm LOAD_ALGO,
+          cub::BlockLoadAlgorithm LOAD_ALGO,
+          // CHECK: hipcub::BlockScanAlgorithm SCAN_ALGO>
+          cub::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;
+  // CHECK: typedef ::hipcub::BlockLoad<int*, BLOCK_THREADS, ITEMS_PER_THREAD, LOAD_ALGO> BlockLoadT;
+  typedef ::cub::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 {