Author: Yaxun (Sam) Liu Date: 2026-05-27T23:29:05-04:00 New Revision: 5acb952cb5d9c899692f44f4280af75a780a599c
URL: https://github.com/llvm/llvm-project/commit/5acb952cb5d9c899692f44f4280af75a780a599c DIFF: https://github.com/llvm/llvm-project/commit/5acb952cb5d9c899692f44f4280af75a780a599c.diff LOG: [HIP][Driver] Forward -fcoverage-mapping flags to device compiler (#198872) Add `-fcoverage-mapping`, `-fno-coverage-mapping`, `-fcoverage-compilation-dir=`, `-ffile-compilation-dir=`, and `-fcoverage-prefix-map=` to the LinkerWrapper `CompilerOptions` forwarding list. Without this, passing `-fprofile-instr-generate -fcoverage-mapping` to clang for a HIP program silently omits the coverage mapping flags from the embedded device recompilation, so `__llvm_covmap`/`__llvm_covfun` sections are never emitted for device code. Added: clang/test/CodeGenHIP/profile-coverage-mapping.hip Modified: clang/lib/Driver/ToolChains/Clang.cpp clang/test/Driver/hip-options.hip Removed: ################################################################################ diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 1bdcbc4e7620b..05e1f6db80a11 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -9502,6 +9502,39 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs, Inputs, Output)); } +// Options that need the profile compiler-rt library on the target toolchain. +// Coverage mapping flags require -fprofile-instr-generate, so they belong here +// too. +static bool requiresProfileRT(unsigned ID) { + switch (ID) { + case options::OPT_fprofile_generate: + case options::OPT_fprofile_generate_EQ: + case options::OPT_fprofile_instr_generate: + case options::OPT_fprofile_instr_generate_EQ: + case options::OPT_fcoverage_mapping: + case options::OPT_fno_coverage_mapping: + case options::OPT_fcoverage_compilation_dir_EQ: + case options::OPT_ffile_compilation_dir_EQ: + case options::OPT_fcoverage_prefix_map_EQ: + return true; + default: + return false; + } +} + +// Options that need the ubsan compiler-rt library on the target toolchain. +static bool requiresUBSanRT(unsigned ID) { + switch (ID) { + case options::OPT_fsanitize_EQ: + case options::OPT_fno_sanitize_EQ: + case options::OPT_fsanitize_minimal_runtime: + case options::OPT_fno_sanitize_minimal_runtime: + return true; + default: + return false; + } +} + void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, const InputInfo &Output, const InputInfoList &Inputs, @@ -9549,6 +9582,11 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, OPT_fprofile_generate_EQ, OPT_fprofile_instr_generate, OPT_fprofile_instr_generate_EQ, + OPT_fcoverage_mapping, + OPT_fno_coverage_mapping, + OPT_fcoverage_compilation_dir_EQ, + OPT_ffile_compilation_dir_EQ, + OPT_fcoverage_prefix_map_EQ, OPT_fsanitize_EQ, OPT_fno_sanitize_EQ, OPT_fsanitize_minimal_runtime, @@ -9556,29 +9594,24 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, OPT_fsanitize_trap_EQ, OPT_fno_sanitize_trap_EQ}; const llvm::DenseSet<unsigned> LinkerOptions{OPT_mllvm, OPT_Zlinker_input}; + auto ToolChainHasRT = [&](const ToolChain &TC, StringRef Name) { + return TC.getVFS().exists( + TC.getCompilerRT(Args, Name, ToolChain::FT_Static)); + }; auto ShouldForwardForToolChain = [&](Arg *A, const ToolChain &TC) { - auto HasProfileRT = TC.getVFS().exists( - TC.getCompilerRT(Args, "profile", ToolChain::FT_Static)); + unsigned ID = A->getOption().getID(); // Don't forward profiling arguments if the toolchain doesn't support it. // Without this check using it on the host would result in linker errors. - if (!HasProfileRT && - (A->getOption().matches(OPT_fprofile_generate) || - A->getOption().matches(OPT_fprofile_generate_EQ) || - A->getOption().matches(OPT_fprofile_instr_generate) || - A->getOption().matches(OPT_fprofile_instr_generate_EQ))) + // Coverage mapping flags require -fprofile-instr-generate, so drop them + // together to avoid a device cc1 diagnostic. + if (requiresProfileRT(ID) && !ToolChainHasRT(TC, "profile")) return false; - auto HasUBSanRT = TC.getVFS().exists( - TC.getCompilerRT(Args, "ubsan_minimal", ToolChain::FT_Static)); // Don't forward sanitizer arguments if the toolchain doesn't support it. // Without this check using it on the host would result in linker errors. - if (!HasUBSanRT && - (A->getOption().matches(OPT_fsanitize_EQ) || - A->getOption().matches(OPT_fno_sanitize_EQ) || - A->getOption().matches(OPT_fsanitize_minimal_runtime) || - A->getOption().matches(OPT_fno_sanitize_minimal_runtime))) + if (requiresUBSanRT(ID) && !ToolChainHasRT(TC, "ubsan_minimal")) return false; // Don't forward -mllvm to toolchains that don't support LLVM. - return TC.HasNativeLLVMSupport() || A->getOption().getID() != OPT_mllvm; + return TC.HasNativeLLVMSupport() || ID != OPT_mllvm; }; auto ShouldForward = [&](const llvm::DenseSet<unsigned> &Set, Arg *A, const ToolChain &TC) { diff --git a/clang/test/CodeGenHIP/profile-coverage-mapping.hip b/clang/test/CodeGenHIP/profile-coverage-mapping.hip new file mode 100644 index 0000000000000..dcdf539849004 --- /dev/null +++ b/clang/test/CodeGenHIP/profile-coverage-mapping.hip @@ -0,0 +1,22 @@ +// REQUIRES: amdgpu-registered-target + +// Device-side compilation must emit coverage mapping sections and call the +// GPU profile instrumentation runtime. + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1100 \ +// RUN: -fcuda-is-device -emit-llvm -o - %s \ +// RUN: -fprofile-instrument=clang -fcoverage-mapping | FileCheck %s + +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) + +// CHECK-DAG: section "__llvm_covfun" +// CHECK-DAG: @__llvm_coverage_mapping = {{.*}}section "__llvm_covmap" +// CHECK-DAG: @__profc__Z3dfni = {{.*}}section "__llvm_prf_cnts" +// CHECK-DAG: @__profc__Z6kerneli = {{.*}}section "__llvm_prf_cnts" +// CHECK-DAG: @__llvm_prf_nm = {{.*}}section "__llvm_prf_names" +// CHECK: call void @__llvm_profile_instrument_gpu + +__device__ int dfn(int i) { return i ? i + 1 : 0; } + +__global__ void kernel(int i) { dfn(i); } diff --git a/clang/test/Driver/hip-options.hip b/clang/test/Driver/hip-options.hip index ed20239605c22..25d5d7f7673d1 100644 --- a/clang/test/Driver/hip-options.hip +++ b/clang/test/Driver/hip-options.hip @@ -261,3 +261,28 @@ // RUN: FileCheck -check-prefix=JOBSV %s // JOBSV: clang-linker-wrapper{{.*}} "--wrapper-jobs=jobserver" +// Check that -fcoverage-mapping and related flags are forwarded to the device compiler. + +// RUN: %clang -### -Werror --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib \ +// RUN: -resource-dir=%S/Inputs/resource_dir_with_per_target_subdir \ +// RUN: --offload-arch=gfx1100 --offload-new-driver \ +// RUN: -fprofile-instr-generate -fcoverage-mapping %s 2>&1 | \ +// RUN: FileCheck -check-prefix=COV %s +// COV: clang-linker-wrapper{{.*}} "--device-compiler=amdgcn-amd-amdhsa=-fprofile-instr-generate" "--device-compiler=amdgcn-amd-amdhsa=-fcoverage-mapping" + +// RUN: %clang -### -Werror --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib \ +// RUN: -resource-dir=%S/Inputs/resource_dir_with_per_target_subdir \ +// RUN: --offload-arch=gfx1100 --offload-new-driver \ +// RUN: -fprofile-instr-generate -fcoverage-mapping \ +// RUN: -fcoverage-compilation-dir=/src %s 2>&1 | \ +// RUN: FileCheck -check-prefix=COV-DIR %s +// COV-DIR: clang-linker-wrapper{{.*}} "--device-compiler=amdgcn-amd-amdhsa=-fcoverage-compilation-dir=/src" + +// RUN: %clang -### -Werror --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib \ +// RUN: -resource-dir=%S/Inputs/resource_dir \ +// RUN: --offload-arch=gfx1100 --offload-new-driver \ +// RUN: -fprofile-instr-generate -fcoverage-mapping %s 2>&1 | \ +// RUN: FileCheck -check-prefix=NO-COV %s +// NO-COV-NOT: --device-compiler=amdgcn-amd-amdhsa=-fprofile-instr-generate +// NO-COV-NOT: --device-compiler=amdgcn-amd-amdhsa=-fcoverage-mapping + _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
