https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/236081
>From 4bd6a03fb527f51e037aa9567e3934740e4442be Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <[email protected]> Date: Mon, 26 Oct 2026 21:43:19 +0000 Subject: [PATCH 1/5] [profile] Preserve device instrumentation layout flags when writing profiles The LLVM options -pgo-instrument-entry and -pgo-instrument-loop-entries control whether function entries and edges entering loops must be instrumented. Both default to false for ordinary IR PGO. These settings affect counter placement and are recorded as flags in the raw profile header so profile use can reconstruct the same counter layout. HIP users can specify different values for host and device compilation. For example, -Xarch_device -mllvm=-pgo-instrument-entry=true enables explicit entry counters for the device while the host keeps the default. The host and device profiles then need different entry flags. The ROCm collector currently writes device counters with the host runtime's version flags. In this example, it incorrectly marks the device profile as having entry instrumentation disabled. Profile use can then reconstruct the wrong layout and associate counts with the wrong blocks. Read and validate the device's profile-version variable separately from the host's, then pass it to the existing raw-profile writer. This preserves the device's own layout flags in its profile header so profile use can interpret the device counters correctly. --- .../profile/InstrProfilingPlatformROCm.cpp | 16 +++- .../GPU/instrprof-hip-layout-flags.hip | 95 +++++++++++++++++++ 2 files changed, 110 insertions(+), 1 deletion(-) create mode 100644 compiler-rt/test/profile/GPU/instrprof-hip-layout-flags.hip diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp index 376957b1177841..2e3cbd4f0428d5 100644 --- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp +++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp @@ -828,6 +828,20 @@ int __prof_rocm::processDeviceOffloadPrf( if (CountersSize == 0 || DataSize == 0) return 0; + // The host and device may use different instrumentation layouts. + uint64_t Version; + if (!HostSections.VersionVar || hipMemcpy(&Version, HostSections.VersionVar, + sizeof(Version), 2 /*DToH*/) != 0) { + PROF_ERR("%s\n", "failed to copy profile version from device"); + return -1; + } + if (GET_VERSION(Version) != INSTR_PROF_RAW_VERSION) { + PROF_ERR("Device runtime and instrumentation version mismatch: " + "expected %d, but get %d\n", + INSTR_PROF_RAW_VERSION, (int)GET_VERSION(Version)); + return -1; + } + int ret = -1; /* Sections using linker-defined __start_/__stop_ bounds are shared across @@ -1122,7 +1136,7 @@ int __prof_rocm::processDeviceOffloadPrf( (__llvm_profile_data *)(BufDataBegin + DataSize), BufCountersBegin, BufCountersBegin + CountersSize, UCnts.HostBegin, UCnts.HostBegin ? UCnts.HostBegin + UniformCountersSize : nullptr, - BufNamesBegin, BufNamesBegin + NamesSize, nullptr); + BufNamesBegin, BufNamesBegin + NamesSize, &Version); if (ret != 0) { PROF_ERR("%s\n", "failed to write device profile using shared API"); diff --git a/compiler-rt/test/profile/GPU/instrprof-hip-layout-flags.hip b/compiler-rt/test/profile/GPU/instrprof-hip-layout-flags.hip new file mode 100644 index 00000000000000..5673d0e7654d30 --- /dev/null +++ b/compiler-rt/test/profile/GPU/instrprof-hip-layout-flags.hip @@ -0,0 +1,95 @@ +// NOTE: Do not autogenerate. Check runtime profile flags and reconstructed entry count. +// Device profile layout flags must be independent of the host flags. +// REQUIRES: hip, amdgpu +// +// RUN: rm -rf %t.dir && mkdir -p %t.dir +// RUN: mkdir %t.dir/entry +// RUN: %clang -x hip -O2 -Xarch_device -nogpulib -Xarch_device -nogpuinc \ +// RUN: -fprofile-generate --offload-arch=%amdgpu_arch \ +// RUN: --rocm-path=%hip_lib_path/.. %s -L%hip_lib_path -lamdhip64 \ +// RUN: -Xarch_device -mllvm=-pgo-instrument-entry \ +// RUN: -o %t.dir/entry/test +// RUN: env LLVM_PROFILE_FILE=%t.dir/entry/host.profraw \ +// RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ +// RUN: HIP_VISIBLE_DEVICES=0 %run %t.dir/entry/test +// RUN: llvm-profdata show %t.dir/entry/gfx*.profraw | FileCheck %s --check-prefix=ENTRY +// RUN: llvm-profdata show %t.dir/entry/host.profraw | FileCheck %s --check-prefix=NEITHER +// RUN: mkdir %t.dir/loop +// RUN: %clang -x hip -O2 -Xarch_device -nogpulib -Xarch_device -nogpuinc \ +// RUN: -fprofile-generate --offload-arch=%amdgpu_arch \ +// RUN: --rocm-path=%hip_lib_path/.. %s -L%hip_lib_path -lamdhip64 \ +// RUN: -Xarch_device -mllvm=-pgo-instrument-loop-entries \ +// RUN: -o %t.dir/loop/test +// RUN: env LLVM_PROFILE_FILE=%t.dir/loop/host.profraw \ +// RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ +// RUN: HIP_VISIBLE_DEVICES=0 %run %t.dir/loop/test +// RUN: llvm-profdata show %t.dir/loop/gfx*.profraw | FileCheck %s --check-prefix=LOOP +// RUN: llvm-profdata show %t.dir/loop/host.profraw | FileCheck %s --check-prefix=NEITHER +// RUN: mkdir %t.dir/both +// RUN: %clang -x hip -O2 -Xarch_device -nogpulib -Xarch_device -nogpuinc \ +// RUN: -fprofile-generate --offload-arch=%amdgpu_arch \ +// RUN: --rocm-path=%hip_lib_path/.. %s -L%hip_lib_path -lamdhip64 \ +// RUN: -Xarch_device -mllvm=-pgo-instrument-entry \ +// RUN: -Xarch_device -mllvm=-pgo-instrument-loop-entries \ +// RUN: -o %t.dir/both/test +// RUN: env LLVM_PROFILE_FILE=%t.dir/both/host.profraw \ +// RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ +// RUN: HIP_VISIBLE_DEVICES=0 %run %t.dir/both/test +// RUN: llvm-profdata show %t.dir/both/gfx*.profraw | FileCheck %s --check-prefix=BOTH +// RUN: llvm-profdata show %t.dir/both/host.profraw | FileCheck %s --check-prefix=NEITHER +// RUN: mkdir %t.dir/neither +// RUN: %clang -x hip -O2 -Xarch_device -nogpulib -Xarch_device -nogpuinc \ +// RUN: -fprofile-generate --offload-arch=%amdgpu_arch \ +// RUN: --rocm-path=%hip_lib_path/.. %s -L%hip_lib_path -lamdhip64 \ +// RUN: -Xarch_host -mllvm=-pgo-instrument-entry \ +// RUN: -Xarch_host -mllvm=-pgo-instrument-loop-entries \ +// RUN: -o %t.dir/neither/test +// RUN: env LLVM_PROFILE_FILE=%t.dir/neither/host.profraw \ +// RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ +// RUN: HIP_VISIBLE_DEVICES=0 %run %t.dir/neither/test +// RUN: llvm-profdata show %t.dir/neither/gfx*.profraw | FileCheck %s --check-prefix=NEITHER +// RUN: llvm-profdata show %t.dir/neither/host.profraw | FileCheck %s --check-prefix=BOTH +// RUN: llvm-profdata merge %t.dir/both/gfx*.profraw -o %t.dir/device.profdata +// RUN: %clang -x hip -O2 -Xarch_device -nogpulib -Xarch_device -nogpuinc \ +// RUN: --offload-arch=%amdgpu_arch --offload-device-only \ +// RUN: --rocm-path=%hip_lib_path/.. -fprofile-use=%t.dir/device.profdata \ +// RUN: -Werror=profile-instr-out-of-date -Werror=profile-instr-unprofiled \ +// RUN: -S -emit-llvm %s -o - | FileCheck %s --check-prefix=USE +// +// ENTRY: Instrumentation level: IR entry_first = 1 instrument_loop_entries = 0 +// LOOP: Instrumentation level: IR entry_first = 0 instrument_loop_entries = 1 +// BOTH: Instrumentation level: IR entry_first = 1 instrument_loop_entries = 1 +// NEITHER: Instrumentation level: IR entry_first = 0 instrument_loop_entries = 0 +// USE-LABEL: define{{.*}} @layout_kernel( +// USE-SAME: !prof ![[ENTRY_COUNT:[0-9]+]] +// USE: ![[ENTRY_COUNT]] = !{!"function_entry_count", i64 64} + +#ifdef __HIP_DEVICE_COMPILE__ +# define __global__ __attribute__((global)) +#else +# include <hip/hip_runtime.h> +#endif + +extern "C" __global__ void layout_kernel(int *Out, int N) { + unsigned I = 0; +#ifdef __HIP_DEVICE_COMPILE__ + I = __builtin_amdgcn_workitem_id_x(); +#endif + Out[I] = (I & 1) ? N : 2 * N; +} + +#ifndef __HIP_DEVICE_COMPILE__ +int main() { + int *Dev, Host[64]; + if (hipMalloc(&Dev, sizeof(Host)) != hipSuccess) + return 1; + layout_kernel<<<1, 64>>>(Dev, 10); + if (hipDeviceSynchronize() != hipSuccess || + hipMemcpy(Host, Dev, sizeof(Host), hipMemcpyDeviceToHost) != hipSuccess) + return 2; + for (int I = 0; I < 64; ++I) + if (Host[I] != ((I & 1) ? 10 : 20)) + return 3; + return hipFree(Dev) != hipSuccess; +} +#endif >From aa2d9f4906c534ec800104d4e0406dac02951e7b Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <[email protected]> Date: Wed, 28 Oct 2026 23:59:25 +0000 Subject: [PATCH 2/5] [HIP][profile] Register the device version on Windows Windows profile collection needs a registered handle for the device version before copying it. Register a handle per translation unit and resolve it through HIP, while keeping Linux on the HSA collection path. Ignore dependent registration callbacks when their table registration fails so they cannot modify the previous module's handles. --- clang/lib/CodeGen/CGCUDANV.cpp | 83 ++++++++----- .../CodeGenHIP/offload-pgo-version-shadow.hip | 72 ++++++++++++ .../profile/InstrProfilingPlatformROCm.cpp | 30 ++++- .../instrprof-rocm-shadow-registration.cpp | 110 ++++++++++++++++++ 4 files changed, 265 insertions(+), 30 deletions(-) create mode 100644 clang/test/CodeGenHIP/offload-pgo-version-shadow.hip create mode 100644 compiler-rt/test/profile/instrprof-rocm-shadow-registration.cpp diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index 7c38bde6a78d8f..e85d7a390265ae 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -80,6 +80,8 @@ class CGNVCUDARuntime : public CGCUDARuntime { /// __hipRegisterVar (non-RDC) or an offloading entry (RDC) so the runtime /// can locate the device-side table by name. llvm::GlobalVariable *OffloadProfShadow = nullptr; + /// Windows needs a registered handle for the device raw-version symbol. + llvm::GlobalVariable *OffloadProfVersionShadow = nullptr; struct OffloadProfSectionShadowInfo { llvm::GlobalVariable *Shadow; std::string DeviceName; @@ -767,6 +769,23 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() { } } + auto RegisterProfileShadow = [&](llvm::GlobalVariable *Shadow, + StringRef DeviceName) { + llvm::Constant *Name = makeConstantString(DeviceName.str()); + llvm::Constant *IntZero = llvm::ConstantInt::get(IntTy, 0); + uint64_t Size = + CGM.getDataLayout().getTypeAllocSize(Shadow->getValueType()); + llvm::Value *Args[] = {&GpuBinaryHandlePtr, + Shadow, + Name, + Name, + IntZero, + llvm::ConstantInt::get(VarSizeTy, Size), + IntZero, + IntZero}; + Builder.CreateCall(RegisterVar, Args); + }; + // Register the per-TU offload-profiling shadow so the host runtime can // locate the matching device-side __llvm_profile_sections_<CUID>. We // emit both __hipRegisterVar (so the HIP runtime can map the host @@ -774,20 +793,7 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() { // __llvm_profile_offload_register_shadow_variable (so the profile // runtime adds the shadow to its drain list). if (OffloadProfShadow) { - llvm::Constant *Name = - makeConstantString(std::string(OffloadProfShadow->getName())); - llvm::Constant *IntZero = llvm::ConstantInt::get(IntTy, 0); - llvm::Value *RegisterVarArgs[] = { - &GpuBinaryHandlePtr, - OffloadProfShadow, - Name, - Name, - IntZero, - llvm::ConstantInt::get(VarSizeTy, - CGM.getDataLayout().getPointerSize(/*AS=*/0)), - IntZero, - IntZero}; - Builder.CreateCall(RegisterVar, RegisterVarArgs); + RegisterProfileShadow(OffloadProfShadow, OffloadProfShadow->getName()); llvm::FunctionCallee RegisterShadow = CGM.CreateRuntimeFunction( llvm::FunctionType::get(VoidTy, {PtrTy}, false), @@ -795,24 +801,21 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() { Builder.CreateCall(RegisterShadow, {OffloadProfShadow}); } + if (OffloadProfVersionShadow) { + RegisterProfileShadow(OffloadProfVersionShadow, + "__llvm_profile_raw_version"); + llvm::FunctionCallee RegisterVersionShadow = CGM.CreateRuntimeFunction( + llvm::FunctionType::get(VoidTy, {PtrTy}, false), + "__llvm_profile_offload_register_version_shadow_variable"); + Builder.CreateCall(RegisterVersionShadow, {OffloadProfVersionShadow}); + } + if (!OffloadProfSectionShadows.empty()) { llvm::FunctionCallee RegisterSectionShadow = CGM.CreateRuntimeFunction( llvm::FunctionType::get(VoidTy, {PtrTy}, false), "__llvm_profile_offload_register_section_shadow_variable"); - llvm::Constant *IntZero = llvm::ConstantInt::get(IntTy, 0); for (const auto &Info : OffloadProfSectionShadows) { - llvm::Constant *Name = makeConstantString(Info.DeviceName); - llvm::Value *RegisterVarArgs[] = { - &GpuBinaryHandlePtr, - Info.Shadow, - Name, - Name, - IntZero, - llvm::ConstantInt::get(VarSizeTy, - CGM.getDataLayout().getPointerSize(/*AS=*/0)), - IntZero, - IntZero}; - Builder.CreateCall(RegisterVar, RegisterVarArgs); + RegisterProfileShadow(Info.Shadow, Info.DeviceName); Builder.CreateCall(RegisterSectionShadow, {Info.Shadow}); } } @@ -1372,6 +1375,21 @@ void CGNVCUDARuntime::createOffloadingEntries() { auto *Entry = llvm::BasicBlock::Create(Ctx, "entry", CtorFn); llvm::IRBuilder<> B(Entry); B.CreateCall(RegisterShadow, {OffloadProfShadow}); + if (OffloadProfVersionShadow) { + uint64_t VersionSize = CGM.getDataLayout().getTypeAllocSize( + OffloadProfVersionShadow->getValueType()); + auto *VersionEntry = llvm::offloading::emitOffloadingEntry( + M, Kind, OffloadProfVersionShadow, "__llvm_profile_raw_version", + VersionSize, llvm::offloading::OffloadGlobalEntry, /*Data=*/0); + // The device symbol is shared, but each TU needs its own host handle. + // Keep the weak offloading entries from coalescing at the host link. + VersionEntry->setName(VersionEntry->getName() + "." + + CGM.getContext().getCUIDHash()); + llvm::FunctionCallee RegisterVersionShadow = CGM.CreateRuntimeFunction( + llvm::FunctionType::get(VoidTy, {PtrTy}, false), + "__llvm_profile_offload_register_version_shadow_variable"); + B.CreateCall(RegisterVersionShadow, {OffloadProfVersionShadow}); + } for (const auto &Info : OffloadProfSectionShadows) { llvm::offloading::emitOffloadingEntry( M, Kind, Info.Shadow, Info.DeviceName, @@ -1438,6 +1456,17 @@ void CGNVCUDARuntime::emitOffloadProfilingSections() { llvm::ConstantPointerNull::get(PtrTy), Name); CGM.addCompilerUsedGlobal(OffloadProfShadow); + // Linux collects the raw version through HSA. Windows resolves it through + // HIP's registered-symbol API, using a distinct handle for each TU. + if (CGM.getTriple().isOSWindows()) { + auto *VersionTy = llvm::Type::getInt64Ty(Ctx); + OffloadProfVersionShadow = new llvm::GlobalVariable( + M, VersionTy, /*isConstant=*/false, llvm::GlobalValue::ExternalLinkage, + llvm::ConstantInt::get(VersionTy, 0), + ("__llvm_profile_shadow_version_" + CUIDHash).str()); + CGM.addCompilerUsedGlobal(OffloadProfVersionShadow); + } + auto AddSectionShadow = [&](StringRef Kind, const Twine &DeviceName) { std::string ShadowName = (Twine("__llvm_profile_shadow_") + Kind + "_" + CUIDHash + "_" + diff --git a/clang/test/CodeGenHIP/offload-pgo-version-shadow.hip b/clang/test/CodeGenHIP/offload-pgo-version-shadow.hip new file mode 100644 index 00000000000000..e9e7cfd08646d5 --- /dev/null +++ b/clang/test/CodeGenHIP/offload-pgo-version-shadow.hip @@ -0,0 +1,72 @@ +// NOTE: Do not autogenerate. Targeted checks express the relationships between +// global handles, registration calls, and entries retained by a two-TU link. +// REQUIRES: x86-registered-target +// +// Windows registers a per-TU host handle for the same module's raw version. +// Its eight-byte size is independent of the host pointer size. +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -cuid=abc \ +// RUN: -fprofile-instrument=clang -emit-llvm -o - -x hip %s \ +// RUN: | FileCheck %s --check-prefixes=WIN,STATIC --implicit-check-not=offloading.entry +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -cuid=abc -fgpu-rdc \ +// RUN: -fprofile-instrument=clang -emit-llvm -o - -x hip %s \ +// RUN: | FileCheck %s --check-prefixes=WIN,RDC +// RUN: %clang_cc1 -triple i686-pc-windows-msvc -cuid=abc \ +// RUN: -fprofile-instrument=clang -emit-llvm -o - -x hip %s \ +// RUN: | FileCheck %s --check-prefixes=WIN,STATIC +// +// Linux uses HSA collection, so it needs no version handle. +// RUN: %clang_cc1 -triple x86_64-linux-gnu -cuid=abc \ +// RUN: -fprofile-instrument=clang -emit-llvm -o - -x hip %s \ +// RUN: | FileCheck %s --check-prefix=NO-VERSION +// RUN: %clang_cc1 -triple x86_64-linux-gnu -cuid=abc -fgpu-rdc \ +// RUN: -fprofile-instrument=clang -emit-llvm -o - -x hip %s \ +// RUN: | FileCheck %s --check-prefix=NO-VERSION +// +// No profile instrumentation or no CUID means no version registration. +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -cuid=abc \ +// RUN: -emit-llvm -o - -x hip %s \ +// RUN: | FileCheck %s --check-prefix=NO-VERSION +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc \ +// RUN: -fprofile-instrument=clang -emit-llvm -o - -x hip %s \ +// RUN: | FileCheck %s --check-prefix=NO-VERSION + +// The shared device symbol must retain both per-TU host registrations at link. +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -cuid=abc -fgpu-rdc \ +// RUN: -fprofile-instrument=clang -emit-llvm -o %t.first.ll -x hip %s +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -cuid=def -fgpu-rdc \ +// RUN: -DSECOND_TU -fprofile-instrument=clang -emit-llvm -o %t.second.ll -x hip %s +// RUN: llvm-link %t.first.ll %t.second.ll -S -o - | FileCheck %s --check-prefix=MULTI + +#ifdef SECOND_TU +#define kernel second_kernel +#endif +extern "C" __attribute__((global)) void kernel(int *p) { *p += 1; } + +// WIN: @__llvm_profile_sections_[[CUID:[0-9a-f]+]] = global ptr null +// WIN: @__llvm_profile_shadow_version_[[CUID]] = global i64 0 +// STATIC: [[NAME:@[^ ]+]] = private unnamed_addr constant [27 x i8] c"__llvm_profile_raw_version\00" +// RDC: [[NAME:@[^ ]+]] = internal unnamed_addr constant [27 x i8] c"__llvm_profile_raw_version\00" +// RDC: @.offloading.entry.{{[^ ]+}} = weak constant %struct.__tgt_offload_entry {{.*}}ptr @__llvm_profile_shadow_version_[[CUID]], ptr [[NAME]], i64 8, +// WIN: @llvm.compiler.used = {{.*}}ptr @__llvm_profile_shadow_version_[[CUID]] +// STATIC-LABEL: define internal void @__hip_register_globals( +// STATIC: call void @__llvm_profile_offload_register_shadow_variable(ptr @__llvm_profile_sections_[[CUID]]) +// STATIC-NEXT: call void @__hipRegisterVar({{.*}}ptr @__llvm_profile_shadow_version_[[CUID]], ptr [[NAME]], ptr [[NAME]], i32 0, i{{32|64}} 8, i32 0, i32 0) +// STATIC-NEXT: call void @__llvm_profile_offload_register_version_shadow_variable(ptr @__llvm_profile_shadow_version_[[CUID]]) +// RDC-LABEL: define internal void @__llvm_profile_register_shadow. +// RDC-SAME: [[CUID]]() +// RDC: call void @__llvm_profile_offload_register_shadow_variable(ptr @__llvm_profile_sections_[[CUID]]) +// RDC-NEXT: call void @__llvm_profile_offload_register_version_shadow_variable(ptr @__llvm_profile_shadow_version_[[CUID]]) +// WIN: call void @__llvm_profile_offload_register_section_shadow_variable(ptr @__llvm_profile_shadow_data_[[CUID]]_0) +// WIN: call void @__llvm_profile_offload_register_section_shadow_variable(ptr @__llvm_profile_shadow_cnts_[[CUID]]_1) +// WIN: call void @__llvm_profile_offload_register_section_shadow_variable(ptr @__llvm_profile_shadow_ucnts_[[CUID]]_2) +// WIN: call void @__llvm_profile_offload_register_section_shadow_variable(ptr @__llvm_profile_shadow_names_[[CUID]]_3) + +// NO-VERSION-NOT: __llvm_profile_shadow_version_ +// NO-VERSION-NOT: __llvm_profile_offload_register_version_shadow_variable + +// MULTI-DAG: @__llvm_profile_shadow_version_[[FIRST:[0-9a-f]+]] = global i64 0 +// MULTI-DAG: @__llvm_profile_shadow_version_[[SECOND:[0-9a-f]+]] = global i64 0 +// MULTI-DAG: @.offloading.entry.{{[^ ]+}} = weak constant %struct.__tgt_offload_entry {{.*}}ptr @__llvm_profile_shadow_version_[[FIRST]], {{.*}}i64 8, +// MULTI-DAG: @.offloading.entry.{{[^ ]+}} = weak constant %struct.__tgt_offload_entry {{.*}}ptr @__llvm_profile_shadow_version_[[SECOND]], {{.*}}i64 8, +// MULTI-DAG: call void @__llvm_profile_offload_register_version_shadow_variable(ptr @__llvm_profile_shadow_version_[[FIRST]]) +// MULTI-DAG: call void @__llvm_profile_offload_register_version_shadow_variable(ptr @__llvm_profile_shadow_version_[[SECOND]]) diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp index 2e3cbd4f0428d5..2d03613e9fa1f6 100644 --- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp +++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp @@ -621,6 +621,7 @@ extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) { static void **OffloadShadowVariables = nullptr; static int NumShadowVariables = 0; static int CapShadowVariables = 0; +static bool CurrentShadowVariableRegistered = false; struct OffloadSectionShadow { void *Data; @@ -630,6 +631,7 @@ struct OffloadSectionShadow { }; struct OffloadSectionShadowGroup { + void *VersionShadow; OffloadSectionShadow *Shadows; int NumShadows; int CapShadows; @@ -652,6 +654,8 @@ static int ensureSectionShadowCapacity(OffloadSectionShadowGroup *Group, } extern "C" void __llvm_profile_offload_register_shadow_variable(void *ptr) { + // Dependent callbacks must not use the previous group if this attempt fails. + CurrentShadowVariableRegistered = false; if (growArray((void **)&OffloadShadowVariables, &CapShadowVariables, NumShadowVariables + 1, 64, sizeof(*OffloadShadowVariables))) return; @@ -661,11 +665,12 @@ extern "C" void __llvm_profile_offload_register_shadow_variable(void *ptr) { OffloadShadowVariables[Index] = ptr; __builtin_memset(&OffloadSectionShadowGroups[Index], 0, sizeof(OffloadSectionShadowGroups[Index])); + CurrentShadowVariableRegistered = true; } extern "C" void __llvm_profile_offload_register_section_shadow_variable(void *ptr) { - if (NumShadowVariables == 0) + if (!CurrentShadowVariableRegistered) return; /* Match CGCUDANV.cpp: data, counters, uniform counters, then names for each @@ -696,6 +701,13 @@ __llvm_profile_offload_register_section_shadow_variable(void *ptr) { ++Group->NumSections; } +extern "C" void +__llvm_profile_offload_register_version_shadow_variable(void *ptr) { + if (!CurrentShadowVariableRegistered) + return; + OffloadSectionShadowGroups[NumShadowVariables - 1].VersionShadow = ptr; +} + namespace { struct ProfileSectionCopy { @@ -829,9 +841,21 @@ int __prof_rocm::processDeviceOffloadPrf( return 0; // The host and device may use different instrumentation layouts. +#ifdef _WIN32 + // Use the host handle to make the version address known to HIP. + void *VersionVar = nullptr; + if (!Sections || !Sections->VersionShadow || + hipGetSymbolAddress(&VersionVar, Sections->VersionShadow) != 0 || + !VersionVar) { + PROF_ERR("%s\n", "failed to get profile version symbol address"); + return -1; + } +#else + const void *VersionVar = HostSections.VersionVar; +#endif uint64_t Version; - if (!HostSections.VersionVar || hipMemcpy(&Version, HostSections.VersionVar, - sizeof(Version), 2 /*DToH*/) != 0) { + if (!VersionVar || + hipMemcpy(&Version, VersionVar, sizeof(Version), 2 /*DToH*/) != 0) { PROF_ERR("%s\n", "failed to copy profile version from device"); return -1; } diff --git a/compiler-rt/test/profile/instrprof-rocm-shadow-registration.cpp b/compiler-rt/test/profile/instrprof-rocm-shadow-registration.cpp new file mode 100644 index 00000000000000..2f48873e183a54 --- /dev/null +++ b/compiler-rt/test/profile/instrprof-rocm-shadow-registration.cpp @@ -0,0 +1,110 @@ +// Check that callbacks following a failed table registration cannot modify +// the previous table's version or section handles. No GPU runtime is needed. +// +// Compile without instrumentation; the link supplies the profile library +// for the collector's writer and padding helpers. +// REQUIRES: linux +// RUN: %clangxx -I%S/../../lib -I%S/../../include -c %s -o %t.o +// RUN: %clangxx_profgen %t.o -pthread -ldl -o %t +// RUN: %run %t | FileCheck %s + +#include <assert.h> +#include <stdio.h> +#include <stdlib.h> + +static int ReallocsUntilFailure = -1; +static void *controlledRealloc(void *Ptr, size_t Size) { + if (ReallocsUntilFailure == 0) { + ReallocsUntilFailure = -1; + return nullptr; + } + if (ReallocsUntilFailure > 0) + --ReallocsUntilFailure; + return realloc(Ptr, Size); +} + +// Include the production registration callbacks with allocation fault +// injection. +#define SANITIZER_COMMON_NO_REDEFINE_BUILTINS +#define realloc controlledRealloc +#include "profile/InstrProfilingPlatformROCm.cpp" +#undef realloc + +// Disable runtime discovery: this test exercises host registration only. +namespace __interception { +bool DynamicLoaderAvailable() { return false; } +void *OpenLibrary(const char *) { + assert(false); + return nullptr; +} +void *LookupSymbol(void *, const char *) { + assert(false); + return nullptr; +} +bool InterceptFunction(const char *, uptr *, uptr, uptr) { + assert(false); + return false; +} +} // namespace __interception + +namespace __prof_rocm { +void profRecordDrainedBounds(const void *, const void *, const void *) {} +int hsaRuntimeAvailable() { return 0; } +int drainDevicesViaHsa() { return 0; } +} // namespace __prof_rocm + +struct Handles { + int Table, Version, Sections[4]; +}; + +static void registerDependents(Handles &H) { + __llvm_profile_offload_register_version_shadow_variable(&H.Version); + for (int &Section : H.Sections) + __llvm_profile_offload_register_section_shadow_variable(&Section); +} + +static void checkGroup(int Index, Handles &H) { + const auto &Group = OffloadSectionShadowGroups[Index]; + assert(OffloadShadowVariables[Index] == &H.Table); + assert(Group.VersionShadow == &H.Version); + assert(Group.NumSections == 4 && Group.NumShadows == 1); + assert(Group.Shadows[0].Data == &H.Sections[0]); + assert(Group.Shadows[0].Counters == &H.Sections[1]); + assert(Group.Shadows[0].UniformCounters == &H.Sections[2]); + assert(Group.Shadows[0].Names == &H.Sections[3]); +} + +int main() { + Handles Tables[65] = {}; + // Dependent callbacks before the first table must also be ignored. + registerDependents(Tables[64]); + assert(NumShadowVariables == 0); + + for (int I = 0; I != 64; ++I) { + __llvm_profile_offload_register_shadow_variable(&Tables[I].Table); + registerDependents(Tables[I]); + } + assert(NumShadowVariables == 64); + assert(CapShadowVariables == 64 && CapSectionShadowGroups == 64); + + for (int Allocation = 0; Allocation != 2; ++Allocation) { + // First fail growing the table array, then fail growing the group array + // after the table array has successfully grown. + ReallocsUntilFailure = Allocation; + __llvm_profile_offload_register_shadow_variable(&Tables[64].Table); + assert(ReallocsUntilFailure == -1 && NumShadowVariables == 64); + registerDependents(Tables[64]); + for (int I = 0; I != 64; ++I) + checkGroup(I, Tables[I]); + } + + // A later successful table must start accepting its own callbacks again. + __llvm_profile_offload_register_shadow_variable(&Tables[64].Table); + registerDependents(Tables[64]); + assert(NumShadowVariables == 65); + for (int I = 0; I != 65; ++I) + checkGroup(I, Tables[I]); + puts("PASS"); +} + +// CHECK: PASS >From 2821bf8c066bd639f00fd04fddeb1822dd7c0a0e Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <[email protected]> Date: Sun, 1 Nov 2026 16:49:32 -0500 Subject: [PATCH 3/5] [Profile] Preserve GPU raw profile version visibility Create protected raw-version definitions on GPUs so HIP profile collection can find the symbol after device LTO. Keep existing CPU visibility and align the runtime declaration and weak fallback with the compiler-generated override. Add compiler producer, runtime declaration/default, and final dynamic-symbol regressions. --- compiler-rt/lib/profile/InstrProfiling.h | 20 ++++++++++++------- .../lib/profile/InstrProfilingVersionVar.c | 4 ++-- .../GPU/instrprof-version-visibility.test | 14 +++++++++++++ .../Instrumentation/PGOInstrumentation.cpp | 5 +++-- .../PGOProfile/gpu-profile-version.ll | 12 +++++++++++ 5 files changed, 44 insertions(+), 11 deletions(-) create mode 100644 compiler-rt/test/profile/GPU/instrprof-version-visibility.test create mode 100644 llvm/test/Transforms/PGOProfile/gpu-profile-version.ll diff --git a/compiler-rt/lib/profile/InstrProfiling.h b/compiler-rt/lib/profile/InstrProfiling.h index 928b045ab87afc..b2166d6fc2c33c 100644 --- a/compiler-rt/lib/profile/InstrProfiling.h +++ b/compiler-rt/lib/profile/InstrProfiling.h @@ -299,14 +299,20 @@ int __llvm_write_custom_profile( COMPILER_RT_VISIBILITY extern int INSTR_PROF_PROFILE_RUNTIME_VAR; /*! - * This variable is defined in InstrProfilingVersionVar.c as a hidden symbol - * (except on Apple platforms where this symbol is checked by TAPI). Its main - * purpose is to encode the raw profile version value and other format related - * information such as whether the profile is from IR based instrumentation. The - * variable is defined as weak so that compiler can emit an overriding - * definition depending on user option. + * This variable has protected visibility on GPUs and the usual runtime + * visibility on other targets. It is defined in InstrProfilingVersionVar.c. Its + * main purpose is to encode the raw profile version value and other format + * related information such as whether the profile is from IR based + * instrumentation. The variable is defined as weak so that compiler can emit an + * overriding definition depending on user option. */ -COMPILER_RT_VISIBILITY extern uint64_t +#if defined(__AMDGCN__) || defined(__NVPTX__) +#define INSTR_PROF_RAW_VERSION_VISIBILITY \ + __attribute__((visibility("protected"))) +#else +#define INSTR_PROF_RAW_VERSION_VISIBILITY COMPILER_RT_VISIBILITY +#endif +INSTR_PROF_RAW_VERSION_VISIBILITY extern uint64_t INSTR_PROF_RAW_VERSION_VAR; /* __llvm_profile_raw_version */ /*! diff --git a/compiler-rt/lib/profile/InstrProfilingVersionVar.c b/compiler-rt/lib/profile/InstrProfilingVersionVar.c index 21400bfb2caa0a..eb5b5cd5efa835 100644 --- a/compiler-rt/lib/profile/InstrProfilingVersionVar.c +++ b/compiler-rt/lib/profile/InstrProfilingVersionVar.c @@ -14,5 +14,5 @@ * user has not specified one. Set this up by moving the runtime's copy of this * symbol to an object file within the archive. */ -COMPILER_RT_VISIBILITY COMPILER_RT_WEAK uint64_t INSTR_PROF_RAW_VERSION_VAR = - INSTR_PROF_RAW_VERSION; +INSTR_PROF_RAW_VERSION_VISIBILITY COMPILER_RT_WEAK uint64_t + INSTR_PROF_RAW_VERSION_VAR = INSTR_PROF_RAW_VERSION; diff --git a/compiler-rt/test/profile/GPU/instrprof-version-visibility.test b/compiler-rt/test/profile/GPU/instrprof-version-visibility.test new file mode 100644 index 00000000000000..cf3f1dc386ad95 --- /dev/null +++ b/compiler-rt/test/profile/GPU/instrprof-version-visibility.test @@ -0,0 +1,14 @@ +# NOTE: Do not autogenerate. Check the runtime version definition and its device export. +# REQUIRES: system-linux, amdgpu, lld-available +# RUN: %clang -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../../lib/profile/InstrProfilingVersionVar.c -o - | FileCheck %s --check-prefix=CPU +# RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../../lib/profile/InstrProfilingVersionVar.c -o %t.ll +# RUN: FileCheck %s --check-prefix=GPU --input-file=%t.ll +# RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../../lib/profile/InstrProfiling.c -o - | FileCheck %s --check-prefix=DECL +# RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -flto -nostdlib -shared %t.ll -o %t.so +# RUN: llvm-readelf --dyn-syms %t.so | FileCheck %s --check-prefix=ELF + +# CPU: @__llvm_profile_raw_version = weak hidden global i64 +# GPU: @__llvm_profile_raw_version = weak protected addrspace(1) global i64 +# ELF: OBJECT WEAK PROTECTED {{[0-9]+}} __llvm_profile_raw_version + +# DECL: @__llvm_profile_raw_version = external protected addrspace(1) global i64 diff --git a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp index fdebfa44c0ae3e..613a1e4afbf59f 100644 --- a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp +++ b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp @@ -307,9 +307,10 @@ createIRLevelProfileFlagVar(const InstrumentationOptions &Opts, Module &M, auto IRLevelVersionVariable = new GlobalVariable( M, IntTy64, true, GlobalValue::WeakAnyLinkage, Constant::getIntegerValue(IntTy64, APInt(64, ProfileVersion)), VarName); - IRLevelVersionVariable->setVisibility(GlobalValue::HiddenVisibility); - Triple TT(M.getTargetTriple()); + IRLevelVersionVariable->setVisibility(TT.isGPU() + ? GlobalValue::ProtectedVisibility + : GlobalValue::HiddenVisibility); if (TT.supportsCOMDAT()) { IRLevelVersionVariable->setLinkage(GlobalValue::ExternalLinkage); IRLevelVersionVariable->setComdat(M.getOrInsertComdat(VarName)); diff --git a/llvm/test/Transforms/PGOProfile/gpu-profile-version.ll b/llvm/test/Transforms/PGOProfile/gpu-profile-version.ll new file mode 100644 index 00000000000000..f3734298b21113 --- /dev/null +++ b/llvm/test/Transforms/PGOProfile/gpu-profile-version.ll @@ -0,0 +1,12 @@ +; REQUIRES: amdgpu-registered-target, x86-registered-target +; NOTE: Do not autogenerate. Check the visibility of the shared profile version. +; RUN: opt -mtriple=amdgpu9.50-amd-amdhsa -passes=pgo-instr-gen -S %s | FileCheck %s --check-prefix=GPU +; RUN: opt -mtriple=x86_64-unknown-linux-gnu -passes=pgo-instr-gen -S %s | FileCheck %s --check-prefix=CPU + +; GPU: @__llvm_profile_raw_version = protected addrspace(1) constant i64 +; CPU: @__llvm_profile_raw_version = hidden constant i64 + +define void @profiled() { +entry: + ret void +} >From d2f8497fb4d4fbb7949712a96245521e228cf484 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <[email protected]> Date: Sun, 1 Nov 2026 20:40:31 -0500 Subject: [PATCH 4/5] [Profile] Run GPU version regressions without GPU hardware Place the runtime visibility regression in the profile suite and detect AMDGPU compiler support separately from HIP runtime availability. Link bitcode directly with LLD to avoid host linker flags from the clang substitution, and reuse opt to inspect the bitcode. Feed the PGO producer test through stdin to follow AMDGPU test conventions. --- ...ty.test => instrprof-gpu-version-visibility.test} | 12 ++++++------ compiler-rt/test/profile/lit.cfg.py | 9 +++++++++ .../Transforms/PGOProfile/gpu-profile-version.ll | 4 ++-- 3 files changed, 17 insertions(+), 8 deletions(-) rename compiler-rt/test/profile/{GPU/instrprof-version-visibility.test => instrprof-gpu-version-visibility.test} (60%) diff --git a/compiler-rt/test/profile/GPU/instrprof-version-visibility.test b/compiler-rt/test/profile/instrprof-gpu-version-visibility.test similarity index 60% rename from compiler-rt/test/profile/GPU/instrprof-version-visibility.test rename to compiler-rt/test/profile/instrprof-gpu-version-visibility.test index cf3f1dc386ad95..96c0db4fa966d1 100644 --- a/compiler-rt/test/profile/GPU/instrprof-version-visibility.test +++ b/compiler-rt/test/profile/instrprof-gpu-version-visibility.test @@ -1,10 +1,10 @@ # NOTE: Do not autogenerate. Check the runtime version definition and its device export. -# REQUIRES: system-linux, amdgpu, lld-available -# RUN: %clang -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../../lib/profile/InstrProfilingVersionVar.c -o - | FileCheck %s --check-prefix=CPU -# RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../../lib/profile/InstrProfilingVersionVar.c -o %t.ll -# RUN: FileCheck %s --check-prefix=GPU --input-file=%t.ll -# RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../../lib/profile/InstrProfiling.c -o - | FileCheck %s --check-prefix=DECL -# RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -flto -nostdlib -shared %t.ll -o %t.so +# REQUIRES: system-linux, amdgpu-registered-target, lld-available +# RUN: %clang -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../lib/profile/InstrProfilingVersionVar.c -o - | FileCheck %s --check-prefix=CPU +# RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -emit-llvm -c %S/../../lib/profile/InstrProfilingVersionVar.c -o %t.bc +# RUN: opt -S < %t.bc | FileCheck %s --check-prefix=GPU +# RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../lib/profile/InstrProfiling.c -o - | FileCheck %s --check-prefix=DECL +# RUN: ld.lld -shared %t.bc -o %t.so # RUN: llvm-readelf --dyn-syms %t.so | FileCheck %s --check-prefix=ELF # CPU: @__llvm_profile_raw_version = weak hidden global i64 diff --git a/compiler-rt/test/profile/lit.cfg.py b/compiler-rt/test/profile/lit.cfg.py index f3cd3f525a58d9..e3f9629807ea2c 100644 --- a/compiler-rt/test/profile/lit.cfg.py +++ b/compiler-rt/test/profile/lit.cfg.py @@ -2,6 +2,7 @@ import os import re +import shlex def get_required_attr(config, attr_name): @@ -50,6 +51,14 @@ def get_required_attr(config, attr_name): # What to exclude. config.excludes = ["Inputs"] +# Compiler-only GPU tests require a registered backend, independently of HIP +# runtime and GPU availability. +registered_targets = lit_config.run_command_cached( + shlex.split(config.clang) + ["--print-targets"], allow_failure=True, text=True +) +if re.search(r"^\s*amdgcn\s+-", registered_targets, re.MULTILINE): + config.available_features.add("amdgpu-registered-target") + # Clang flags. target_cflags = [get_required_attr(config, "target_cflags")] clang_cflags = target_cflags + extra_link_flags diff --git a/llvm/test/Transforms/PGOProfile/gpu-profile-version.ll b/llvm/test/Transforms/PGOProfile/gpu-profile-version.ll index f3734298b21113..9298582f6cc0d0 100644 --- a/llvm/test/Transforms/PGOProfile/gpu-profile-version.ll +++ b/llvm/test/Transforms/PGOProfile/gpu-profile-version.ll @@ -1,7 +1,7 @@ ; REQUIRES: amdgpu-registered-target, x86-registered-target ; NOTE: Do not autogenerate. Check the visibility of the shared profile version. -; RUN: opt -mtriple=amdgpu9.50-amd-amdhsa -passes=pgo-instr-gen -S %s | FileCheck %s --check-prefix=GPU -; RUN: opt -mtriple=x86_64-unknown-linux-gnu -passes=pgo-instr-gen -S %s | FileCheck %s --check-prefix=CPU +; RUN: opt -mtriple=amdgpu9.50-amd-amdhsa -passes=pgo-instr-gen -S < %s | FileCheck %s --check-prefix=GPU +; RUN: opt -mtriple=x86_64-unknown-linux-gnu -passes=pgo-instr-gen -S < %s | FileCheck %s --check-prefix=CPU ; GPU: @__llvm_profile_raw_version = protected addrspace(1) constant i64 ; CPU: @__llvm_profile_raw_version = hidden constant i64 >From 07a771349d889ec65e743bde5706d0ca51ecdcc0 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <[email protected]> Date: Mon, 2 Nov 2026 07:44:23 -0500 Subject: [PATCH 5/5] [compiler-rt][profile] Avoid libc headers in GPU visibility test Compile a minimal consumer of InstrProfiling.h for the declaration check. Compiling all of InstrProfiling.c requires string.h and target libc headers that are unavailable in the Linux CI cross-target configuration. --- compiler-rt/test/profile/instrprof-gpu-version-visibility.test | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/compiler-rt/test/profile/instrprof-gpu-version-visibility.test b/compiler-rt/test/profile/instrprof-gpu-version-visibility.test index 96c0db4fa966d1..352ee6d448ab23 100644 --- a/compiler-rt/test/profile/instrprof-gpu-version-visibility.test +++ b/compiler-rt/test/profile/instrprof-gpu-version-visibility.test @@ -3,7 +3,8 @@ # RUN: %clang -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../lib/profile/InstrProfilingVersionVar.c -o - | FileCheck %s --check-prefix=CPU # RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -emit-llvm -c %S/../../lib/profile/InstrProfilingVersionVar.c -o %t.bc # RUN: opt -S < %t.bc | FileCheck %s --check-prefix=GPU -# RUN: %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -S -emit-llvm %S/../../lib/profile/InstrProfiling.c -o - | FileCheck %s --check-prefix=DECL +# Compile only the header to avoid host libc headers when targeting AMDGPU. +# RUN: echo 'uint64_t get_version(void) { return __llvm_profile_raw_version; }' | %clang --target=amdgpu9.50-amd-amdhsa -nogpulib -ffreestanding -DCOMPILER_RT_PROFILE_BAREMETAL=1 -DCOMPILER_RT_HAS_ATOMICS=1 -include %S/../../lib/profile/InstrProfiling.h -x c -S -emit-llvm - -o - | FileCheck %s --check-prefix=DECL # RUN: ld.lld -shared %t.bc -o %t.so # RUN: llvm-readelf --dyn-syms %t.so | FileCheck %s --check-prefix=ELF _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
