diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 0413d76022..1893f6ce5b 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -1461,23 +1461,6 @@ const std::map 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 diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp index 192dd00949..f4c7fd7764 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp @@ -358,6 +358,8 @@ std::unique_ptr 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 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; }