[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.
[ROCm/clr commit: ccc959fdc9]
Этот коммит содержится в:
1 родитель
e3e37db0d7
Коммит
d8f8cc2838
2 files changed
+4
-19
No files matched your search
@@ -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
|
||||||
|
|||||||
@@ -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;
|
||||||
}
|
}
|
||||||
Reference in new issue
Block a user