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

[hliao]

Summary

• 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.

[tra]

LGTM for CUDA.

clang/lib/CodeGen/CodeGenPGO.cpp

839–840

We will still have around some functions that may never be used on the host side (HD functions referenced from device code only). I'm not sure if that's a problem for profiling, though. I wonder if we can somehow tie skipRegionMappingForDecl to whether we've actually codegen'ed the function.

[hliao]

clang/lib/CodeGen/CodeGenPGO.cpp

839–840

Skipping wrong-side functions here just makes the report not confusing as these functions are not emitted at all and are supposed never running on the host/device side. If we still create the mapping for them, e.g., we may report they have 0 runs instead of reporting nothing (just like comments between function.) That looks a little bit confusing.
It seems the current PGO adds everything for coverage mapping and late prune them based on checks here. Just try to follow that logic to skip wrong-side functions. If we need to revise the original logic and generate coverage mapping for emitted functions only, the change here is unnecessary.

[tra]

clang/lib/CodeGen/CodeGenPGO.cpp

839–840

I'd add a comment here that this 'filter' is just a rough best-effort approximation that still allows some effectively device-only Decls through.
The output should still be correct, even though the functions will never be used. Maybe add a TODO to deal with it if/when we know if the Decl was codegen'ed.

[yaxunl]

Do we need to disable pgo and coverage mapping for device compilation? Or it is already disabled?

[tra]

In D85276#2199655, @yaxunl wrote:

Do we need to disable pgo and coverage mapping for device compilation? Or it is already disabled?

We already disable profiling during device compilation for NVIDIA and AMD GPUs:
https://github.com/llvm/llvm-project/blob/394db2259575ef3cac8d3d37836b11eb2373c435/clang/lib/Driver/ToolChains/Clang.cpp#L4876

[hliao]

In D85276#2200108, @tra wrote:

In D85276#2199655, @yaxunl wrote:

Do we need to disable pgo and coverage mapping for device compilation? Or it is already disabled?

We already disable profiling during device compilation for NVIDIA and AMD GPUs:
https://github.com/llvm/llvm-project/blob/394db2259575ef3cac8d3d37836b11eb2373c435/clang/lib/Driver/ToolChains/Clang.cpp#L4876

Anyway, this patch just fixes the caused by that device stub function. As it's "emitted" in the host compilation, we need to skip generating instrumentation on it explicitly.

[yaxunl]

LGTM. thanks

[hliao]

Revise the comment.

[hliao]

clang/lib/CodeGen/CodeGenPGO.cpp

839–840

Add that comment. But, I tend to not deal that "effectively" host-only/device-only ones as that should be developers' responsibility to handle them. The additional zero coverage mapping may be useful as well. If a function is really device-only but is attributed with HD, the 0 coverage may help developers correcting them.

[tra]

clang/lib/CodeGen/CodeGenPGO.cpp

839–840

It will be rather noisy in practice. A lot of code has either has been written for NVCC or has to compile with it. NVCC does not have target overloads, so sticking HD everywhere is pretty much the only practical way to do it in complicated enough C++ code. Anything that uses Eigen or Thrust will have tons of HD functions that are actually used only on one side.

[hliao]

clang/lib/CodeGen/CodeGenPGO.cpp

839–840

Most HD interfaces in Eigen are designed to be used in both CPU and GPU. For GPU only ones, they are marked with __device__ only. Thrust has a similar situation.

clang/lib/CodeGen/CodeGenPGO.cpp

uint64_t PGOHash::finalize() {
   return Result.low();
 }
 
 void CodeGenPGO::assignRegionCounters(GlobalDecl GD, llvm::Function *Fn) {
   const Decl *D = GD.getDecl();
   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)
     return;
   if (D->isImplicit())
     return;
   // Constructors and destructors may be represented by several functions in IR.
   // If so, instrument only base variant, others are implemented by delegation
void CodeGenPGO::mapRegionCounters(const Decl *D) {
   NumRegionCounters = Walker.NextCounter;
   FunctionHash = Walker.Hash.finalize();
 }
 
 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();
   return SM.isInSystemHeader(Loc);
 }
 
 void CodeGenPGO::emitCounterRegionMapping(const Decl *D) {
   if (skipRegionMappingForDecl(D))

clang/test/CodeGenCUDA/profile-coverage-mapping.cu

+// 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);
+}