From 3889be30316822ca2599ed822d683d43cab3195e Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 5 Nov 2019 14:00:13 +0300 Subject: [PATCH] [HIPIFY][#1409] Fix for kernel launch macro expansion + Add a corresponding test kernel_launch_01.cu + Add isBefore() check to avoid crash on Replacement with negative length TODO: + Compatibility with former LLVM versions + More complicated kernel launch tests [ROCm/clr commit: a1c380f38df84e75f840c88c07787c560bf5c628] --- .../hipamd/hipify-clang/src/HipifyAction.cpp | 28 +++++++---- .../kernel_launch/kernel_launch_01.cu | 46 +++++++++++++++++++ 2 files changed, 64 insertions(+), 10 deletions(-) create mode 100644 projects/clr/hipamd/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp index 510d91978a..0e776df63e 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp @@ -198,7 +198,7 @@ clang::SourceRange getWriteRange(clang::SourceManager &SM, const clang::SourceRa // If the range is contained within a macro, update the macro definition. // Otherwise, use the file location and hope for the best. if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) { - return {SM.getFileLoc(begin), SM.getFileLoc(end)}; + return {SM.getExpansionLoc(begin), SM.getExpansionLoc(end)}; } return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; } @@ -390,16 +390,24 @@ bool HipifyAction::cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result) OS << readSourceText(*SM, {argStart, argEnd}); } OS << ")"; - clang::SourceRange replacementRange = getWriteRange(*SM, {llcompat::getBeginLoc(launchKernel), llcompat::getEndLoc(launchKernel)}); - clang::SourceLocation launchStart = replacementRange.getBegin(); + clang::SourceLocation launchKernelExprLocBeg = launchKernel->getExprLoc(); + clang::SourceLocation launchKernelExprLocEnd = launchKernelExprLocBeg.isMacroID() ? SM->getExpansionRange(launchKernelExprLocBeg).getEnd() : llcompat::getEndLoc(launchKernel); + clang::SourceLocation launchKernelEnd = llcompat::getEndLoc(launchKernel); + clang::BeforeThanCompare isBefore(*SM); + launchKernelExprLocEnd = isBefore(launchKernelEnd, launchKernelExprLocEnd) ? launchKernelExprLocEnd : launchKernelEnd; + clang::SourceRange replacementRange = getWriteRange(*SM, {launchKernelExprLocBeg, launchKernelExprLocEnd}); + clang::SourceLocation launchBeg = replacementRange.getBegin(); clang::SourceLocation launchEnd = replacementRange.getEnd(); - size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart); - ct::Replacement Rep(*SM, launchStart, length, OS.str()); - clang::FullSourceLoc fullSL(launchStart, *SM); - insertReplacement(Rep, fullSL); - hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME}; - Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str()); - return true; + if (isBefore(launchBeg, launchEnd)) { + size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchBeg); + ct::Replacement Rep(*SM, launchBeg, length, OS.str()); + clang::FullSourceLoc fullSL(launchBeg, *SM); + insertReplacement(Rep, fullSL); + hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME}; + Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str()); + return true; + } + return false; } bool HipifyAction::cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result) { diff --git a/projects/clr/hipamd/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu new file mode 100644 index 0000000000..3795d3c799 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu @@ -0,0 +1,46 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// Synthetic test to warn only on device functions umin and umax as unsupported, but not on user defined ones. +// ToDo: change lit testing in order to parse the output. + +#define LEN 1024 +#define SIZE LEN * sizeof(float) +#define ITER 1024*1024 + +// CHECK: #include +#include + +#define CUDA_LAUNCH(cuda_call,dimGrid,dimBlock, ...) \ + cuda_call<<>>(__VA_ARGS__); + +__global__ void Inc1(float *Ad, float *Bd) { + int tx = threadIdx.x + blockIdx.x * blockDim.x; + if (tx < 1) { + for (int i = 0; i < ITER; ++i) { + Ad[tx] = Ad[tx] + 1.0f; + for (int j = 0; j < 256; ++j) { + Bd[tx] = Ad[tx]; + } + } + } +} + +int main() { + float *A, *Ad, *Bd; + A = new float[LEN]; + for (int i = 0; i < LEN; ++i) { + A[i] = 0.0f; + } + // CHECK: hipError_t status; + cudaError_t status; + // CHECK: status = hipHostRegister(A, SIZE, hipHostRegisterMapped); + status = cudaHostRegister(A, SIZE, cudaHostRegisterMapped); + // CHECK: hipHostGetDevicePointer(&Ad, A, 0); + cudaHostGetDevicePointer(&Ad, A, 0); + // CHECK: hipMalloc((void**)&Bd, SIZE); + cudaMalloc((void**)&Bd, SIZE); + dim3 dimGrid(LEN / 512, 1, 1); + dim3 dimBlock(512, 1, 1); + + // CHECK: hipLaunchKernelGGL(Inc1, dim3(dimGrid), dim3(dimBlock), 0, 0, Ad, Bd); + CUDA_LAUNCH(Inc1, dimGrid, dimBlock, Ad, Bd); +}