diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp index ff16b188..9258bca9 100644 --- a/src/HipifyAction.cpp +++ b/src/HipifyAction.cpp @@ -2660,6 +2660,7 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) auto it = FuncArgCasts.find(sName); if (it == FuncArgCasts.end()) return false; auto castStructs = it->second; + auto &SM = *Result.SourceManager; for (auto cc : castStructs) { if (cc.isToMIOpen != TranslateToMIOpen || cc.isToRoc != TranslateToRoc) continue; clang::LangOptions DefaultLangOptions; @@ -2668,17 +2669,17 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) unsigned int argNum = c.first; clang::SmallString<40> XStr; llvm::raw_svector_ostream OS(XStr); - auto &SM = *Result.SourceManager; clang::SourceRange sr, replacementRange; clang::SourceLocation s, e; if (argNum < call->getNumArgs()) { sr = call->getArg(argNum)->getSourceRange(); replacementRange = getWriteRange(SM, { sr.getBegin(), sr.getEnd() }); - s = replacementRange.getBegin(); - e = replacementRange.getEnd(); } else { s = e = call->getEndLoc(); + replacementRange = getWriteRange(SM, { s, e }); } + s = replacementRange.getBegin(); + e = replacementRange.getEnd(); switch (c.second.castType) { case e_remove_argument: { @@ -2695,6 +2696,8 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) s = prevComma->getLocation(); } } + replacementRange = getWriteRange(SM, { s, e }); + e = replacementRange.getEnd(); length = SM.getCharacterData(e) - SM.getCharacterData(s); break; } @@ -2727,6 +2730,8 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) e = call->getArg(argNum + c.second.numberToMoveOrCopy)->getBeginLoc(); else e = call->getEndLoc(); + replacementRange = getWriteRange(SM, { s, e }); + e = replacementRange.getEnd(); length = SM.getCharacterData(e) - SM.getCharacterData(s); break; } @@ -2736,8 +2741,6 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) OS << c.second.constValToAddOrReplace << ", "; else OS << ", " << c.second.constValToAddOrReplace; - clang::SourceRange replacementRange = getWriteRange(*Result.SourceManager, { s, s }); - s = replacementRange.getBegin(); break; } case e_add_var_argument: diff --git a/tests/unit_tests/samples/cudaMallocHost.cu b/tests/unit_tests/synthetic/transforming_matchers.cu similarity index 67% rename from tests/unit_tests/samples/cudaMallocHost.cu rename to tests/unit_tests/synthetic/transforming_matchers.cu index 13f85ec8..9ade8811 100644 --- a/tests/unit_tests/samples/cudaMallocHost.cu +++ b/tests/unit_tests/synthetic/transforming_matchers.cu @@ -15,6 +15,7 @@ void check(T result, char const *const func, const char *const file, int const l int main(int argc, const char *argv[]) { int *input = nullptr; + void *input_ptr = nullptr; int deviceCount = 0; // CHECK: checkErrors(hipGetDeviceCount(&deviceCount)); checkErrors(cudaGetDeviceCount(&deviceCount)); @@ -30,9 +31,25 @@ int main(int argc, const char *argv[]) { checkErrors(cudaSetDevice(deviceID)); // CHECK: checkErrors(hipHostMalloc(&input, sizeof(int) * num * 2, hipHostMallocDefault)); checkErrors(cudaMallocHost(&input, sizeof(int) * num * 2)); + // CHECK: checkErrors(hipHostMalloc(&input_ptr, sizeof(int) * num * 2, hipHostMallocDefault)); + checkErrors(cudaMallocHost(&input_ptr, sizeof(int) * num * 2)); for (int i = 0; i < num * 2; ++i) { input[i] = i; } + + int *value = 0; + int *value_2 = 0; + int iBlockSize = 0; + int iBlockSize_2 = 0; + size_t bytes = 0; + // CHECK: hipFunction_t function; + CUfunction function; + // CHECK: void* occupancyB2DSize; + CUoccupancyB2DSize occupancyB2DSize; + + // CHECK: checkErrors(hipModuleOccupancyMaxPotentialBlockSizeWithFlags(value, value_2, function, bytes, iBlockSize, iBlockSize_2)); + checkErrors(cuOccupancyMaxPotentialBlockSizeWithFlags(value, value_2, function, occupancyB2DSize, bytes, iBlockSize, iBlockSize_2)); + // CHECK: checkErrors(hipHostFree(input)); checkErrors(cudaFreeHost(input)); // CHECK: checkErrors(hipDeviceSynchronize());