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

[HIPIFY] Disable cudaBuiltin matcher.

Browse Source 
As HIP has started to support vanilla CUDA syntax for threadIdx, blockIdx, blockDim and gridDim.
Other CUDA builtins are not tracked for now.
This commit is contained in:
Evgeny Mankov
2017-12-05 20:28:51 +03:00
parent f24dfc6f36
commit 71d2fb20c8
2 changed files with 4 additions and 19 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipify-clang/src/CUDA2HipMap.cpp
-17
View File
@@ -1461,23 +1461,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
{"cudaHostRegisterMapped", {"hipHostRegisterMapped", CONV_MEM, API_RUNTIME}}, {"cudaHostRegisterMapped", {"hipHostRegisterMapped", CONV_MEM, API_RUNTIME}},
{"cudaHostRegisterIoMemory", {"hipHostRegisterIoMemory", CONV_MEM, API_RUNTIME}}, {"cudaHostRegisterIoMemory", {"hipHostRegisterIoMemory", CONV_MEM, API_RUNTIME}},
// Coordinate Indexing and Dimensions
{"threadIdx.x", {"hipThreadIdx_x", CONV_COORD_FUNC, API_RUNTIME}},
{"threadIdx.y", {"hipThreadIdx_y", CONV_COORD_FUNC, API_RUNTIME}},
{"threadIdx.z", {"hipThreadIdx_z", CONV_COORD_FUNC, API_RUNTIME}},
{"blockIdx.x", {"hipBlockIdx_x", CONV_COORD_FUNC, API_RUNTIME}},
{"blockIdx.y", {"hipBlockIdx_y", CONV_COORD_FUNC, API_RUNTIME}},
{"blockIdx.z", {"hipBlockIdx_z", CONV_COORD_FUNC, API_RUNTIME}},
{"blockDim.x", {"hipBlockDim_x", CONV_COORD_FUNC, API_RUNTIME}},
{"blockDim.y", {"hipBlockDim_y", CONV_COORD_FUNC, API_RUNTIME}},
{"blockDim.z", {"hipBlockDim_z", CONV_COORD_FUNC, API_RUNTIME}},
{"gridDim.x", {"hipGridDim_x", CONV_COORD_FUNC, API_RUNTIME}},
{"gridDim.y", {"hipGridDim_y", CONV_COORD_FUNC, API_RUNTIME}},
{"gridDim.z", {"hipGridDim_z", CONV_COORD_FUNC, API_RUNTIME}},
{"warpSize", {"hipWarpSize", CONV_SPECIAL_FUNC, API_RUNTIME}}, {"warpSize", {"hipWarpSize", CONV_SPECIAL_FUNC, API_RUNTIME}},
// Events // Events
hipify-clang/src/HipifyAction.cpp
+4 -2
View File
@@ -358,6 +358,8 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this); Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this);
// Replace cuda builtins. // Replace cuda builtins.
/* Disable as HIP has started to support vanilla CUDA syntax for threadIdx, blockIdx, blockDim and gridDim. */
/* Other CUDA builtins are not tracked for now.
Finder->addMatcher( Finder->addMatcher(
mat::memberExpr( mat::memberExpr(
mat::isExpansionInMainFile(), mat::isExpansionInMainFile(),
@@ -371,7 +373,7 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
).bind("cudaBuiltin"), ).bind("cudaBuiltin"),
this this
); );
*/
Finder->addMatcher( Finder->addMatcher(
mat::varDecl( mat::varDecl(
mat::isExpansionInMainFile(), mat::isExpansionInMainFile(),
@@ -454,7 +456,7 @@ void HipifyAction::ExecuteAction() {
} }
void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) { void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
if (cudaBuiltin(Result)) return; // if (cudaBuiltin(Result)) return;
if (cudaLaunchKernel(Result)) return; if (cudaLaunchKernel(Result)) return;
if (cudaSharedIncompleteArrayVar(Result)) return; if (cudaSharedIncompleteArrayVar(Result)) return;
} }