maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[PGO][CUDA][HIP] Skip generating profile on the device stub and wrong-side functions. 

- Skip generating profile data on `__global__` function in the host
  compilation. It's a host-side stub function only and don't have
  profile instrumentation generated on the real function body. The extra
  profile data results in the malformed instrumentation profile data.
- Skip generating region mapping on functions in the wrong-side, i.e.,
  + For the device compilation, skip host-only functions; and,
  + For the host compilation, skip device-only functions (including
    `__global__` functions.)
- As the device-side profiling is not ready yet, only host-side profile
  code generation is checked.

Differential Revision: https://reviews.llvm.org/D85276
This commit is contained in:
Michael Liao 2020-08-04 13:56:36 -04:00
parent 7866442b3f
commit c7b683c126
2 changed files with 37 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

17
clang/lib/CodeGen/CodeGenPGO.cpp
View File

@ -773,6 +773,11 @@ void CodeGenPGO::assignRegionCounters(GlobalDecl GD, llvm::Function *Fn) {
if (!D->hasBody())
return;
// Skip CUDA/HIP kernel launch stub functions.
if (CGM.getLangOpts().CUDA && !CGM.getLangOpts().CUDAIsDevice &&
D->hasAttr<CUDAGlobalAttr>())
return;
bool InstrumentRegions = CGM.getCodeGenOpts().hasProfileClangInstr();
llvm::IndexedInstrProfReader *PGOReader = CGM.getPGOReader();
if (!InstrumentRegions && !PGOReader)
@ -831,6 +836,18 @@ bool CodeGenPGO::skipRegionMappingForDecl(const Decl *D) {
if (!D->getBody())
return true;
// Skip host-only functions in the CUDA device compilation and device-only
// functions in the host compilation. Just roughly filter them out based on
// the function attributes. If there are effectively host-only or device-only
// ones, their coverage mapping may still be generated.
if (CGM.getLangOpts().CUDA &&
((CGM.getLangOpts().CUDAIsDevice && !D->hasAttr<CUDADeviceAttr>() &&
!D->hasAttr<CUDAGlobalAttr>()) ||
(!CGM.getLangOpts().CUDAIsDevice &&
(D->hasAttr<CUDAGlobalAttr>() ||
(!D->hasAttr<CUDAHostAttr>() && D->hasAttr<CUDADeviceAttr>())))))
return true;
// Don't map the functions in system headers.
const auto &SM = CGM.getContext().getSourceManager();
auto Loc = D->getBody()->getBeginLoc();

20
clang/test/CodeGenCUDA/profile-coverage-mapping.cu Normal file
View File

@ -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);
}