diff --git a/clang/lib/CodeGen/CodeGenPGO.cpp b/clang/lib/CodeGen/CodeGenPGO.cpp --- a/clang/lib/CodeGen/CodeGenPGO.cpp +++ b/clang/lib/CodeGen/CodeGenPGO.cpp @@ -773,6 +773,11 @@ if (!D->hasBody()) return; + // Skip CUDA/HIP kernel launch stub functions. + if (CGM.getLangOpts().CUDA && !CGM.getLangOpts().CUDAIsDevice && + D->hasAttr()) + return; + bool InstrumentRegions = CGM.getCodeGenOpts().hasProfileClangInstr(); llvm::IndexedInstrProfReader *PGOReader = CGM.getPGOReader(); if (!InstrumentRegions && !PGOReader) @@ -831,6 +836,16 @@ if (!D->getBody()) return true; + // Skip host-only functions in the CUDA device compilation and device-only + // functions in the host compilation. + if (CGM.getLangOpts().CUDA && + ((CGM.getLangOpts().CUDAIsDevice && !D->hasAttr() && + !D->hasAttr()) || + (!CGM.getLangOpts().CUDAIsDevice && + (D->hasAttr() || + (!D->hasAttr() && D->hasAttr()))))) + return true; + // Don't map the functions in system headers. const auto &SM = CGM.getContext().getSourceManager(); auto Loc = D->getBody()->getBeginLoc(); diff --git a/clang/test/CodeGenCUDA/profile-coverage-mapping.cu b/clang/test/CodeGenCUDA/profile-coverage-mapping.cu new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenCUDA/profile-coverage-mapping.cu @@ -0,0 +1,20 @@ +// RUN: echo "GPU binary would be here" > %t +// RUN: %clang_cc1 -fprofile-instrument=clang -triple x86_64-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm -o - %s | FileCheck --check-prefix=PGOGEN %s +// RUN: %clang_cc1 -fprofile-instrument=clang -fcoverage-mapping -triple x86_64-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm -o - %s | FileCheck --check-prefix=COVMAP %s +// RUN: %clang_cc1 -fprofile-instrument=clang -fcoverage-mapping -dump-coverage-mapping -triple x86_64-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm-only -o - %s | FileCheck --check-prefix=MAPPING %s + +#include "Inputs/cuda.h" + +// PGOGEN-NOT: @__profn_{{.*kernel.*}} = +// COVMAP-COUNT-2: section "__llvm_covfun", comdat +// COVMAP-NOT: section "__llvm_covfun", comdat +// MAPPING-NOT: {{.*dfn.*}}: +// MAPPING-NOT: {{.*kernel.*}}: + +__device__ void dfn(int i) {} + +__global__ void kernel(int i) { dfn(i); } + +void host(void) { + kernel<<<1, 1>>>(1); +}