[HIPIFY] Disable cudaBuiltin matcher.

As HIP has started to support vanilla CUDA syntax for threadIdx, blockIdx, blockDim and gridDim.
Other CUDA builtins are not tracked for now.
Esse commit está contido em:
Evgeny Mankov committed 2017-12-05 20:28:51 +03:00
1 pai 43f0a14005
commit ccc959fdc9
2 files changed
+4 -19

No files matched your search

-17
Ver Arquivo
@@ -1461,23 +1461,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
{"cudaHostRegisterMapped", {"hipHostRegisterMapped", 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}},
// Events
+4 -2
Ver Arquivo
@@ -358,6 +358,8 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this);
// 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(
mat::memberExpr(
mat::isExpansionInMainFile(),
@@ -371,7 +373,7 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
).bind("cudaBuiltin"),
this
);
*/
Finder->addMatcher(
mat::varDecl(
mat::isExpansionInMainFile(),
@@ -454,7 +456,7 @@ void HipifyAction::ExecuteAction() {
}
void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
if (cudaBuiltin(Result)) return;
// if (cudaBuiltin(Result)) return;
if (cudaLaunchKernel(Result)) return;
if (cudaSharedIncompleteArrayVar(Result)) return;
}