AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge pull request #1624 from emankov/hipify

Browse Source 
[HIPIFY][#1409] Fix for kernel launch macro expansion

[ROCm/clr commit: cf5d1caab3]
This commit is contained in:
Evgeny Mankov
2019-11-05 14:04:51 +03:00
committed by GitHub
parent 92393d6ecf 3889be3031
commit ecad81b794
2 changed files with 64 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
+18 -10
View File
@@ -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<clang::SourceLocation> 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) {
projects/clr/hipamd/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu
+46
View File
@@ -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 <hip/hip_runtime.h>
#include <algorithm>
#define CUDA_LAUNCH(cuda_call,dimGrid,dimBlock, ...) \
cuda_call<<<dimGrid,dimBlock>>>(__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);
}