AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[HIPIFY] Disable delayed template parsing

Browse Source 
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.
This commit is contained in:
Evgeny Mankov
2019-10-22 19:07:37 +03:00
parent 76c8406449
commit b6e6f12b54
2 changed files with 7 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipify-clang/src/main.cpp
+1
View File
@@ -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));
tests/hipify-clang/unit_tests/libraries/CUB/cub_03.cu
+6 -6
View File
@@ -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 {