[clang] [compiler-rt] [llvm] [CodeGen] Add OffloadBlockUniformityAnalysis for offload PGO (PR #178417)
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Mon Jun 1 06:39:23 PDT 2026
https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/178417
>From 1f532f6a31673857f3835040380767a824d05eb3 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Mon, 6 Apr 2026 18:58:11 -0400
Subject: [PATCH 1/2] [PGO][AMDGPU] Add uniformity-aware offload profile format
and instrumentation
Add the producer-side format and instrumentation pieces that layer on top
of the basic HIP offload PGO support.
This extends the raw/indexed profile format with wave size and uniform
counters, teaches llvm-profdata to read and merge them, and adds the AMDGPU
sampling runtime.
The uniform counters follow the same optional-section design as the MC/DC
bitmap: they live in their own section (__llvm_prf_ucnts) with header
descriptor fields that are always present, but the section is empty
(NumUniformCounters == 0) for non-GPU targets, so host profiles are
unaffected. Because this inserts a new section into the baseline layout,
the format version is bumped (raw 10->11, indexed 13->14).
---
clang/lib/CodeGen/CGCUDANV.cpp | 18 +-
compiler-rt/include/profile/InstrProfData.inc | 27 ++-
compiler-rt/lib/profile/InstrProfiling.h | 30 ++-
.../lib/profile/InstrProfilingBuffer.c | 34 +--
compiler-rt/lib/profile/InstrProfilingFile.c | 24 ++-
.../lib/profile/InstrProfilingInternal.h | 19 +-
.../lib/profile/InstrProfilingPlatformGPU.c | 48 ++++-
.../profile/InstrProfilingPlatformROCm.cpp | 102 +++++++--
.../lib/profile/InstrProfilingWriter.c | 38 +++-
llvm/include/llvm/IR/RuntimeLibcalls.td | 5 +-
llvm/include/llvm/ProfileData/InstrProf.h | 43 +++-
.../llvm/ProfileData/InstrProfData.inc | 27 ++-
.../llvm/ProfileData/InstrProfReader.h | 5 +
.../llvm/ProfileData/InstrProfWriter.h | 1 +
llvm/lib/ProfileData/InstrProf.cpp | 10 +-
llvm/lib/ProfileData/InstrProfCorrelator.cpp | 2 +
llvm/lib/ProfileData/InstrProfReader.cpp | 113 +++++++++-
llvm/lib/ProfileData/InstrProfWriter.cpp | 45 +++-
.../Instrumentation/InstrProfiling.cpp | 202 ++++++++++++++++--
.../Instrumentation/PGOInstrumentation.cpp | 31 ++-
.../AArch64/global-merge-profile-sections.ll | 4 +-
.../InstrProfiling/amdgpu-3d-grid.ll | 23 ++
...rrays.ll => amdgpu-contiguous-counters.ll} | 14 +-
.../InstrProfiling/amdgpu-instrumentation.ll | 45 +++-
.../InstrProfiling/amdgpu-uniform-counters.ll | 22 ++
.../InstrProfiling/amdgpu-wave32.ll | 35 +++
.../InstrProfiling/amdgpu-wave64.ll | 34 +++
.../InstrProfiling/coverage.ll | 8 +-
.../InstrProfiling/inline-data-var-create.ll | 23 +-
.../InstrProfiling/platform.ll | 14 ++
.../thinlto_indirect_call_promotion.profraw | Bin 544 -> 600 bytes
.../Transforms/PGOProfile/comdat_internal.ll | 4 +-
.../instrprof_burst_sampling_fast.ll | 2 +-
.../Transforms/PGOProfile/vtable_profile.ll | 2 +-
.../llvm-profdata/Inputs/c-general.profraw | Bin 2032 -> 2152 bytes
.../llvm-profdata/Inputs/compressed.profraw | Bin 1984 -> 2104 bytes
.../llvm-profdata/binary-ids-padding.test | 19 +-
.../llvm-profdata/large-binary-id-size.test | 9 +-
...alformed-not-space-for-another-header.test | 7 +-
.../malformed-num-counters-zero.test | 8 +-
.../malformed-ptr-to-counter-array.test | 8 +-
.../misaligned-binary-ids-size.test | 2 +-
.../mismatched-raw-profile-header.test | 3 +
.../tools/llvm-profdata/profile-version.test | 2 +-
.../tools/llvm-profdata/raw-32-bits-be.test | 19 +-
.../tools/llvm-profdata/raw-32-bits-le.test | 19 +-
.../tools/llvm-profdata/raw-64-bits-be.test | 21 +-
.../tools/llvm-profdata/raw-64-bits-le.test | 19 +-
.../tools/llvm-profdata/raw-two-profiles.test | 12 +-
llvm/tools/llvm-profdata/llvm-profdata.cpp | 27 +++
.../common/include/GlobalHandler.h | 14 +-
.../common/src/GlobalHandler.cpp | 4 +-
52 files changed, 1035 insertions(+), 212 deletions(-)
create mode 100644 llvm/test/Instrumentation/InstrProfiling/amdgpu-3d-grid.ll
rename llvm/test/Instrumentation/InstrProfiling/{amdgpu-profc-arrays.ll => amdgpu-contiguous-counters.ll} (52%)
create mode 100644 llvm/test/Instrumentation/InstrProfiling/amdgpu-uniform-counters.ll
create mode 100644 llvm/test/Instrumentation/InstrProfiling/amdgpu-wave32.ll
create mode 100644 llvm/test/Instrumentation/InstrProfiling/amdgpu-wave64.ll
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index a727a006969d8..518e0c3db5ed6 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -1328,12 +1328,14 @@ void CGNVCUDARuntime::createOffloadingEntries() {
}
// For HIP host+device compiles with PGO enabled, emit the per-TU global
-// __llvm_profile_sections_<CUID>. Device side: a 7-pointer struct holding
-// section start/stop bounds for the names/counters/data sections plus the
-// raw-version variable. Host side: an opaque void* shadow whose only
-// purpose is to give the host-runtime a registered symbol name to look up
-// via hipGetSymbolAddress; the actual device-side data lives in the
-// matching device-side global.
+// __llvm_profile_sections_<CUID>. Device side: a 9-pointer struct holding
+// section start/stop bounds for the names/counters/data/uniform-counters
+// sections plus the raw-version variable; the field order must match
+// INSTR_PROF_GPU_SECT in InstrProfData.inc, which the host runtime's
+// __llvm_profile_gpu_sections struct is generated from. Host side: an opaque
+// void* shadow whose only purpose is to give the host-runtime a registered
+// symbol name to look up via hipGetSymbolAddress; the actual device-side data
+// lives in the matching device-side global.
void CGNVCUDARuntime::emitOffloadProfilingSections() {
if (!CGM.getLangOpts().HIP)
return;
@@ -1381,7 +1383,7 @@ void CGNVCUDARuntime::emitOffloadProfilingSections() {
}
auto *StructTy = llvm::StructType::get(
- Ctx, {PtrTy, PtrTy, PtrTy, PtrTy, PtrTy, PtrTy, PtrTy});
+ Ctx, {PtrTy, PtrTy, PtrTy, PtrTy, PtrTy, PtrTy, PtrTy, PtrTy, PtrTy});
llvm::Constant *Fields[] = {
getOrDeclare("__start___llvm_prf_names"),
getOrDeclare("__stop___llvm_prf_names"),
@@ -1389,6 +1391,8 @@ void CGNVCUDARuntime::emitOffloadProfilingSections() {
getOrDeclare("__stop___llvm_prf_cnts"),
getOrDeclare("__start___llvm_prf_data"),
getOrDeclare("__stop___llvm_prf_data"),
+ getOrDeclare("__start___llvm_prf_ucnts"),
+ getOrDeclare("__stop___llvm_prf_ucnts"),
VersionGV,
};
auto *Init = llvm::ConstantStruct::get(StructTy, Fields);
diff --git a/compiler-rt/include/profile/InstrProfData.inc b/compiler-rt/include/profile/InstrProfData.inc
index 3117cb2fddf36..1f45eb0bc8b2c 100644
--- a/compiler-rt/include/profile/InstrProfData.inc
+++ b/compiler-rt/include/profile/InstrProfData.inc
@@ -78,6 +78,8 @@ INSTR_PROF_DATA(const uint64_t, llvm::Type::getInt64Ty(Ctx), FuncHash, \
ConstantInt::get(llvm::Type::getInt64Ty(Ctx), \
Inc->getHash()->getZExtValue()))
INSTR_PROF_DATA(const IntPtrT, IntPtrTy, CounterPtr, RelativeCounterPtr)
+INSTR_PROF_DATA(const IntPtrT, IntPtrTy, UniformCounterPtr, \
+ RelativeUniformCounterPtr)
INSTR_PROF_DATA(const IntPtrT, IntPtrTy, BitmapPtr, RelativeBitmapPtr)
/* This is used to map function pointers for the indirect call targets to
* function name hashes during the conversion from raw to merged profile
@@ -91,6 +93,10 @@ INSTR_PROF_DATA(const uint32_t, llvm::Type::getInt32Ty(Ctx), NumCounters, \
ConstantInt::get(llvm::Type::getInt32Ty(Ctx), NumCounters))
INSTR_PROF_DATA(const uint16_t, Int16ArrayTy, NumValueSites[IPVK_Last+1], \
ConstantArray::get(Int16ArrayTy, Int16ArrayVals)) \
+INSTR_PROF_DATA(const uint16_t, llvm::Type::getInt16Ty(Ctx), \
+ OffloadDeviceWaveSize, \
+ ConstantInt::get(llvm::Type::getInt16Ty(Ctx), \
+ OffloadDeviceWaveSizeVal))
INSTR_PROF_DATA(const uint32_t, llvm::Type::getInt32Ty(Ctx), NumBitmapBytes, \
ConstantInt::get(llvm::Type::getInt32Ty(Ctx), NumBitmapBytes))
#undef INSTR_PROF_DATA
@@ -168,6 +174,12 @@ INSTR_PROF_GPU_SECT(const __llvm_profile_data *, llvm::PointerType::getUnqual( \
INSTR_PROF_GPU_SECT(const __llvm_profile_data *, llvm::PointerType::getUnqual( \
Ctx), DataStop, \
ConstantPointerNull::get(llvm::PointerType::getUnqual(Ctx)))
+INSTR_PROF_GPU_SECT(char *, llvm::PointerType::getUnqual(Ctx), \
+ UniformCountersStart, \
+ ConstantPointerNull::get(llvm::PointerType::getUnqual(Ctx)))
+INSTR_PROF_GPU_SECT(char *, llvm::PointerType::getUnqual(Ctx), \
+ UniformCountersStop, \
+ ConstantPointerNull::get(llvm::PointerType::getUnqual(Ctx)))
INSTR_PROF_GPU_SECT(uint64_t *, llvm::PointerType::getUnqual(Ctx), \
VersionVar, \
ConstantPointerNull::get(llvm::PointerType::getUnqual(Ctx)))
@@ -192,6 +204,12 @@ INSTR_PROF_RAW_HEADER(uint64_t, NumCounters, NumCounters)
INSTR_PROF_RAW_HEADER(uint64_t, PaddingBytesAfterCounters, PaddingBytesAfterCounters)
INSTR_PROF_RAW_HEADER(uint64_t, NumBitmapBytes, NumBitmapBytes)
INSTR_PROF_RAW_HEADER(uint64_t, PaddingBytesAfterBitmapBytes, PaddingBytesAfterBitmapBytes)
+INSTR_PROF_RAW_HEADER(uint64_t, NumUniformCounters, NumUniformCounters)
+INSTR_PROF_RAW_HEADER(uint64_t, PaddingBytesAfterUniformCounters, PaddingBytesAfterUniformCounters)
+INSTR_PROF_RAW_HEADER(uint64_t, UniformCountersDelta,
+ UniformCountersBegin ? (uintptr_t)UniformCountersBegin -
+ (uintptr_t)DataBegin
+ : 0)
INSTR_PROF_RAW_HEADER(uint64_t, NamesSize, NamesSize)
INSTR_PROF_RAW_HEADER(uint64_t, CountersDelta,
(uintptr_t)CountersBegin - (uintptr_t)DataBegin)
@@ -356,6 +374,9 @@ INSTR_PROF_SECT_ENTRY(IPSK_data, \
INSTR_PROF_SECT_ENTRY(IPSK_cnts, \
INSTR_PROF_QUOTE(INSTR_PROF_CNTS_COMMON), \
INSTR_PROF_CNTS_COFF, "__DATA,")
+INSTR_PROF_SECT_ENTRY(IPSK_ucnts, \
+ INSTR_PROF_QUOTE(INSTR_PROF_UCNTS_COMMON), \
+ INSTR_PROF_UCNTS_COFF, "__DATA,")
INSTR_PROF_SECT_ENTRY(IPSK_bitmap, \
INSTR_PROF_QUOTE(INSTR_PROF_BITS_COMMON), \
INSTR_PROF_BITS_COFF, "__DATA,")
@@ -752,9 +773,9 @@ serializeValueProfDataFrom(ValueProfRecordClosure *Closure,
(uint64_t)'f' << 16 | (uint64_t)'R' << 8 | (uint64_t)129
/* Raw profile format version (start from 1). */
-#define INSTR_PROF_RAW_VERSION 10
+#define INSTR_PROF_RAW_VERSION 11
/* Indexed profile format version (start from 1). */
-#define INSTR_PROF_INDEX_VERSION 13
+#define INSTR_PROF_INDEX_VERSION 14
/* Coverage mapping format version (start from 0). */
#define INSTR_PROF_COVMAP_VERSION 6
@@ -803,6 +824,7 @@ serializeValueProfDataFrom(ValueProfRecordClosure *Closure,
#define INSTR_PROF_NAME_COMMON __llvm_prf_names
#define INSTR_PROF_VNAME_COMMON __llvm_prf_vns
#define INSTR_PROF_CNTS_COMMON __llvm_prf_cnts
+#define INSTR_PROF_UCNTS_COMMON __llvm_prf_ucnts
#define INSTR_PROF_BITS_COMMON __llvm_prf_bits
#define INSTR_PROF_VALS_COMMON __llvm_prf_vals
#define INSTR_PROF_VNODES_COMMON __llvm_prf_vnds
@@ -820,6 +842,7 @@ serializeValueProfDataFrom(ValueProfRecordClosure *Closure,
#define INSTR_PROF_NAME_COFF ".lprfn$M"
#define INSTR_PROF_VNAME_COFF ".lprfvn$M"
#define INSTR_PROF_CNTS_COFF ".lprfc$M"
+#define INSTR_PROF_UCNTS_COFF ".lprfuc$M"
#define INSTR_PROF_BITS_COFF ".lprfb$M"
#define INSTR_PROF_VALS_COFF ".lprfv$M"
#define INSTR_PROF_VNODES_COFF ".lprfnd$M"
diff --git a/compiler-rt/lib/profile/InstrProfiling.h b/compiler-rt/lib/profile/InstrProfiling.h
index 6fc26b59d3cdc..595c258d19abd 100644
--- a/compiler-rt/lib/profile/InstrProfiling.h
+++ b/compiler-rt/lib/profile/InstrProfiling.h
@@ -308,10 +308,12 @@ uint64_t __llvm_profile_get_vtable_section_size(const VTableProfData *Begin,
*/
int __llvm_profile_get_padding_sizes_for_counters(
uint64_t DataSize, uint64_t CountersSize, uint64_t NumBitmapBytes,
- uint64_t NamesSize, uint64_t VTableSize, uint64_t VNameSize,
- uint64_t *PaddingBytesBeforeCounters, uint64_t *PaddingBytesAfterCounters,
- uint64_t *PaddingBytesAfterBitmap, uint64_t *PaddingBytesAfterNames,
- uint64_t *PaddingBytesAfterVTable, uint64_t *PaddingBytesAfterVNames);
+ uint64_t NumUniformCounters, uint64_t NamesSize, uint64_t VTableSize,
+ uint64_t VNameSize, uint64_t *PaddingBytesBeforeCounters,
+ uint64_t *PaddingBytesAfterCounters, uint64_t *PaddingBytesAfterBitmap,
+ uint64_t *PaddingBytesAfterUniformCounters,
+ uint64_t *PaddingBytesAfterNames, uint64_t *PaddingBytesAfterVTable,
+ uint64_t *PaddingBytesAfterVNames);
/*!
* \brief Set the flag that profile data has been dumped to the file.
@@ -323,14 +325,27 @@ void __llvm_profile_set_dumped(void);
/*!
* \brief Write custom target-specific profiling data to a separate file.
- * Used by offload PGO.
+ * Used by offload PGO (HIP and OpenMP).
+ *
+ * \param Target Target triple (e.g., "amdgcn-amd-amdhsa")
+ * \param DataBegin Start of profile data records
+ * \param DataEnd End of profile data records
+ * \param CountersBegin Start of counter data
+ * \param CountersEnd End of counter data
+ * \param UniformCountersBegin Start of uniform counters (NULL if not used)
+ * \param UniformCountersEnd End of uniform counters (NULL if not used)
+ * \param NamesBegin Start of names data
+ * \param NamesEnd End of names data
+ * \param VersionOverride Profile version override (NULL to use default)
*/
int __llvm_write_custom_profile(const char *Target,
const __llvm_profile_data *DataBegin,
const __llvm_profile_data *DataEnd,
const char *CountersBegin,
- const char *CountersEnd, const char *NamesBegin,
- const char *NamesEnd,
+ const char *CountersEnd,
+ const char *UniformCountersBegin,
+ const char *UniformCountersEnd,
+ const char *NamesBegin, const char *NamesEnd,
const uint64_t *VersionOverride);
/*!
@@ -364,4 +379,5 @@ extern char INSTR_PROF_PROFILE_NAME_VAR[1]; /* __llvm_profile_filename. */
const __llvm_gcov_init_func_struct *__llvm_profile_begin_covinit();
const __llvm_gcov_init_func_struct *__llvm_profile_end_covinit();
+
#endif /* PROFILE_INSTRPROFILING_H_ */
diff --git a/compiler-rt/lib/profile/InstrProfilingBuffer.c b/compiler-rt/lib/profile/InstrProfilingBuffer.c
index b406e8db74f3f..5776a8643abcc 100644
--- a/compiler-rt/lib/profile/InstrProfilingBuffer.c
+++ b/compiler-rt/lib/profile/InstrProfilingBuffer.c
@@ -147,10 +147,12 @@ static int needsCounterPadding(void) {
COMPILER_RT_VISIBILITY
int __llvm_profile_get_padding_sizes_for_counters(
uint64_t DataSize, uint64_t CountersSize, uint64_t NumBitmapBytes,
- uint64_t NamesSize, uint64_t VTableSize, uint64_t VNameSize,
- uint64_t *PaddingBytesBeforeCounters, uint64_t *PaddingBytesAfterCounters,
- uint64_t *PaddingBytesAfterBitmapBytes, uint64_t *PaddingBytesAfterNames,
- uint64_t *PaddingBytesAfterVTable, uint64_t *PaddingBytesAfterVName) {
+ uint64_t NumUniformCounters, uint64_t NamesSize, uint64_t VTableSize,
+ uint64_t VNameSize, uint64_t *PaddingBytesBeforeCounters,
+ uint64_t *PaddingBytesAfterCounters, uint64_t *PaddingBytesAfterBitmapBytes,
+ uint64_t *PaddingBytesAfterUniformCounters,
+ uint64_t *PaddingBytesAfterNames, uint64_t *PaddingBytesAfterVTable,
+ uint64_t *PaddingBytesAfterVName) {
// Counter padding is needed only if continuous mode is enabled.
if (!needsCounterPadding()) {
*PaddingBytesBeforeCounters = 0;
@@ -158,6 +160,9 @@ int __llvm_profile_get_padding_sizes_for_counters(
__llvm_profile_get_num_padding_bytes(CountersSize);
*PaddingBytesAfterBitmapBytes =
__llvm_profile_get_num_padding_bytes(NumBitmapBytes);
+ if (PaddingBytesAfterUniformCounters != NULL)
+ *PaddingBytesAfterUniformCounters = __llvm_profile_get_num_padding_bytes(
+ NumUniformCounters * sizeof(uint64_t));
*PaddingBytesAfterNames = __llvm_profile_get_num_padding_bytes(NamesSize);
if (PaddingBytesAfterVTable != NULL)
*PaddingBytesAfterVTable =
@@ -179,6 +184,8 @@ int __llvm_profile_get_padding_sizes_for_counters(
*PaddingBytesAfterCounters = calculateBytesNeededToPageAlign(CountersSize);
*PaddingBytesAfterBitmapBytes =
calculateBytesNeededToPageAlign(NumBitmapBytes);
+ if (PaddingBytesAfterUniformCounters != NULL)
+ *PaddingBytesAfterUniformCounters = 0;
*PaddingBytesAfterNames = calculateBytesNeededToPageAlign(NamesSize);
// Set these two variables to zero to avoid uninitialized variables
// even if VTableSize and VNameSize are known to be zero.
@@ -212,20 +219,22 @@ uint64_t __llvm_profile_get_size_for_buffer_internal(
* the names. */
uint64_t PaddingBytesBeforeCounters, PaddingBytesAfterCounters,
PaddingBytesAfterNames, PaddingBytesAfterBitmapBytes,
- PaddingBytesAfterVTable, PaddingBytesAfterVNames;
+ PaddingBytesAfterUniformCounters, PaddingBytesAfterVTable,
+ PaddingBytesAfterVNames;
__llvm_profile_get_padding_sizes_for_counters(
- DataSize, CountersSize, NumBitmapBytes, NamesSize, 0 /* VTableSize */,
- 0 /* VNameSize */, &PaddingBytesBeforeCounters,
- &PaddingBytesAfterCounters, &PaddingBytesAfterBitmapBytes,
+ DataSize, CountersSize, NumBitmapBytes, 0 /* NumUniformCounters */,
+ NamesSize, 0 /* VTableSize */, 0 /* VNameSize */,
+ &PaddingBytesBeforeCounters, &PaddingBytesAfterCounters,
+ &PaddingBytesAfterBitmapBytes, &PaddingBytesAfterUniformCounters,
&PaddingBytesAfterNames, &PaddingBytesAfterVTable,
&PaddingBytesAfterVNames);
return sizeof(__llvm_profile_header) + __llvm_write_binary_ids(NULL) +
DataSize + PaddingBytesBeforeCounters + CountersSize +
PaddingBytesAfterCounters + NumBitmapBytes +
- PaddingBytesAfterBitmapBytes + NamesSize + PaddingBytesAfterNames +
- VTableSize + PaddingBytesAfterVTable + VNameSize +
- PaddingBytesAfterVNames;
+ PaddingBytesAfterBitmapBytes + PaddingBytesAfterUniformCounters +
+ NamesSize + PaddingBytesAfterNames + VTableSize +
+ PaddingBytesAfterVTable + VNameSize + PaddingBytesAfterVNames;
}
COMPILER_RT_VISIBILITY
@@ -250,7 +259,8 @@ COMPILER_RT_VISIBILITY int __llvm_profile_write_buffer_internal(
// Set virtual table arguments to NULL since they are not supported yet.
return lprofWriteDataImpl(
&BufferWriter, DataBegin, DataEnd, CountersBegin, CountersEnd,
- BitmapBegin, BitmapEnd, /*VPDataReader=*/0, NamesBegin, NamesEnd,
+ BitmapBegin, BitmapEnd, /*UniformCountersBegin=*/NULL,
+ /*UniformCountersEnd=*/NULL, /*VPDataReader=*/0, NamesBegin, NamesEnd,
/*VTableBegin=*/NULL, /*VTableEnd=*/NULL, /*VNamesBegin=*/NULL,
/*VNamesEnd=*/NULL, /*SkipNameDataWrite=*/0,
__llvm_profile_get_version());
diff --git a/compiler-rt/lib/profile/InstrProfilingFile.c b/compiler-rt/lib/profile/InstrProfilingFile.c
index 9ea5a2638fac9..99cb601b9ce65 100644
--- a/compiler-rt/lib/profile/InstrProfilingFile.c
+++ b/compiler-rt/lib/profile/InstrProfilingFile.c
@@ -168,11 +168,13 @@ static int mmapForContinuousMode(uint64_t CurrentFileOffset, FILE *File) {
* after the names. */
uint64_t PaddingBytesBeforeCounters, PaddingBytesAfterCounters,
PaddingBytesAfterNames, PaddingBytesAfterBitmapBytes,
- PaddingBytesAfterVTable, PaddingBytesAfterVNames;
+ PaddingBytesAfterUniformCounters, PaddingBytesAfterVTable,
+ PaddingBytesAfterVNames;
__llvm_profile_get_padding_sizes_for_counters(
- DataSize, CountersSize, NumBitmapBytes, NamesSize, /*VTableSize=*/0,
- /*VNameSize=*/0, &PaddingBytesBeforeCounters, &PaddingBytesAfterCounters,
- &PaddingBytesAfterBitmapBytes, &PaddingBytesAfterNames,
+ DataSize, CountersSize, NumBitmapBytes, 0 /* NumUniformCounters */,
+ NamesSize, /*VTableSize=*/0, /*VNameSize=*/0, &PaddingBytesBeforeCounters,
+ &PaddingBytesAfterCounters, &PaddingBytesAfterBitmapBytes,
+ &PaddingBytesAfterUniformCounters, &PaddingBytesAfterNames,
&PaddingBytesAfterVTable, &PaddingBytesAfterVNames);
uint64_t PageAlignedCountersLength = CountersSize + PaddingBytesAfterCounters;
@@ -1316,8 +1318,10 @@ int __llvm_write_custom_profile(const char *Target,
const __llvm_profile_data *DataBegin,
const __llvm_profile_data *DataEnd,
const char *CountersBegin,
- const char *CountersEnd, const char *NamesBegin,
- const char *NamesEnd,
+ const char *CountersEnd,
+ const char *UniformCountersBegin,
+ const char *UniformCountersEnd,
+ const char *NamesBegin, const char *NamesEnd,
const uint64_t *VersionOverride) {
int ReturnValue = 0, FilenameLength, TargetLength;
char *FilenameBuf, *TargetFilename;
@@ -1405,10 +1409,10 @@ int __llvm_write_custom_profile(const char *Target,
Version = *VersionOverride;
/* Write custom data to the file */
- ReturnValue =
- lprofWriteDataImpl(&fileWriter, DataBegin, DataEnd, CountersBegin,
- CountersEnd, NULL, NULL, lprofGetVPDataReader(), NULL,
- NULL, NULL, NULL, NamesBegin, NamesEnd, 0, Version);
+ ReturnValue = lprofWriteDataImpl(
+ &fileWriter, DataBegin, DataEnd, CountersBegin, CountersEnd, NULL, NULL,
+ UniformCountersBegin, UniformCountersEnd, lprofGetVPDataReader(), NULL,
+ NULL, NULL, NULL, NamesBegin, NamesEnd, 0, Version);
closeFileObject(OutputFile);
// Restore SIGKILL.
diff --git a/compiler-rt/lib/profile/InstrProfilingInternal.h b/compiler-rt/lib/profile/InstrProfilingInternal.h
index 5647782527eb7..61485915bddad 100644
--- a/compiler-rt/lib/profile/InstrProfilingInternal.h
+++ b/compiler-rt/lib/profile/InstrProfilingInternal.h
@@ -152,16 +152,15 @@ typedef struct VPDataReaderType {
the name data is already in destination, we just skip over it. */
int lprofWriteData(ProfDataWriter *Writer, VPDataReaderType *VPDataReader,
int SkipNameDataWrite);
-int lprofWriteDataImpl(ProfDataWriter *Writer,
- const __llvm_profile_data *DataBegin,
- const __llvm_profile_data *DataEnd,
- const char *CountersBegin, const char *CountersEnd,
- const char *BitmapBegin, const char *BitmapEnd,
- VPDataReaderType *VPDataReader, const char *NamesBegin,
- const char *NamesEnd, const VTableProfData *VTableBegin,
- const VTableProfData *VTableEnd, const char *VNamesBegin,
- const char *VNamesEnd, int SkipNameDataWrite,
- uint64_t Version);
+int lprofWriteDataImpl(
+ ProfDataWriter *Writer, const __llvm_profile_data *DataBegin,
+ const __llvm_profile_data *DataEnd, const char *CountersBegin,
+ const char *CountersEnd, const char *BitmapBegin, const char *BitmapEnd,
+ const char *UniformCountersBegin, const char *UniformCountersEnd,
+ VPDataReaderType *VPDataReader, const char *NamesBegin,
+ const char *NamesEnd, const VTableProfData *VTableBegin,
+ const VTableProfData *VTableEnd, const char *VNamesBegin,
+ const char *VNamesEnd, int SkipNameDataWrite, uint64_t Version);
/* Merge value profile data pointed to by SrcValueProfData into
* in-memory profile counters pointed by to DstData. */
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformGPU.c b/compiler-rt/lib/profile/InstrProfilingPlatformGPU.c
index 3c67e2c7089d4..eb1507ecd25cb 100644
--- a/compiler-rt/lib/profile/InstrProfilingPlatformGPU.c
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformGPU.c
@@ -16,6 +16,7 @@
#include "InstrProfiling.h"
#include <gpuintrin.h>
+#include <stdint.h>
// Symbols exported to the GPU runtime need to be visible in the .dynsym table.
#define COMPILER_RT_GPU_VISIBILITY __attribute__((visibility("protected")))
@@ -42,6 +43,34 @@ COMPILER_RT_VISIBILITY void INSTR_PROF_INSTRUMENT_GPU_FUNC(uint64_t *counter,
}
}
+// Block-level sampling for offload PGO. For GPU kernels with stationary
+// behavior (where all thread blocks execute the same code paths regardless of
+// block ID), partial sampling significantly reduces instrumentation overhead
+// without losing PGO performance gains.
+//
+// Returns 1 if this block should be instrumented, 0 to skip. Samples by
+// matching lower bits of the x-dimension block ID to zero.
+// sampling_bits=0: all blocks (100%)
+// sampling_bits=3: every 8th block in x (12.5%, default)
+//
+// Note: We use only block_id_x rather than a fully linearized 3D block ID.
+// The 3D linearization requires __gpu_num_blocks_x/y which expands to
+// __builtin_amdgcn_workgroup_size_x/y. With -mcode-object-version=none (used
+// to build compiler-rt profile runtime), the compiler emits a load of
+// __oclc_ABI_version to select the correct ABI path. Since the profile runtime
+// is linked after device libs are internalized, __oclc_ABI_version is no longer
+// available. Using block_id_x directly avoids this dependency. For typical
+// kernels with large 1D or x-dominant grids this is sufficient; blocks sharing
+// the same x-index are sampled together in 3D grids (minor uniformity loss).
+COMPILER_RT_VISIBILITY int __llvm_profile_sampling_gpu(uint32_t sampling_bits) {
+ if (sampling_bits == 0)
+ return 1;
+
+ uint32_t block_id = __gpu_block_id_x();
+ uint32_t mask = (1u << sampling_bits) - 1;
+ return (block_id & mask) == 0;
+}
+
#if defined(__AMDGPU__)
#define PROF_NAME_START INSTR_PROF_SECT_START(INSTR_PROF_NAME_COMMON)
@@ -50,6 +79,8 @@ COMPILER_RT_VISIBILITY void INSTR_PROF_INSTRUMENT_GPU_FUNC(uint64_t *counter,
#define PROF_CNTS_STOP INSTR_PROF_SECT_STOP(INSTR_PROF_CNTS_COMMON)
#define PROF_DATA_START INSTR_PROF_SECT_START(INSTR_PROF_DATA_COMMON)
#define PROF_DATA_STOP INSTR_PROF_SECT_STOP(INSTR_PROF_DATA_COMMON)
+#define PROF_UCNTS_START INSTR_PROF_SECT_START(INSTR_PROF_UCNTS_COMMON)
+#define PROF_UCNTS_STOP INSTR_PROF_SECT_STOP(INSTR_PROF_UCNTS_COMMON)
extern char PROF_NAME_START[] COMPILER_RT_VISIBILITY COMPILER_RT_WEAK;
extern char PROF_NAME_STOP[] COMPILER_RT_VISIBILITY COMPILER_RT_WEAK;
@@ -59,17 +90,15 @@ extern __llvm_profile_data PROF_DATA_START[] COMPILER_RT_VISIBILITY
COMPILER_RT_WEAK;
extern __llvm_profile_data PROF_DATA_STOP[] COMPILER_RT_VISIBILITY
COMPILER_RT_WEAK;
+extern char PROF_UCNTS_START[] COMPILER_RT_VISIBILITY COMPILER_RT_WEAK;
+extern char PROF_UCNTS_STOP[] COMPILER_RT_VISIBILITY COMPILER_RT_WEAK;
// AMDGPU is a proper ELF target and exports the linker-defined section bounds.
COMPILER_RT_GPU_VISIBILITY
__llvm_profile_gpu_sections INSTR_PROF_SECT_BOUNDS_TABLE = {
- PROF_NAME_START,
- PROF_NAME_STOP,
- PROF_CNTS_START,
- PROF_CNTS_STOP,
- PROF_DATA_START,
- PROF_DATA_STOP,
- &INSTR_PROF_RAW_VERSION_VAR};
+ PROF_NAME_START, PROF_NAME_STOP, PROF_CNTS_START,
+ PROF_CNTS_STOP, PROF_DATA_START, PROF_DATA_STOP,
+ PROF_UCNTS_START, PROF_UCNTS_STOP, &INSTR_PROF_RAW_VERSION_VAR};
#elif defined(__NVPTX__)
@@ -79,8 +108,9 @@ __llvm_profile_gpu_sections INSTR_PROF_SECT_BOUNDS_TABLE = {
// globals next to each other.
COMPILER_RT_GPU_VISIBILITY
__llvm_profile_gpu_sections INSTR_PROF_SECT_BOUNDS_TABLE = {
- NULL, NULL, NULL, NULL, NULL, NULL, &INSTR_PROF_RAW_VERSION_VAR};
-
+ NULL, NULL, NULL,
+ NULL, NULL, NULL,
+ NULL, NULL, &INSTR_PROF_RAW_VERSION_VAR};
#endif
#endif
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
index ee00c572e3a42..410fcfdd8d0c8 100644
--- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
@@ -560,29 +560,35 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
const void *DevCntsBegin = HostSections.CountersStart;
const void *DevDataBegin = HostSections.DataStart;
const void *DevNamesBegin = HostSections.NamesStart;
+ const void *DevUniformCntsBegin = HostSections.UniformCountersStart;
const void *DevCntsEnd = HostSections.CountersStop;
const void *DevDataEnd = HostSections.DataStop;
const void *DevNamesEnd = HostSections.NamesStop;
+ const void *DevUniformCntsEnd = HostSections.UniformCountersStop;
size_t CountersSize = (const char *)DevCntsEnd - (const char *)DevCntsBegin;
size_t DataSize = (const char *)DevDataEnd - (const char *)DevDataBegin;
size_t NamesSize = (const char *)DevNamesEnd - (const char *)DevNamesBegin;
+ size_t UniformCountersSize =
+ (const char *)DevUniformCntsEnd - (const char *)DevUniformCntsBegin;
if (isVerboseMode())
PROF_NOTE("Section pointers: Cnts=[%p,%p]=%zu Data=[%p,%p]=%zu "
- "Names=[%p,%p]=%zu\n",
+ "Names=[%p,%p]=%zu UCnts=[%p,%p]=%zu\n",
DevCntsBegin, DevCntsEnd, CountersSize, DevDataBegin, DevDataEnd,
- DataSize, DevNamesBegin, DevNamesEnd, NamesSize);
+ DataSize, DevNamesBegin, DevNamesEnd, NamesSize,
+ DevUniformCntsBegin, DevUniformCntsEnd, UniformCountersSize);
if (CountersSize == 0 || DataSize == 0)
return 0;
int ret = -1;
- int NamesReused = 0, CntsReused = 0, DataReused = 0;
+ int NamesReused = 0, CntsReused = 0, UCntsReused = 0, DataReused = 0;
char *HostDataBegin = nullptr;
char *HostCountersBegin = nullptr;
char *HostNamesBegin = nullptr;
+ char *HostUniformCountersBegin = nullptr;
/* Sections using linker-defined __start_/__stop_ bounds are shared across
TU structs in RDC mode. Deduplicate by caching the last copied range. */
@@ -598,8 +604,12 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
static char *CachedHostData = nullptr;
static size_t CachedDataSize = 0;
+ static const void *CachedDevUCntsBegin = nullptr;
+ static char *CachedHostUCnts = nullptr;
+ static size_t CachedUCntsSize = 0;
+
// Owns freshly malloc'd buffers; release() transfers ownership to the cache.
- UniqueFree CntsOwner, DataOwner, NamesOwner;
+ UniqueFree CntsOwner, DataOwner, NamesOwner, UCntsOwner;
if (CountersSize > 0 && DevCntsBegin == CachedDevCntsBegin &&
CountersSize == CachedCntsSize) {
@@ -634,9 +644,22 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
NamesOwner.reset(HostNamesBegin);
}
+ if (UniformCountersSize > 0 && DevUniformCntsBegin == CachedDevUCntsBegin &&
+ UniformCountersSize == CachedUCntsSize) {
+ HostUniformCountersBegin = CachedHostUCnts;
+ UCntsReused = 1;
+ if (isVerboseMode())
+ PROF_NOTE("Reusing cached ucnts section (%zu bytes)\n",
+ UniformCountersSize);
+ } else if (UniformCountersSize > 0) {
+ HostUniformCountersBegin = (char *)malloc(UniformCountersSize);
+ UCntsOwner.reset(HostUniformCountersBegin);
+ }
+
if ((DataSize > 0 && !HostDataBegin) ||
(CountersSize > 0 && !HostCountersBegin) ||
- (NamesSize > 0 && !HostNamesBegin)) {
+ (NamesSize > 0 && !HostNamesBegin) ||
+ (UniformCountersSize > 0 && !HostUniformCountersBegin)) {
PROF_ERR("%s\n", "failed to allocate host memory for device sections");
return -1;
}
@@ -647,7 +670,10 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
memcpyDeviceToHost(HostCountersBegin, DevCntsBegin, CountersSize) !=
0) ||
(NamesSize > 0 && !NamesReused &&
- memcpyDeviceToHost(HostNamesBegin, DevNamesBegin, NamesSize) != 0)) {
+ memcpyDeviceToHost(HostNamesBegin, DevNamesBegin, NamesSize) != 0) ||
+ (UniformCountersSize > 0 && !UCntsReused &&
+ memcpyDeviceToHost(HostUniformCountersBegin, DevUniformCntsBegin,
+ UniformCountersSize) != 0)) {
PROF_ERR("%s\n", "failed to copy profile sections from device");
return -1;
}
@@ -672,25 +698,34 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
CachedNamesSize = NamesSize;
NamesOwner.release();
}
+ if (!UCntsReused && UniformCountersSize > 0) {
+ CachedDevUCntsBegin = DevUniformCntsBegin;
+ CachedHostUCnts = HostUniformCountersBegin;
+ CachedUCntsSize = UniformCountersSize;
+ UCntsOwner.release();
+ }
if (isVerboseMode())
- PROF_NOTE("Copied device sections: Counters=%zu, Data=%zu, Names=%zu\n",
- CountersSize, DataSize, NamesSize);
+ PROF_NOTE("Copied device sections: Counters=%zu, Data=%zu, Names=%zu, "
+ "UniformCounters=%zu\n",
+ CountersSize, DataSize, NamesSize, UniformCountersSize);
// Arrange buffer as [Data][Padding][Counters][Names] to match the layout
// expected by lprofWriteDataImpl (CountersDelta = CountersBegin - DataBegin).
const uint64_t NumData = DataSize / sizeof(__llvm_profile_data);
const uint64_t NumBitmapBytes = 0;
+ const uint64_t NumUniformCounters = UniformCountersSize / sizeof(uint64_t);
const uint64_t VTableSectionSize = 0;
const uint64_t VNamesSize = 0;
uint64_t PaddingBytesBeforeCounters, PaddingBytesAfterCounters,
- PaddingBytesAfterBitmapBytes, PaddingBytesAfterNames,
- PaddingBytesAfterVTable, PaddingBytesAfterVNames;
+ PaddingBytesAfterBitmapBytes, PaddingBytesAfterUniformCounters,
+ PaddingBytesAfterNames, PaddingBytesAfterVTable, PaddingBytesAfterVNames;
if (__llvm_profile_get_padding_sizes_for_counters(
- DataSize, CountersSize, NumBitmapBytes, NamesSize, VTableSectionSize,
- VNamesSize, &PaddingBytesBeforeCounters, &PaddingBytesAfterCounters,
- &PaddingBytesAfterBitmapBytes, &PaddingBytesAfterNames,
+ DataSize, CountersSize, NumBitmapBytes, NumUniformCounters, NamesSize,
+ VTableSectionSize, VNamesSize, &PaddingBytesBeforeCounters,
+ &PaddingBytesAfterCounters, &PaddingBytesAfterBitmapBytes,
+ &PaddingBytesAfterUniformCounters, &PaddingBytesAfterNames,
&PaddingBytesAfterVTable, &PaddingBytesAfterVNames) != 0) {
PROF_ERR("%s\n", "failed to get padding sizes");
return -1;
@@ -715,19 +750,25 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
__builtin_memcpy(BufCountersBegin, HostCountersBegin, CountersSize);
__builtin_memcpy(BufNamesBegin, HostNamesBegin, NamesSize);
- // CounterPtr is a device-relative offset; relocate it for the file layout
- // where the Data section precedes Counters.
+ // CounterPtr and UniformCounterPtr are device-relative offsets; relocate
+ // them for the file layout where the Data section precedes the Counters and
+ // UniformCounters sections. Uniform counters are copied in linker (section)
+ // order and located via their relative pointer, exactly like the regular
+ // counters: llvm-profdata reads them through UniformCounterPtr (decrementing
+ // UniformCountersDelta per record, just like CountersDelta) and does not
+ // assume data-record order, so no reordering is needed.
+ ptrdiff_t UCFileOffset =
+ DataSize + PaddingBytesBeforeCounters + CountersSize +
+ PaddingBytesAfterCounters + NumBitmapBytes + PaddingBytesAfterBitmapBytes;
__llvm_profile_data *RelocatedData = (__llvm_profile_data *)BufDataBegin;
for (uint64_t i = 0; i < NumData; ++i) {
+ const char *DeviceDataStructAddr =
+ (const char *)DevDataBegin + (i * sizeof(__llvm_profile_data));
if (RelocatedData[i].CounterPtr) {
- ptrdiff_t DeviceCounterPtrOffset = (ptrdiff_t)RelocatedData[i].CounterPtr;
- const char *DeviceDataStructAddr =
- (const char *)DevDataBegin + (i * sizeof(__llvm_profile_data));
const char *DeviceCountersAddr =
- DeviceDataStructAddr + DeviceCounterPtrOffset;
+ DeviceDataStructAddr + (ptrdiff_t)RelocatedData[i].CounterPtr;
ptrdiff_t OffsetIntoCountersSection =
DeviceCountersAddr - (const char *)DevCntsBegin;
-
ptrdiff_t NewRelativeOffset = DataSize + PaddingBytesBeforeCounters +
OffsetIntoCountersSection -
(i * sizeof(__llvm_profile_data));
@@ -735,6 +776,21 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
offsetof(__llvm_profile_data, CounterPtr),
&NewRelativeOffset, sizeof(NewRelativeOffset));
}
+ if (HostUniformCountersBegin && RelocatedData[i].UniformCounterPtr) {
+ const char *DeviceUCAddr =
+ DeviceDataStructAddr + (ptrdiff_t)RelocatedData[i].UniformCounterPtr;
+ ptrdiff_t OffsetIntoUCSection =
+ DeviceUCAddr - (const char *)DevUniformCntsBegin;
+ ptrdiff_t NewUCRelativeOffset = UCFileOffset + OffsetIntoUCSection -
+ (i * sizeof(__llvm_profile_data));
+ __builtin_memcpy((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
+ offsetof(__llvm_profile_data, UniformCounterPtr),
+ &NewUCRelativeOffset, sizeof(NewUCRelativeOffset));
+ } else {
+ __builtin_memset((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
+ offsetof(__llvm_profile_data, UniformCounterPtr),
+ 0, sizeof(RelocatedData[i].UniformCounterPtr));
+ }
__builtin_memset((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
offsetof(__llvm_profile_data, BitmapPtr),
0,
@@ -749,8 +805,10 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
ret = __llvm_write_custom_profile(
Target, (__llvm_profile_data *)BufDataBegin,
(__llvm_profile_data *)(BufDataBegin + DataSize), BufCountersBegin,
- BufCountersBegin + CountersSize, BufNamesBegin, BufNamesBegin + NamesSize,
- nullptr);
+ BufCountersBegin + CountersSize, HostUniformCountersBegin,
+ HostUniformCountersBegin ? HostUniformCountersBegin + UniformCountersSize
+ : nullptr,
+ BufNamesBegin, BufNamesBegin + NamesSize, nullptr);
if (ret != 0) {
PROF_ERR("%s\n", "failed to write device profile using shared API");
diff --git a/compiler-rt/lib/profile/InstrProfilingWriter.c b/compiler-rt/lib/profile/InstrProfilingWriter.c
index 633fdb9661162..2c459acd39af0 100644
--- a/compiler-rt/lib/profile/InstrProfilingWriter.c
+++ b/compiler-rt/lib/profile/InstrProfilingWriter.c
@@ -256,7 +256,9 @@ COMPILER_RT_VISIBILITY int lprofWriteData(ProfDataWriter *Writer,
const char *VNamesEnd = __llvm_profile_end_vtabnames();
uint64_t Version = __llvm_profile_get_version();
return lprofWriteDataImpl(Writer, DataBegin, DataEnd, CountersBegin,
- CountersEnd, BitmapBegin, BitmapEnd, VPDataReader,
+ CountersEnd, BitmapBegin, BitmapEnd,
+ /*UniformCountersBegin=*/NULL,
+ /*UniformCountersEnd=*/NULL, VPDataReader,
NamesBegin, NamesEnd, VTableBegin, VTableEnd,
VNamesBegin, VNamesEnd, SkipNameDataWrite, Version);
}
@@ -265,6 +267,7 @@ COMPILER_RT_VISIBILITY int lprofWriteDataImpl(
ProfDataWriter *Writer, const __llvm_profile_data *DataBegin,
const __llvm_profile_data *DataEnd, const char *CountersBegin,
const char *CountersEnd, const char *BitmapBegin, const char *BitmapEnd,
+ const char *UniformCountersBegin, const char *UniformCountersEnd,
VPDataReaderType *VPDataReader, const char *NamesBegin,
const char *NamesEnd, const VTableProfData *VTableBegin,
const VTableProfData *VTableEnd, const char *VNamesBegin,
@@ -286,6 +289,13 @@ COMPILER_RT_VISIBILITY int lprofWriteDataImpl(
__llvm_profile_get_vtable_section_size(VTableBegin, VTableEnd);
const uint64_t VNamesSize =
__llvm_profile_get_name_size(VNamesBegin, VNamesEnd);
+ const uint64_t NumUniformCounters =
+ (UniformCountersBegin && UniformCountersEnd &&
+ UniformCountersEnd > UniformCountersBegin)
+ ? (UniformCountersEnd - UniformCountersBegin) / sizeof(uint64_t)
+ : 0;
+ const uint64_t UniformCountersSectionSize =
+ NumUniformCounters * sizeof(uint64_t);
/* Create the header. */
__llvm_profile_header Header;
@@ -293,12 +303,13 @@ COMPILER_RT_VISIBILITY int lprofWriteDataImpl(
/* Determine how much padding is needed before/after the counters and after
* the names. */
uint64_t PaddingBytesBeforeCounters, PaddingBytesAfterCounters,
- PaddingBytesAfterBitmapBytes, PaddingBytesAfterNames,
- PaddingBytesAfterVTable, PaddingBytesAfterVNames;
+ PaddingBytesAfterBitmapBytes, PaddingBytesAfterUniformCounters,
+ PaddingBytesAfterNames, PaddingBytesAfterVTable, PaddingBytesAfterVNames;
if (__llvm_profile_get_padding_sizes_for_counters(
- DataSectionSize, CountersSectionSize, NumBitmapBytes, NamesSize,
- VTableSectionSize, VNamesSize, &PaddingBytesBeforeCounters,
- &PaddingBytesAfterCounters, &PaddingBytesAfterBitmapBytes,
+ DataSectionSize, CountersSectionSize, NumBitmapBytes,
+ NumUniformCounters, NamesSize, VTableSectionSize, VNamesSize,
+ &PaddingBytesBeforeCounters, &PaddingBytesAfterCounters,
+ &PaddingBytesAfterBitmapBytes, &PaddingBytesAfterUniformCounters,
&PaddingBytesAfterNames, &PaddingBytesAfterVTable,
&PaddingBytesAfterVNames) == -1)
return -1;
@@ -315,12 +326,25 @@ COMPILER_RT_VISIBILITY int lprofWriteDataImpl(
#ifdef _WIN64
Header.CountersDelta = (uint32_t)Header.CountersDelta;
Header.BitmapDelta = (uint32_t)Header.BitmapDelta;
+ Header.UniformCountersDelta = (uint32_t)Header.UniformCountersDelta;
#endif
+ /* Recompute UniformCountersDelta from file layout. The macro initializer
+ uses in-memory pointer arithmetic which is wrong when sections are in
+ separate allocations (e.g., custom profiles from device PGO). */
+ if (NumUniformCounters > 0)
+ Header.UniformCountersDelta = DataSectionSize + PaddingBytesBeforeCounters +
+ CountersSectionSize +
+ PaddingBytesAfterCounters + NumBitmapBytes +
+ PaddingBytesAfterBitmapBytes;
+ else
+ Header.UniformCountersDelta = 0;
+
/* The data and names sections are omitted in lightweight mode. */
if (NumData == 0 && NamesSize == 0) {
Header.CountersDelta = 0;
Header.NamesDelta = 0;
+ Header.UniformCountersDelta = 0;
}
/* Write the profile header. */
@@ -340,6 +364,8 @@ COMPILER_RT_VISIBILITY int lprofWriteDataImpl(
{NULL, sizeof(uint8_t), PaddingBytesAfterCounters, 1},
{BitmapBegin, sizeof(uint8_t), NumBitmapBytes, 0},
{NULL, sizeof(uint8_t), PaddingBytesAfterBitmapBytes, 1},
+ {UniformCountersBegin, sizeof(uint8_t), UniformCountersSectionSize, 0},
+ {NULL, sizeof(uint8_t), PaddingBytesAfterUniformCounters, 1},
{SkipNameDataWrite ? NULL : NamesBegin, sizeof(uint8_t), NamesSize, 0},
{NULL, sizeof(uint8_t), PaddingBytesAfterNames, 1},
{VTableBegin, sizeof(uint8_t), VTableSectionSize, 0},
diff --git a/llvm/include/llvm/IR/RuntimeLibcalls.td b/llvm/include/llvm/IR/RuntimeLibcalls.td
index 37bad559f49e7..a703fe773ef88 100644
--- a/llvm/include/llvm/IR/RuntimeLibcalls.td
+++ b/llvm/include/llvm/IR/RuntimeLibcalls.td
@@ -502,6 +502,7 @@ def RETURN_ADDRESS : RuntimeLibcall;
// GPU profiling
def PROFILE_INSTRUMENT_GPU : RuntimeLibcall;
+def PROFILE_SAMPLING_GPU : RuntimeLibcall;
// Clear cache
def CLEAR_CACHE : RuntimeLibcall;
@@ -2245,9 +2246,11 @@ def WindowsARM64ECSystemLibrary
def isAMDGPU : RuntimeLibcallPredicate<"TT.isAMDGPU()">;
def __llvm_profile_instrument_gpu : RuntimeLibcallImpl<PROFILE_INSTRUMENT_GPU>;
+def __llvm_profile_sampling_gpu : RuntimeLibcallImpl<PROFILE_SAMPLING_GPU>;
def AMDGPUSystemLibrary
- : SystemRuntimeLibrary<isAMDGPU, (add __llvm_profile_instrument_gpu)>;
+ : SystemRuntimeLibrary<isAMDGPU, (add __llvm_profile_instrument_gpu,
+ __llvm_profile_sampling_gpu)>;
//===----------------------------------------------------------------------===//
// ARM Runtime Libcalls
diff --git a/llvm/include/llvm/ProfileData/InstrProf.h b/llvm/include/llvm/ProfileData/InstrProf.h
index b7f917ec39b70..7bf3932074ebc 100644
--- a/llvm/include/llvm/ProfileData/InstrProf.h
+++ b/llvm/include/llvm/ProfileData/InstrProf.h
@@ -902,6 +902,16 @@ struct InstrProfValueSiteRecord {
struct InstrProfRecord {
std::vector<uint64_t> Counts;
std::vector<uint8_t> BitmapBytes;
+ /// For AMDGPU offload profiling: raw uniform counters embedded in the
+ /// profraw file. One uint64_t per instrumented block, tracking wave-uniform
+ /// execution counts on the GPU.
+ std::vector<uint64_t> UniformCounts;
+ /// For AMDGPU offload profiling: 1 bit per basic block indicating whether
+ /// the block is entered via a wave-uniform branch. Set during merge when
+ /// per-slot counters are reduced. If a counter value is a multiple of the
+ /// wave size, the branch is considered wave-uniform.
+ std::vector<uint8_t> UniformityBits;
+ uint16_t OffloadDeviceWaveSize = 0;
InstrProfRecord() = default;
InstrProfRecord(std::vector<uint64_t> Counts) : Counts(std::move(Counts)) {}
@@ -911,6 +921,8 @@ struct InstrProfRecord {
InstrProfRecord(InstrProfRecord &&) = default;
InstrProfRecord(const InstrProfRecord &RHS)
: Counts(RHS.Counts), BitmapBytes(RHS.BitmapBytes),
+ UniformCounts(RHS.UniformCounts), UniformityBits(RHS.UniformityBits),
+ OffloadDeviceWaveSize(RHS.OffloadDeviceWaveSize),
ValueData(RHS.ValueData
? std::make_unique<ValueProfData>(*RHS.ValueData)
: nullptr) {}
@@ -918,6 +930,8 @@ struct InstrProfRecord {
InstrProfRecord &operator=(const InstrProfRecord &RHS) {
Counts = RHS.Counts;
BitmapBytes = RHS.BitmapBytes;
+ UniformityBits = RHS.UniformityBits;
+ OffloadDeviceWaveSize = RHS.OffloadDeviceWaveSize;
if (!RHS.ValueData) {
ValueData = nullptr;
return *this;
@@ -929,6 +943,17 @@ struct InstrProfRecord {
return *this;
}
+ /// Check if a basic block is entered via a wave-uniform branch.
+ /// Returns true if uniform (safe for PGO spill optimization) or if no
+ /// uniformity data is available (conservative default).
+ bool isBlockUniform(unsigned BlockIdx) const {
+ if (UniformityBits.empty())
+ return true; // No uniformity data, assume uniform (conservative)
+ if (BlockIdx / 8 >= UniformityBits.size())
+ return true; // Out of range, assume uniform
+ return (UniformityBits[BlockIdx / 8] >> (BlockIdx % 8)) & 1;
+ }
+
/// Return the number of value profile kinds with non-zero number
/// of profile sites.
inline uint32_t getNumValueKinds() const;
@@ -953,8 +978,12 @@ struct InstrProfRecord {
/// Merge the counts in \p Other into this one.
/// Optionally scale merged counts by \p Weight.
+ /// If \p WaveSize is non-zero and Other has offload profiling slots,
+ /// compute uniformity bits based on whether counter values are multiples
+ /// of WaveSize.
LLVM_ABI void merge(InstrProfRecord &Other, uint64_t Weight,
- function_ref<void(instrprof_error)> Warn);
+ function_ref<void(instrprof_error)> Warn,
+ unsigned WaveSize = 0);
/// Scale up profile counts (including value profile data) by
/// a factor of (N / D).
@@ -1079,6 +1108,14 @@ struct NamedInstrProfRecord : InstrProfRecord {
std::vector<uint8_t> BitmapBytes)
: InstrProfRecord(std::move(Counts), std::move(BitmapBytes)), Name(Name),
Hash(Hash) {}
+ NamedInstrProfRecord(StringRef Name, uint64_t Hash,
+ std::vector<uint64_t> Counts,
+ std::vector<uint8_t> BitmapBytes,
+ std::vector<uint8_t> UniformityBits)
+ : InstrProfRecord(std::move(Counts), std::move(BitmapBytes)), Name(Name),
+ Hash(Hash) {
+ this->UniformityBits = std::move(UniformityBits);
+ }
static bool hasCSFlagInHash(uint64_t FuncHash) {
return ((FuncHash >> CS_FLAG_IN_FUNC_HASH) & 1);
@@ -1185,7 +1222,9 @@ enum ProfVersion {
Version12 = 12,
// In this version, the frontend PGO stable hash algorithm defaults to V4.
Version13 = 13,
- // The current version is 13.
+ // UniformityBits added for AMDGPU offload profiling divergence detection.
+ Version14 = 14,
+ // The current version is 14.
CurrentVersion = INSTR_PROF_INDEX_VERSION
};
const uint64_t Version = ProfVersion::CurrentVersion;
diff --git a/llvm/include/llvm/ProfileData/InstrProfData.inc b/llvm/include/llvm/ProfileData/InstrProfData.inc
index 3117cb2fddf36..1f45eb0bc8b2c 100644
--- a/llvm/include/llvm/ProfileData/InstrProfData.inc
+++ b/llvm/include/llvm/ProfileData/InstrProfData.inc
@@ -78,6 +78,8 @@ INSTR_PROF_DATA(const uint64_t, llvm::Type::getInt64Ty(Ctx), FuncHash, \
ConstantInt::get(llvm::Type::getInt64Ty(Ctx), \
Inc->getHash()->getZExtValue()))
INSTR_PROF_DATA(const IntPtrT, IntPtrTy, CounterPtr, RelativeCounterPtr)
+INSTR_PROF_DATA(const IntPtrT, IntPtrTy, UniformCounterPtr, \
+ RelativeUniformCounterPtr)
INSTR_PROF_DATA(const IntPtrT, IntPtrTy, BitmapPtr, RelativeBitmapPtr)
/* This is used to map function pointers for the indirect call targets to
* function name hashes during the conversion from raw to merged profile
@@ -91,6 +93,10 @@ INSTR_PROF_DATA(const uint32_t, llvm::Type::getInt32Ty(Ctx), NumCounters, \
ConstantInt::get(llvm::Type::getInt32Ty(Ctx), NumCounters))
INSTR_PROF_DATA(const uint16_t, Int16ArrayTy, NumValueSites[IPVK_Last+1], \
ConstantArray::get(Int16ArrayTy, Int16ArrayVals)) \
+INSTR_PROF_DATA(const uint16_t, llvm::Type::getInt16Ty(Ctx), \
+ OffloadDeviceWaveSize, \
+ ConstantInt::get(llvm::Type::getInt16Ty(Ctx), \
+ OffloadDeviceWaveSizeVal))
INSTR_PROF_DATA(const uint32_t, llvm::Type::getInt32Ty(Ctx), NumBitmapBytes, \
ConstantInt::get(llvm::Type::getInt32Ty(Ctx), NumBitmapBytes))
#undef INSTR_PROF_DATA
@@ -168,6 +174,12 @@ INSTR_PROF_GPU_SECT(const __llvm_profile_data *, llvm::PointerType::getUnqual( \
INSTR_PROF_GPU_SECT(const __llvm_profile_data *, llvm::PointerType::getUnqual( \
Ctx), DataStop, \
ConstantPointerNull::get(llvm::PointerType::getUnqual(Ctx)))
+INSTR_PROF_GPU_SECT(char *, llvm::PointerType::getUnqual(Ctx), \
+ UniformCountersStart, \
+ ConstantPointerNull::get(llvm::PointerType::getUnqual(Ctx)))
+INSTR_PROF_GPU_SECT(char *, llvm::PointerType::getUnqual(Ctx), \
+ UniformCountersStop, \
+ ConstantPointerNull::get(llvm::PointerType::getUnqual(Ctx)))
INSTR_PROF_GPU_SECT(uint64_t *, llvm::PointerType::getUnqual(Ctx), \
VersionVar, \
ConstantPointerNull::get(llvm::PointerType::getUnqual(Ctx)))
@@ -192,6 +204,12 @@ INSTR_PROF_RAW_HEADER(uint64_t, NumCounters, NumCounters)
INSTR_PROF_RAW_HEADER(uint64_t, PaddingBytesAfterCounters, PaddingBytesAfterCounters)
INSTR_PROF_RAW_HEADER(uint64_t, NumBitmapBytes, NumBitmapBytes)
INSTR_PROF_RAW_HEADER(uint64_t, PaddingBytesAfterBitmapBytes, PaddingBytesAfterBitmapBytes)
+INSTR_PROF_RAW_HEADER(uint64_t, NumUniformCounters, NumUniformCounters)
+INSTR_PROF_RAW_HEADER(uint64_t, PaddingBytesAfterUniformCounters, PaddingBytesAfterUniformCounters)
+INSTR_PROF_RAW_HEADER(uint64_t, UniformCountersDelta,
+ UniformCountersBegin ? (uintptr_t)UniformCountersBegin -
+ (uintptr_t)DataBegin
+ : 0)
INSTR_PROF_RAW_HEADER(uint64_t, NamesSize, NamesSize)
INSTR_PROF_RAW_HEADER(uint64_t, CountersDelta,
(uintptr_t)CountersBegin - (uintptr_t)DataBegin)
@@ -356,6 +374,9 @@ INSTR_PROF_SECT_ENTRY(IPSK_data, \
INSTR_PROF_SECT_ENTRY(IPSK_cnts, \
INSTR_PROF_QUOTE(INSTR_PROF_CNTS_COMMON), \
INSTR_PROF_CNTS_COFF, "__DATA,")
+INSTR_PROF_SECT_ENTRY(IPSK_ucnts, \
+ INSTR_PROF_QUOTE(INSTR_PROF_UCNTS_COMMON), \
+ INSTR_PROF_UCNTS_COFF, "__DATA,")
INSTR_PROF_SECT_ENTRY(IPSK_bitmap, \
INSTR_PROF_QUOTE(INSTR_PROF_BITS_COMMON), \
INSTR_PROF_BITS_COFF, "__DATA,")
@@ -752,9 +773,9 @@ serializeValueProfDataFrom(ValueProfRecordClosure *Closure,
(uint64_t)'f' << 16 | (uint64_t)'R' << 8 | (uint64_t)129
/* Raw profile format version (start from 1). */
-#define INSTR_PROF_RAW_VERSION 10
+#define INSTR_PROF_RAW_VERSION 11
/* Indexed profile format version (start from 1). */
-#define INSTR_PROF_INDEX_VERSION 13
+#define INSTR_PROF_INDEX_VERSION 14
/* Coverage mapping format version (start from 0). */
#define INSTR_PROF_COVMAP_VERSION 6
@@ -803,6 +824,7 @@ serializeValueProfDataFrom(ValueProfRecordClosure *Closure,
#define INSTR_PROF_NAME_COMMON __llvm_prf_names
#define INSTR_PROF_VNAME_COMMON __llvm_prf_vns
#define INSTR_PROF_CNTS_COMMON __llvm_prf_cnts
+#define INSTR_PROF_UCNTS_COMMON __llvm_prf_ucnts
#define INSTR_PROF_BITS_COMMON __llvm_prf_bits
#define INSTR_PROF_VALS_COMMON __llvm_prf_vals
#define INSTR_PROF_VNODES_COMMON __llvm_prf_vnds
@@ -820,6 +842,7 @@ serializeValueProfDataFrom(ValueProfRecordClosure *Closure,
#define INSTR_PROF_NAME_COFF ".lprfn$M"
#define INSTR_PROF_VNAME_COFF ".lprfvn$M"
#define INSTR_PROF_CNTS_COFF ".lprfc$M"
+#define INSTR_PROF_UCNTS_COFF ".lprfuc$M"
#define INSTR_PROF_BITS_COFF ".lprfb$M"
#define INSTR_PROF_VALS_COFF ".lprfv$M"
#define INSTR_PROF_VNODES_COFF ".lprfnd$M"
diff --git a/llvm/include/llvm/ProfileData/InstrProfReader.h b/llvm/include/llvm/ProfileData/InstrProfReader.h
index 134195059f9e6..4d6351f6425ef 100644
--- a/llvm/include/llvm/ProfileData/InstrProfReader.h
+++ b/llvm/include/llvm/ProfileData/InstrProfReader.h
@@ -349,6 +349,7 @@ class RawInstrProfReader : public InstrProfReader {
uint64_t Version;
uint64_t CountersDelta;
uint64_t BitmapDelta;
+ uint64_t UniformCountersDelta;
uint64_t NamesDelta;
const RawInstrProf::ProfileData<IntPtrT> *Data;
const RawInstrProf::ProfileData<IntPtrT> *DataEnd;
@@ -358,6 +359,8 @@ class RawInstrProfReader : public InstrProfReader {
const char *CountersEnd;
const char *BitmapStart;
const char *BitmapEnd;
+ const char *UniformCountersStart;
+ const char *UniformCountersEnd;
const char *NamesStart;
const char *NamesEnd;
const char *VNamesStart = nullptr;
@@ -469,6 +472,7 @@ class RawInstrProfReader : public InstrProfReader {
Error readFuncHash(NamedInstrProfRecord &Record);
Error readRawCounts(InstrProfRecord &Record);
Error readRawBitmapBytes(InstrProfRecord &Record);
+ Error readRawUniformCounters(InstrProfRecord &Record);
Error readValueProfilingData(InstrProfRecord &Record);
bool atEnd() const { return Data == DataEnd; }
@@ -482,6 +486,7 @@ class RawInstrProfReader : public InstrProfReader {
// with respect to the next record.
CountersDelta -= sizeof(*Data);
BitmapDelta -= sizeof(*Data);
+ UniformCountersDelta -= sizeof(*Data);
}
Data++;
ValueDataStart += CurValueDataSize;
diff --git a/llvm/include/llvm/ProfileData/InstrProfWriter.h b/llvm/include/llvm/ProfileData/InstrProfWriter.h
index 1b24425e68a9e..93212c19769e0 100644
--- a/llvm/include/llvm/ProfileData/InstrProfWriter.h
+++ b/llvm/include/llvm/ProfileData/InstrProfWriter.h
@@ -215,6 +215,7 @@ class InstrProfWriter {
MemProfVersionRequested = Version;
}
void setMemProfFullSchema(bool Full) { MemProfFullSchema = Full; }
+
// Compute the overlap b/w this object and Other. Program level result is
// stored in Overlap and function level result is stored in FuncLevelOverlap.
LLVM_ABI void overlapRecord(NamedInstrProfRecord &&Other,
diff --git a/llvm/lib/ProfileData/InstrProf.cpp b/llvm/lib/ProfileData/InstrProf.cpp
index d9ae1a6437bb1..a9ef7f64da95f 100644
--- a/llvm/lib/ProfileData/InstrProf.cpp
+++ b/llvm/lib/ProfileData/InstrProf.cpp
@@ -950,7 +950,8 @@ void InstrProfRecord::mergeValueProfData(
}
void InstrProfRecord::merge(InstrProfRecord &Other, uint64_t Weight,
- function_ref<void(instrprof_error)> Warn) {
+ function_ref<void(instrprof_error)> Warn,
+ unsigned WaveSize) {
// If the number of counters doesn't match we either have bad data
// or a hash collision.
if (Counts.size() != Other.Counts.size()) {
@@ -975,7 +976,7 @@ void InstrProfRecord::merge(InstrProfRecord &Other, uint64_t Weight,
setPseudoCount(PseudoWarm);
return;
}
-
+ OffloadDeviceWaveSize = Other.OffloadDeviceWaveSize;
for (size_t I = 0, E = Other.Counts.size(); I < E; ++I) {
bool Overflowed;
uint64_t Value =
@@ -1715,7 +1716,7 @@ Expected<Header> Header::readFromBuffer(const unsigned char *Buffer) {
IndexedInstrProf::ProfVersion::CurrentVersion)
return make_error<InstrProfError>(instrprof_error::unsupported_version);
- static_assert(IndexedInstrProf::ProfVersion::CurrentVersion == Version13,
+ static_assert(IndexedInstrProf::ProfVersion::CurrentVersion == Version14,
"Please update the reader as needed when a new field is added "
"or when indexed profile version gets bumped.");
@@ -1748,10 +1749,11 @@ size_t Header::size() const {
// of the header, and byte offset of existing fields shouldn't change when
// indexed profile version gets incremented.
static_assert(
- IndexedInstrProf::ProfVersion::CurrentVersion == Version13,
+ IndexedInstrProf::ProfVersion::CurrentVersion == Version14,
"Please update the size computation below if a new field has "
"been added to the header; for a version bump without new "
"fields, add a case statement to fall through to the latest version.");
+ case 14ull: // UniformityBits added in record data, no header change
case 13ull:
case 12ull:
return 72;
diff --git a/llvm/lib/ProfileData/InstrProfCorrelator.cpp b/llvm/lib/ProfileData/InstrProfCorrelator.cpp
index c571c338c8d32..c732382eb399a 100644
--- a/llvm/lib/ProfileData/InstrProfCorrelator.cpp
+++ b/llvm/lib/ProfileData/InstrProfCorrelator.cpp
@@ -311,6 +311,7 @@ void InstrProfCorrelatorImpl<IntPtrT>::addDataProbe(uint64_t NameRef,
// In this mode, CounterPtr actually stores the section relative address
// of the counter.
maybeSwap<IntPtrT>(CounterOffset),
+ /*UniformCounterPtr=*/maybeSwap<IntPtrT>(0),
// TODO: MC/DC is not yet supported.
/*BitmapOffset=*/maybeSwap<IntPtrT>(0),
maybeSwap<IntPtrT>(FunctionPtr),
@@ -318,6 +319,7 @@ void InstrProfCorrelatorImpl<IntPtrT>::addDataProbe(uint64_t NameRef,
/*ValuesPtr=*/maybeSwap<IntPtrT>(0),
maybeSwap<uint32_t>(NumCounters),
/*NumValueSites=*/{maybeSwap<uint16_t>(0), maybeSwap<uint16_t>(0)},
+ /*OffloadDeviceWaveSize=*/maybeSwap<uint16_t>(0),
// TODO: MC/DC is not yet supported.
/*NumBitmapBytes=*/maybeSwap<uint32_t>(0),
});
diff --git a/llvm/lib/ProfileData/InstrProfReader.cpp b/llvm/lib/ProfileData/InstrProfReader.cpp
index 8147ee8d0e816..b0bf80c0a7a6e 100644
--- a/llvm/lib/ProfileData/InstrProfReader.cpp
+++ b/llvm/lib/ProfileData/InstrProfReader.cpp
@@ -607,6 +607,7 @@ Error RawInstrProfReader<IntPtrT>::readHeader(
CountersDelta = swap(Header.CountersDelta);
BitmapDelta = swap(Header.BitmapDelta);
+ UniformCountersDelta = swap(Header.UniformCountersDelta);
NamesDelta = swap(Header.NamesDelta);
auto NumData = swap(Header.NumData);
auto PaddingBytesBeforeCounters = swap(Header.PaddingBytesBeforeCounters);
@@ -614,6 +615,9 @@ Error RawInstrProfReader<IntPtrT>::readHeader(
auto PaddingBytesAfterCounters = swap(Header.PaddingBytesAfterCounters);
auto NumBitmapBytes = swap(Header.NumBitmapBytes);
auto PaddingBytesAfterBitmapBytes = swap(Header.PaddingBytesAfterBitmapBytes);
+ auto NumUniformCounters = swap(Header.NumUniformCounters);
+ auto PaddingBytesAfterUniformCounters =
+ swap(Header.PaddingBytesAfterUniformCounters);
auto NamesSize = swap(Header.NamesSize);
auto VTableNameSize = swap(Header.VNamesSize);
auto NumVTables = swap(Header.NumVTables);
@@ -626,14 +630,17 @@ Error RawInstrProfReader<IntPtrT>::readHeader(
auto VTableSectionSize =
NumVTables * sizeof(RawInstrProf::VTableProfileData<IntPtrT>);
auto PaddingBytesAfterVTableProfData = getNumPaddingBytes(VTableSectionSize);
+ auto UniformCountersSectionSize = NumUniformCounters * sizeof(uint64_t);
// Profile data starts after profile header and binary ids if exist.
ptrdiff_t DataOffset = sizeof(RawInstrProf::Header) + BinaryIdSize;
ptrdiff_t CountersOffset = DataOffset + DataSize + PaddingBytesBeforeCounters;
ptrdiff_t BitmapOffset =
CountersOffset + CountersSize + PaddingBytesAfterCounters;
- ptrdiff_t NamesOffset =
+ ptrdiff_t UniformCountersOffset =
BitmapOffset + NumBitmapBytes + PaddingBytesAfterBitmapBytes;
+ ptrdiff_t NamesOffset = UniformCountersOffset + UniformCountersSectionSize +
+ PaddingBytesAfterUniformCounters;
ptrdiff_t VTableProfDataOffset =
NamesOffset + NamesSize + PaddingBytesAfterNames;
ptrdiff_t VTableNameOffset = VTableProfDataOffset + VTableSectionSize +
@@ -694,6 +701,8 @@ Error RawInstrProfReader<IntPtrT>::readHeader(
CountersEnd = CountersStart + CountersSize;
BitmapStart = Start + BitmapOffset;
BitmapEnd = BitmapStart + NumBitmapBytes;
+ UniformCountersStart = Start + UniformCountersOffset;
+ UniformCountersEnd = UniformCountersStart + UniformCountersSectionSize;
ValueDataStart = reinterpret_cast<const uint8_t *>(Start + ValueDataOffset);
std::unique_ptr<InstrProfSymtab> NewSymtab = std::make_unique<InstrProfSymtab>();
@@ -827,6 +836,52 @@ Error RawInstrProfReader<IntPtrT>::readRawBitmapBytes(InstrProfRecord &Record) {
return success();
}
+template <class IntPtrT>
+Error RawInstrProfReader<IntPtrT>::readRawUniformCounters(
+ InstrProfRecord &Record) {
+ Record.UniformCounts.clear();
+
+ if (UniformCountersStart == UniformCountersEnd)
+ return success();
+
+ uint32_t NumCounters = swap(Data->NumCounters);
+
+ ptrdiff_t UniformCounterOffset =
+ swap(Data->UniformCounterPtr) - UniformCountersDelta;
+ if (UniformCounterOffset < 0)
+ return error(instrprof_error::malformed,
+ ("uniform counter offset " + Twine(UniformCounterOffset) +
+ " is negative")
+ .str());
+
+ if (UniformCounterOffset >= UniformCountersEnd - UniformCountersStart)
+ return error(instrprof_error::malformed,
+ ("uniform counter offset " + Twine(UniformCounterOffset) +
+ " is greater than the maximum uniform counter offset " +
+ Twine(UniformCountersEnd - UniformCountersStart - 1))
+ .str());
+
+ uint64_t MaxNumCounters =
+ (UniformCountersEnd - (UniformCountersStart + UniformCounterOffset)) /
+ sizeof(uint64_t);
+ if (NumCounters > MaxNumCounters)
+ return error(instrprof_error::malformed,
+ ("number of uniform counters " + Twine(NumCounters) +
+ " is greater than the maximum number of uniform counters " +
+ Twine(MaxNumCounters))
+ .str());
+
+ Record.UniformCounts.reserve(NumCounters);
+ for (uint32_t I = 0; I < NumCounters; I++) {
+ const char *Ptr =
+ UniformCountersStart + UniformCounterOffset + I * sizeof(uint64_t);
+ uint64_t CounterValue = swap(*reinterpret_cast<const uint64_t *>(Ptr));
+ Record.UniformCounts.push_back(CounterValue);
+ }
+
+ return success();
+}
+
template <class IntPtrT>
Error RawInstrProfReader<IntPtrT>::readValueProfilingData(
InstrProfRecord &Record) {
@@ -873,6 +928,8 @@ Error RawInstrProfReader<IntPtrT>::readNextRecord(NamedInstrProfRecord &Record)
if (Error E = readFuncHash(Record))
return error(std::move(E));
+ Record.OffloadDeviceWaveSize = swap(Data->OffloadDeviceWaveSize);
+
// Read raw counts and set Record.
if (Error E = readRawCounts(Record))
return error(std::move(E));
@@ -881,6 +938,10 @@ Error RawInstrProfReader<IntPtrT>::readNextRecord(NamedInstrProfRecord &Record)
if (Error E = readRawBitmapBytes(Record))
return error(std::move(E));
+ // Read raw uniform counters and set Record.
+ if (Error E = readRawUniformCounters(Record))
+ return error(std::move(E));
+
// Read value data and set Record.
if (Error E = readValueProfilingData(Record))
return error(std::move(E));
@@ -945,11 +1006,12 @@ data_type InstrProfLookupTrait::ReadData(StringRef K, const unsigned char *D,
DataBuffer.clear();
std::vector<uint64_t> CounterBuffer;
std::vector<uint8_t> BitmapByteBuffer;
+ std::vector<uint8_t> UniformityBitsBuffer;
const unsigned char *End = D + N;
while (D < End) {
// Read hash.
- if (D + sizeof(uint64_t) >= End)
+ if (D + sizeof(uint64_t) > End)
return data_type();
uint64_t Hash = endian::readNext<uint64_t, llvm::endianness::little>(D);
@@ -977,18 +1039,51 @@ data_type InstrProfLookupTrait::ReadData(StringRef K, const unsigned char *D,
if (D + sizeof(uint64_t) > End)
return data_type();
BitmapBytes = endian::readNext<uint64_t, llvm::endianness::little>(D);
- // Read bitmap byte values.
- if (D + BitmapBytes * sizeof(uint8_t) > End)
- return data_type();
BitmapByteBuffer.clear();
BitmapByteBuffer.reserve(BitmapBytes);
- for (uint64_t J = 0; J < BitmapBytes; ++J)
- BitmapByteBuffer.push_back(static_cast<uint8_t>(
- endian::readNext<uint64_t, llvm::endianness::little>(D)));
+
+ if (GET_VERSION(FormatVersion) >=
+ IndexedInstrProf::ProfVersion::Version14) {
+ // Version 14+: bitmap bytes stored as uint8_t with padding.
+ uint64_t PaddedSize = alignTo(BitmapBytes, sizeof(uint64_t));
+ if (D + PaddedSize > End)
+ return data_type();
+ for (uint64_t J = 0; J < BitmapBytes; ++J)
+ BitmapByteBuffer.push_back(
+ endian::readNext<uint8_t, llvm::endianness::little>(D));
+ for (uint64_t J = BitmapBytes; J < PaddedSize; ++J)
+ (void)endian::readNext<uint8_t, llvm::endianness::little>(D);
+
+ // Read uniformity bits (AMDGPU offload profiling).
+ uint64_t UniformityBitsSize = 0;
+ if (D + sizeof(uint64_t) > End)
+ return data_type();
+ UniformityBitsSize =
+ endian::readNext<uint64_t, llvm::endianness::little>(D);
+ uint64_t PaddedUniformitySize =
+ alignTo(UniformityBitsSize, sizeof(uint64_t));
+ if (D + PaddedUniformitySize > End)
+ return data_type();
+ UniformityBitsBuffer.clear();
+ UniformityBitsBuffer.reserve(UniformityBitsSize);
+ for (uint64_t J = 0; J < UniformityBitsSize; ++J)
+ UniformityBitsBuffer.push_back(
+ endian::readNext<uint8_t, llvm::endianness::little>(D));
+ for (uint64_t J = UniformityBitsSize; J < PaddedUniformitySize; ++J)
+ (void)endian::readNext<uint8_t, llvm::endianness::little>(D);
+ } else {
+ // Version 11-13: each bitmap byte stored as a uint64_t.
+ if (D + BitmapBytes * sizeof(uint64_t) > End)
+ return data_type();
+ for (uint64_t J = 0; J < BitmapBytes; ++J)
+ BitmapByteBuffer.push_back(static_cast<uint8_t>(
+ endian::readNext<uint64_t, llvm::endianness::little>(D)));
+ }
}
DataBuffer.emplace_back(K, Hash, std::move(CounterBuffer),
- std::move(BitmapByteBuffer));
+ std::move(BitmapByteBuffer),
+ std::move(UniformityBitsBuffer));
// Read value profiling data.
if (GET_VERSION(FormatVersion) > IndexedInstrProf::ProfVersion::Version2 &&
diff --git a/llvm/lib/ProfileData/InstrProfWriter.cpp b/llvm/lib/ProfileData/InstrProfWriter.cpp
index 0f15ca8ff6df7..96611612184e7 100644
--- a/llvm/lib/ProfileData/InstrProfWriter.cpp
+++ b/llvm/lib/ProfileData/InstrProfWriter.cpp
@@ -51,6 +51,7 @@ class InstrProfRecordWriterTrait {
llvm::endianness ValueProfDataEndianness = llvm::endianness::little;
InstrProfSummaryBuilder *SummaryBuilder;
InstrProfSummaryBuilder *CSSummaryBuilder;
+ bool WritePrevVersion = false;
InstrProfRecordWriterTrait() = default;
@@ -58,7 +59,7 @@ class InstrProfRecordWriterTrait {
return IndexedInstrProf::ComputeHash(K);
}
- static std::pair<offset_type, offset_type>
+ std::pair<offset_type, offset_type>
EmitKeyDataLength(raw_ostream &Out, key_type_ref K, data_type_ref V) {
using namespace support;
@@ -74,7 +75,16 @@ class InstrProfRecordWriterTrait {
M += sizeof(uint64_t); // The size of the Counts vector
M += ProfRecord.Counts.size() * sizeof(uint64_t);
M += sizeof(uint64_t); // The size of the Bitmap vector
- M += ProfRecord.BitmapBytes.size() * sizeof(uint64_t);
+ if (WritePrevVersion) {
+ // Version 13: each bitmap byte stored as a uint64_t.
+ M += ProfRecord.BitmapBytes.size() * sizeof(uint64_t);
+ } else {
+ // Version 14+: bitmap bytes as uint8_t with padding, plus
+ // uniformity bits.
+ M += alignTo(ProfRecord.BitmapBytes.size(), sizeof(uint64_t));
+ M += sizeof(uint64_t); // The size of the UniformityBits vector
+ M += alignTo(ProfRecord.UniformityBits.size(), sizeof(uint64_t));
+ }
// Value data
M += ValueProfData::getSize(ProfileData.second);
@@ -88,7 +98,8 @@ class InstrProfRecordWriterTrait {
Out.write(K.data(), N);
}
- void EmitData(raw_ostream &Out, key_type_ref, data_type_ref V, offset_type) {
+ void EmitData(raw_ostream &Out, key_type_ref K, data_type_ref V,
+ offset_type) {
using namespace support;
endian::Writer LE(Out, llvm::endianness::little);
@@ -105,8 +116,27 @@ class InstrProfRecordWriterTrait {
LE.write<uint64_t>(I);
LE.write<uint64_t>(ProfRecord.BitmapBytes.size());
- for (uint64_t I : ProfRecord.BitmapBytes)
- LE.write<uint64_t>(I);
+ if (WritePrevVersion) {
+ // Version 13: each bitmap byte stored as a uint64_t.
+ for (uint8_t I : ProfRecord.BitmapBytes)
+ LE.write<uint64_t>(I);
+ } else {
+ // Version 14+: bitmap bytes as uint8_t with padding.
+ for (uint8_t I : ProfRecord.BitmapBytes)
+ LE.write<uint8_t>(I);
+ for (size_t I = ProfRecord.BitmapBytes.size();
+ I < alignTo(ProfRecord.BitmapBytes.size(), sizeof(uint64_t)); ++I)
+ LE.write<uint8_t>(0);
+
+ // Write uniformity bits (AMDGPU offload profiling).
+ LE.write<uint64_t>(ProfRecord.UniformityBits.size());
+ for (uint8_t I : ProfRecord.UniformityBits)
+ LE.write<uint8_t>(I);
+ for (size_t I = ProfRecord.UniformityBits.size();
+ I < alignTo(ProfRecord.UniformityBits.size(), sizeof(uint64_t));
+ ++I)
+ LE.write<uint8_t>(0);
+ }
// Write value data
std::unique_ptr<ValueProfData> VDataPtr =
@@ -209,7 +239,7 @@ void InstrProfWriter::addRecord(StringRef Name, uint64_t Hash,
Dest.scale(Weight, 1, MapWarn);
} else {
// We're updating a function we've seen before.
- Dest.merge(I, Weight, MapWarn);
+ Dest.merge(I, Weight, MapWarn, I.OffloadDeviceWaveSize);
}
Dest.sortValueData();
@@ -524,6 +554,7 @@ Error InstrProfWriter::writeImpl(ProfOStream &OS) {
InfoObj->SummaryBuilder = &ISB;
InstrProfSummaryBuilder CSISB(ProfileSummaryBuilder::DefaultCutoffs);
InfoObj->CSSummaryBuilder = &CSISB;
+ InfoObj->WritePrevVersion = WritePrevVersion;
// Populate the hash table generator.
SmallVector<std::pair<StringRef, const ProfilingData *>> OrderedData;
@@ -542,7 +573,7 @@ Error InstrProfWriter::writeImpl(ProfOStream &OS) {
// The WritePrevVersion handling will either need to be removed or updated
// if the version is advanced beyond 12.
static_assert(IndexedInstrProf::ProfVersion::CurrentVersion ==
- IndexedInstrProf::ProfVersion::Version13);
+ IndexedInstrProf::ProfVersion::Version14);
if (static_cast<bool>(ProfileKind & InstrProfKind::IRInstrumentation))
Header.Version |= VARIANT_MASK_IR_PROF;
if (static_cast<bool>(ProfileKind & InstrProfKind::ContextSensitive))
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index 8e4ba41919768..58aac27f95225 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -164,6 +164,14 @@ cl::opt<bool> SpeculativeCounterPromotionToLoop(
" update can be further/iteratively promoted into an acyclic "
" region."));
+static cl::opt<unsigned> OffloadPGOSampling(
+ "offload-pgo-sampling",
+ cl::desc("Log2 of the sampling period for offload PGO instrumentation. "
+ "Only 1 in every 2^N blocks is instrumented. "
+ "0 = all blocks, 1 = 50%, 2 = 25%, 3 = 12.5% (default). "
+ "Higher values reduce overhead at the cost of sparser profiles."),
+ cl::init(3));
+
cl::opt<bool> IterativeCounterPromotion(
"iterative-counter-promotion", cl::init(true),
cl::desc("Allow counter promotion across the whole loop nest."));
@@ -269,6 +277,8 @@ class InstrLowerer final {
struct PerFunctionProfileData {
uint32_t NumValueSites[IPVK_Last + 1] = {};
GlobalVariable *RegionCounters = nullptr;
+ GlobalVariable *UniformCounters =
+ nullptr; // Per-block uniform-entry counters
GlobalVariable *DataVar = nullptr;
GlobalVariable *RegionBitmaps = nullptr;
uint32_t NumBitmapBytes = 0;
@@ -298,6 +308,18 @@ class InstrLowerer final {
int64_t TotalCountersPromoted = 0;
+ // Per-function cache of invariant values for GPU PGO instrumentation.
+ // Computed once at the function entry and reused across all instrumentation
+ // points to avoid redundant IR and help the optimizer.
+ struct GPUPGOInvariants {
+ Value *Matched = nullptr;
+ bool WaveSizeStored = false;
+ };
+ DenseMap<Function *, GPUPGOInvariants> GPUInvariantsCache;
+
+ /// Emit invariant PGO values at the function entry block and cache them.
+ GPUPGOInvariants &getOrCreateGPUInvariants(Function *F);
+
/// Lower instrumentation intrinsics in the function. Returns true if there
/// any lowering.
bool lowerIntrinsics(Function *F);
@@ -354,6 +376,10 @@ class InstrLowerer final {
/// referring to them will also be created.
GlobalVariable *getOrCreateRegionCounters(InstrProfCntrInstBase *Inc);
+ /// Get the uniform entry counters for GPU divergence tracking.
+ /// These counters track how often blocks are entered with all lanes active.
+ GlobalVariable *getOrCreateUniformCounters(InstrProfCntrInstBase *Inc);
+
/// Create the region counters.
GlobalVariable *createRegionCounters(InstrProfCntrInstBase *Inc,
StringRef Name,
@@ -1197,26 +1223,125 @@ void InstrLowerer::lowerTimestamp(
TimestampInstruction->eraseFromParent();
}
+InstrLowerer::GPUPGOInvariants &
+InstrLowerer::getOrCreateGPUInvariants(Function *F) {
+ auto It = GPUInvariantsCache.find(F);
+ if (It != GPUInvariantsCache.end())
+ return It->second;
+
+ LLVMContext &Context = M.getContext();
+ auto *Int32Ty = Type::getInt32Ty(Context);
+
+ BasicBlock &EntryBB = F->getEntryBlock();
+ IRBuilder<> Builder(&*EntryBB.getFirstInsertionPt());
+
+ Value *Matched = ConstantInt::getTrue(Context);
+ if (OffloadPGOSampling > 0) {
+ FunctionCallee IsSampledFn =
+ M.getOrInsertFunction(RTLIB::RuntimeLibcallsInfo::getLibcallImplName(
+ RTLIB::impl___llvm_profile_sampling_gpu),
+ Int32Ty, Int32Ty);
+ Value *SampledInt = Builder.CreateCall(
+ IsSampledFn, {ConstantInt::get(Int32Ty, OffloadPGOSampling)},
+ "pgo.sampled");
+ Matched = Builder.CreateICmpNE(SampledInt, ConstantInt::get(Int32Ty, 0),
+ "pgo.matched");
+ }
+
+ auto &Inv = GPUInvariantsCache[F];
+ Inv.Matched = Matched;
+ return Inv;
+}
+
void InstrLowerer::lowerIncrement(InstrProfIncrementInst *Inc) {
- auto *Addr = getCounterAddress(Inc);
IRBuilder<> Builder(Inc);
if (isGPUProfTarget(M)) {
- auto *Int64Ty = Builder.getInt64Ty();
- auto *PtrTy = Builder.getPtrTy();
- auto *CalleeTy = FunctionType::get(Type::getVoidTy(M.getContext()),
+ Function *F = Inc->getFunction();
+ auto &Inv = getOrCreateGPUInvariants(F);
+
+ LLVMContext &Context = M.getContext();
+ auto *Int64Ty = Type::getInt64Ty(Context);
+ auto *PtrTy = PointerType::getUnqual(Context);
+
+ auto *Addr = getCounterAddress(Inc);
+
+ // Store the device wave/warp size into the profile data struct once per
+ // function. The size lets the host-side merge classify wave-uniform
+ // counters. AMDGPU folds llvm.amdgcn.wavefrontsize to the subtarget's
+ // constant; other GPUs use their fixed warp size.
+ if (!Inv.WaveSizeStored) {
+ Inv.WaveSizeStored = true;
+ GlobalVariable *NamePtr = Inc->getName();
+ auto &PD = ProfileDataMap[NamePtr];
+ if (PD.DataVar) {
+ IRBuilder<> EntryBuilder(&*F->getEntryBlock().getFirstInsertionPt());
+ Value *WaveSize16 = nullptr;
+ // Look the intrinsic up by name so this target-agnostic pass does not
+ // pull in IntrinsicsAMDGPU.h. AMDGPU folds the intrinsic to the
+ // subtarget's wavefront size; other GPUs fall back to a 32-lane warp.
+ if (TT.isAMDGPU()) {
+ Intrinsic::ID WaveSizeID =
+ Intrinsic::lookupIntrinsicID("llvm.amdgcn.wavefrontsize");
+ if (WaveSizeID != Intrinsic::not_intrinsic) {
+ Function *WaveSizeFn =
+ Intrinsic::getOrInsertDeclaration(&M, WaveSizeID);
+ Value *WaveSize = EntryBuilder.CreateCall(WaveSizeFn);
+ WaveSize16 = EntryBuilder.CreateTrunc(
+ WaveSize, Type::getInt16Ty(Context), "wavesize.i16");
+ }
+ }
+ if (!WaveSize16)
+ WaveSize16 = ConstantInt::get(Type::getInt16Ty(Context), 32);
+ Value *WaveSizeAddr = EntryBuilder.CreateStructGEP(
+ PD.DataVar->getValueType(), PD.DataVar, 9, "profd.wavesize");
+ EntryBuilder.CreateStore(WaveSize16, WaveSizeAddr);
+ }
+ }
+
+ GlobalVariable *UniformCounters = getOrCreateUniformCounters(Inc);
+ Value *UniformAddrArg = ConstantPointerNull::get(PtrTy);
+ if (UniformCounters) {
+ Value *UniformIndices[] = {Builder.getInt32(0), Inc->getIndex()};
+ Value *UniformAddr = Builder.CreateInBoundsGEP(
+ UniformCounters->getValueType(), UniformCounters, UniformIndices,
+ "unifctr.addr");
+ UniformAddrArg =
+ Builder.CreatePointerBitCastOrAddrSpaceCast(UniformAddr, PtrTy);
+ }
+ Value *CastAddr = Builder.CreatePointerBitCastOrAddrSpaceCast(Addr, PtrTy);
+ Value *StepI64 =
+ Builder.CreateZExtOrTrunc(Inc->getStep(), Int64Ty, "step.i64");
+
+ auto *CalleeTy = FunctionType::get(Type::getVoidTy(Context),
{PtrTy, PtrTy, Int64Ty}, false);
FunctionCallee Callee =
M.getOrInsertFunction(RTLIB::RuntimeLibcallsInfo::getLibcallImplName(
RTLIB::impl___llvm_profile_instrument_gpu),
CalleeTy);
- Value *CastAddr = Builder.CreatePointerBitCastOrAddrSpaceCast(Addr, PtrTy);
- Value *Uniform =
- ConstantPointerNull::get(PointerType::getUnqual(M.getContext()));
- Value *StepI64 =
- Builder.CreateZExtOrTrunc(Inc->getStep(), Int64Ty, "step.i64");
- Builder.CreateCall(Callee, {CastAddr, Uniform, StepI64});
- } else if (Options.Atomic || AtomicCounterUpdateAll ||
- (Inc->getIndex()->isNullValue() && AtomicFirstCounter)) {
+
+ if (OffloadPGOSampling > 0) {
+ BasicBlock *CurBB = Builder.GetInsertBlock();
+ BasicBlock *ContBB =
+ CurBB->splitBasicBlock(BasicBlock::iterator(Inc), "po_cont");
+ BasicBlock *ThenBB = BasicBlock::Create(Context, "po_then", F);
+
+ CurBB->getTerminator()->eraseFromParent();
+ IRBuilder<> HeadBuilder(CurBB);
+ HeadBuilder.CreateCondBr(Inv.Matched, ThenBB, ContBB);
+
+ IRBuilder<> ThenBuilder(ThenBB);
+ ThenBuilder.CreateCall(Callee, {CastAddr, UniformAddrArg, StepI64});
+ ThenBuilder.CreateBr(ContBB);
+ } else {
+ Builder.CreateCall(Callee, {CastAddr, UniformAddrArg, StepI64});
+ }
+ Inc->eraseFromParent();
+ return;
+ }
+
+ auto *Addr = getCounterAddress(Inc);
+ if (Options.Atomic || AtomicCounterUpdateAll ||
+ (Inc->getIndex()->isNullValue() && AtomicFirstCounter)) {
Builder.CreateAtomicRMW(AtomicRMWInst::Add, Addr, Inc->getStep(),
MaybeAlign(), AtomicOrdering::Monotonic);
} else {
@@ -1753,12 +1878,54 @@ InstrLowerer::getOrCreateRegionCounters(InstrProfCntrInstBase *Inc) {
CompilerUsedVars.push_back(PD.RegionCounters);
}
+ // Create uniform counters before the data variable so that
+ // UniformCounterPtr can reference them in createDataVariable().
+ getOrCreateUniformCounters(Inc);
+
// Create the data variable (if it doesn't already exist).
createDataVariable(Inc);
return PD.RegionCounters;
}
+GlobalVariable *
+InstrLowerer::getOrCreateUniformCounters(InstrProfCntrInstBase *Inc) {
+ // Uniform counters are only meaningful for GPU profile targets.
+ if (!isGPUProfTarget(M))
+ return nullptr;
+
+ GlobalVariable *NamePtr = Inc->getName();
+ auto &PD = ProfileDataMap[NamePtr];
+ if (PD.UniformCounters)
+ return PD.UniformCounters;
+
+ // Ensure RegionCounters exists first (we need the same size)
+ getOrCreateRegionCounters(Inc);
+
+ uint64_t NumCounters = Inc->getNumCounters()->getZExtValue();
+
+ LLVMContext &Ctx = M.getContext();
+ ArrayType *CounterTy = ArrayType::get(Type::getInt64Ty(Ctx), NumCounters);
+
+ bool Renamed;
+ std::string VarName = getVarName(Inc, "__llvm_prf_unifcnt_", Renamed);
+
+ auto *GV = new GlobalVariable(M, CounterTy, false, NamePtr->getLinkage(),
+ Constant::getNullValue(CounterTy), VarName);
+ GV->setAlignment(Align(8));
+
+ GV->setSection(getInstrProfSectionName(IPSK_ucnts, TT.getObjectFormat()));
+
+ GV->setComdat(M.getOrInsertComdat(VarName));
+ GV->setLinkage(GlobalValue::LinkOnceODRLinkage);
+ GV->setVisibility(GlobalValue::ProtectedVisibility);
+
+ PD.UniformCounters = GV;
+ CompilerUsedVars.push_back(GV);
+
+ return PD.UniformCounters;
+}
+
void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
// When debug information is correlated to profile data, a data variable
// is not needed.
@@ -1823,6 +1990,7 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
uint64_t NumCounters = Inc->getNumCounters()->getZExtValue();
Constant *CounterPtr = PD.RegionCounters;
+ Constant *UniformCounterPtr = PD.UniformCounters;
uint64_t NumBitmapBytes = PD.NumBitmapBytes;
@@ -1838,6 +2006,8 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
for (uint32_t Kind = IPVK_First; Kind <= IPVK_Last; ++Kind)
Int16ArrayVals[Kind] = ConstantInt::get(Int16Ty, PD.NumValueSites[Kind]);
+ uint16_t OffloadDeviceWaveSizeVal = 0;
+
if (isGPUProfTarget(M)) {
// For GPU targets, weak functions need weak linkage for their profile data
// aliases to allow linker deduplication across TUs
@@ -1875,6 +2045,7 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
new GlobalVariable(M, DataTy, false, Linkage, nullptr, DataVarName);
Constant *RelativeCounterPtr;
+ Constant *RelativeUniformCounterPtr = ConstantInt::get(IntPtrTy, 0);
GlobalVariable *BitmapPtr = PD.RegionBitmaps;
Constant *RelativeBitmapPtr = ConstantInt::get(IntPtrTy, 0);
InstrProfSectKind DataSectionKind;
@@ -1885,6 +2056,9 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
RelativeCounterPtr = ConstantExpr::getPtrToInt(CounterPtr, IntPtrTy);
if (BitmapPtr != nullptr)
RelativeBitmapPtr = ConstantExpr::getPtrToInt(BitmapPtr, IntPtrTy);
+ if (UniformCounterPtr != nullptr)
+ RelativeUniformCounterPtr =
+ ConstantExpr::getPtrToInt(UniformCounterPtr, IntPtrTy);
} else if (TT.isNVPTX()) {
// The NVPTX target cannot handle self-referencing constant expressions in
// global initializers at all. Use absolute pointers and have the runtime
@@ -1902,6 +2076,10 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
RelativeBitmapPtr =
ConstantExpr::getSub(ConstantExpr::getPtrToInt(BitmapPtr, IntPtrTy),
ConstantExpr::getPtrToInt(Data, IntPtrTy));
+ if (UniformCounterPtr != nullptr)
+ RelativeUniformCounterPtr = ConstantExpr::getSub(
+ ConstantExpr::getPtrToInt(UniformCounterPtr, IntPtrTy),
+ ConstantExpr::getPtrToInt(Data, IntPtrTy));
}
Constant *DataVals[] = {
diff --git a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
index b6d07aa821e7f..ba0d5187f3c08 100644
--- a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
+++ b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
@@ -118,6 +118,7 @@
#include <algorithm>
#include <cassert>
#include <cstdint>
+#include <cstdlib>
#include <memory>
#include <numeric>
#include <optional>
@@ -379,9 +380,15 @@ class FunctionInstrumenter final {
// values. Supporting other values is relatively straight-forward - just
// another counter range within the context.
bool isValueProfilingDisabled() const {
+ // Value profiling is disabled for GPU targets because the device-side
+ // profiling runtime does not yet implement
+ // __llvm_profile_instrument_target. The existing compiler-rt implementation
+ // uses a linked-list with locks and eviction policy that is not efficient
+ // for massively parallel GPU execution. A GPU-optimized implementation is
+ // left as future work.
return DisableValueProfiling ||
InstrumentationType == PGOInstrumentationType::CTXPROF ||
- M.getTargetTriple().isGPU();
+ isGPUProfTarget(M);
}
bool shouldInstrumentEntryBB() const {
@@ -1305,11 +1312,25 @@ bool PGOUseFunc::setInstrumentedCounts(
setupBBInfoEdges(FuncInfo);
- unsigned NumCounters =
- InstrumentBBs.size() + FuncInfo.SIVisitor.getNumOfSelectInsts();
+ unsigned NumInstrumentedBBs = InstrumentBBs.size();
+ unsigned NumSelects = FuncInfo.SIVisitor.getNumOfSelectInsts();
+ unsigned NumCounters = NumInstrumentedBBs + NumSelects;
// The number of counters here should match the number of counters
// in profile. Return if they mismatch.
if (NumCounters != CountFromProfile.size()) {
+ LLVM_DEBUG({
+ dbgs() << "PGO COUNTER MISMATCH for function " << F.getName() << ":\n";
+ dbgs() << " Expected counters: " << NumCounters << "\n";
+ dbgs() << " - From instrumented edges: " << NumInstrumentedBBs << "\n";
+ for (size_t i = 0; i < InstrumentBBs.size(); ++i) {
+ dbgs() << " " << i << ": ";
+ InstrumentBBs[i]->printAsOperand(dbgs(), false);
+ dbgs() << "\n";
+ }
+ dbgs() << " - From select instructions: " << NumSelects << "\n";
+ dbgs() << " Actual counters from profile: " << CountFromProfile.size()
+ << "\n";
+ });
return false;
}
auto *FuncEntry = &*F.begin();
@@ -1317,6 +1338,7 @@ bool PGOUseFunc::setInstrumentedCounts(
// Set the profile count to the Instrumented BBs.
uint32_t I = 0;
for (BasicBlock *InstrBB : InstrumentBBs) {
+
uint64_t CountValue = CountFromProfile[I++];
PGOUseBBInfo &Info = getBBInfo(InstrBB);
// If we reach here, we know that we have some nonzero count
@@ -2410,11 +2432,12 @@ void llvm::setProfMetadata(Instruction *TI, ArrayRef<uint64_t> EdgeCounts,
LLVM_DEBUG(dbgs() << "Weight is: "; for (const auto &W
: Weights) {
dbgs() << W << " ";
- } dbgs() << "\n";);
+ } dbgs() << "\n");
misexpect::checkExpectAnnotations(*TI, Weights, /*IsFrontend=*/false);
setBranchWeights(*TI, Weights, /*IsExpected=*/false);
+
if (EmitBranchProbability) {
std::string BrCondStr = getBranchCondString(TI);
if (BrCondStr.empty())
diff --git a/llvm/test/CodeGen/AArch64/global-merge-profile-sections.ll b/llvm/test/CodeGen/AArch64/global-merge-profile-sections.ll
index 7108d82ddaa2e..37e5295c91aee 100644
--- a/llvm/test/CodeGen/AArch64/global-merge-profile-sections.ll
+++ b/llvm/test/CodeGen/AArch64/global-merge-profile-sections.ll
@@ -5,7 +5,7 @@ $__profc_begin = comdat nodeduplicate
$__profc_end = comdat nodeduplicate
@__profc_begin = private global [2 x i64] zeroinitializer, section "__llvm_prf_cnts", comdat, align 8
- at __profd_begin = private global { i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i32 } { i64 -1301828029439649651, i64 172590168, i64 sub (i64 ptrtoint (ptr @__profc_begin to i64), i64 ptrtoint (ptr @__profd_begin to i64)), i64 0, ptr null, ptr null, i32 2, [3 x i16] zeroinitializer, i32 0 }, section "__llvm_prf_data", comdat($__profc_begin), align 8
+ at __profd_begin = private global { i64, i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i16, i32 } { i64 -1301828029439649651, i64 172590168, i64 sub (i64 ptrtoint (ptr @__profc_begin to i64), i64 ptrtoint (ptr @__profd_begin to i64)), i64 0, i64 0, ptr null, ptr null, i32 2, [3 x i16] zeroinitializer, i16 0, i32 0 }, section "__llvm_prf_data", comdat($__profc_begin), align 8
@__profc_end = private global [2 x i64] zeroinitializer, section "__llvm_prf_cnts", comdat, align 8
- at __profd_end = private global { i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i32 } { i64 3274037854792712831, i64 172590168, i64 sub (i64 ptrtoint (ptr @__profc_end to i64), i64 ptrtoint (ptr @__profd_end to i64)), i64 0, ptr null, ptr null, i32 2, [3 x i16] zeroinitializer, i32 0 }, section "__llvm_prf_data", comdat($__profc_end), align 8
+ at __profd_end = private global { i64, i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i16, i32 } { i64 3274037854792712831, i64 172590168, i64 sub (i64 ptrtoint (ptr @__profc_end to i64), i64 ptrtoint (ptr @__profd_end to i64)), i64 0, i64 0, ptr null, ptr null, i32 2, [3 x i16] zeroinitializer, i16 0, i32 0 }, section "__llvm_prf_data", comdat($__profc_end), align 8
diff --git a/llvm/test/Instrumentation/InstrProfiling/amdgpu-3d-grid.ll b/llvm/test/Instrumentation/InstrProfiling/amdgpu-3d-grid.ll
new file mode 100644
index 0000000000000..f775b9756e878
--- /dev/null
+++ b/llvm/test/Instrumentation/InstrProfiling/amdgpu-3d-grid.ll
@@ -0,0 +1,23 @@
+;; Test that AMDGPU PGO instrumentation generates contiguous counter arrays
+;; and profile section symbols with CUID-based naming. The __llvm_profile_sampling_gpu
+;; library function handles 3D block linearization internally.
+
+; RUN: opt %s -mtriple=amdgcn-amd-amdhsa -passes=instrprof -S | FileCheck %s
+
+ at __hip_cuid_abcdef789 = addrspace(1) global i8 0
+ at __profn_kernel_3d = private constant [9 x i8] c"kernel_3d"
+
+define amdgpu_kernel void @kernel_3d() {
+ call void @llvm.instrprof.increment(ptr @__profn_kernel_3d, i64 12345, i32 1, i32 0)
+ ret void
+}
+
+declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
+
+;; Per-function comdat counters (3D grid linearization is handled in the runtime library)
+; CHECK: @__profc_kernel_3d = linkonce_odr protected addrspace(1) global [1 x i64]
+; CHECK: @__llvm_prf_unifcnt_kernel_3d = linkonce_odr protected addrspace(1) global [1 x i64]
+
+;; Check sampling guard calls library function
+; CHECK: call i32 @__llvm_profile_sampling_gpu(i32 3)
+; CHECK: call void @__llvm_profile_instrument_gpu(ptr addrspacecast (ptr addrspace(1) @__profc_kernel_3d to ptr), ptr addrspacecast (ptr addrspace(1) @__llvm_prf_unifcnt_kernel_3d to ptr), i64 1)
diff --git a/llvm/test/Instrumentation/InstrProfiling/amdgpu-profc-arrays.ll b/llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll
similarity index 52%
rename from llvm/test/Instrumentation/InstrProfiling/amdgpu-profc-arrays.ll
rename to llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll
index eab78fb3591b1..8d48372a12de6 100644
--- a/llvm/test/Instrumentation/InstrProfiling/amdgpu-profc-arrays.ll
+++ b/llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll
@@ -1,13 +1,17 @@
;; Per-kernel __profc_* arrays land in section __llvm_prf_cnts with one slot
-;; per counter, and counter increments lower to __llvm_profile_instrument_gpu
-;; calls whose pointer argument is a GEP into the per-kernel array.
+;; per counter, and a parallel __llvm_prf_unifcnt_* array lands in
+;; __llvm_prf_ucnts. Counter increments lower to __llvm_profile_instrument_gpu
+;; calls whose first pointer argument is a GEP into the per-kernel counter
+;; array and whose second is a GEP into the matching uniform-counter array.
+;; Sampling is disabled so the increment lowers to a direct call.
-; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -passes=instrprof < %s | FileCheck %s
+; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -offload-pgo-sampling=0 -passes=instrprof < %s | FileCheck %s
@__profn_kernel1 = private constant [7 x i8] c"kernel1"
@__profn_kernel2 = private constant [7 x i8] c"kernel2"
; CHECK: @__profc_kernel1 = linkonce_odr protected addrspace(1) global [2 x i64] zeroinitializer, section "__llvm_prf_cnts"
+; CHECK: @__llvm_prf_unifcnt_kernel1 = linkonce_odr protected addrspace(1) global [2 x i64] zeroinitializer, section "__llvm_prf_ucnts"
; CHECK: @__profc_kernel2 = linkonce_odr protected addrspace(1) global [1 x i64] zeroinitializer, section "__llvm_prf_cnts"
define amdgpu_kernel void @kernel1() {
@@ -23,4 +27,6 @@ define amdgpu_kernel void @kernel2() {
declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
-; CHECK: call void @__llvm_profile_instrument_gpu(ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ([2 x i64], ptr addrspace(1) @__profc_kernel1, i32 0, i32 1) to ptr), ptr null, i64 1)
+;; Second counter slot uses a GEP into the per-kernel counter array and the
+;; matching uniform-counter array.
+; CHECK: call void @__llvm_profile_instrument_gpu(ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ([2 x i64], ptr addrspace(1) @__profc_kernel1, i32 0, i32 1) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ([2 x i64], ptr addrspace(1) @__llvm_prf_unifcnt_kernel1, i32 0, i32 1) to ptr), i64 1)
diff --git a/llvm/test/Instrumentation/InstrProfiling/amdgpu-instrumentation.ll b/llvm/test/Instrumentation/InstrProfiling/amdgpu-instrumentation.ll
index efe53ab1ebdfb..963e141934df1 100644
--- a/llvm/test/Instrumentation/InstrProfiling/amdgpu-instrumentation.ll
+++ b/llvm/test/Instrumentation/InstrProfiling/amdgpu-instrumentation.ll
@@ -1,8 +1,16 @@
-;; Test basic AMDGPU PGO instrumentation lowering.
-;; Verifies that each instrumentation point lowers directly to a call to
-;; __llvm_profile_instrument_gpu with a null uniform-counter argument.
+;; Test AMDGPU PGO instrumentation lowering with multiple basic blocks.
+;; Verifies:
+;; 1. Sampling decision is computed once in the entry block.
+;; 2. Each instrumentation point calls __llvm_profile_instrument_gpu behind the
+;; sampling guard branch.
+;; 3. No-sampling mode (sampling=0) calls __llvm_profile_instrument_gpu unconditionally.
-; RUN: opt %s -mtriple=amdgcn-amd-amdhsa -passes=instrprof -S | FileCheck %s
+; RUN: opt %s -mtriple=amdgcn-amd-amdhsa -passes=instrprof \
+; RUN: -offload-pgo-sampling=3 -S \
+; RUN: | FileCheck %s --check-prefix=SAMPLED
+; RUN: opt %s -mtriple=amdgcn-amd-amdhsa -passes=instrprof \
+; RUN: -offload-pgo-sampling=0 -S \
+; RUN: | FileCheck %s --check-prefix=NOSAMPLE
@__hip_cuid_test01 = addrspace(1) global i8 0
@__profn_test_kernel = private constant [11 x i8] c"test_kernel"
@@ -24,9 +32,26 @@ if.end:
declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
-; CHECK-LABEL: define {{.*}} @test_kernel
-; CHECK-NOT: @__llvm_profile_sampling_gpu
-; CHECK: call void @__llvm_profile_instrument_gpu(
-; CHECK-SAME: ptr addrspacecast (ptr addrspace(1) @__profc_test_kernel to ptr), ptr null, i64 1)
-; CHECK: call void @__llvm_profile_instrument_gpu(
-; CHECK-SAME: ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ([4 x i64], ptr addrspace(1) @__profc_test_kernel, i32 0, i32 1) to ptr), ptr null, i64 1)
+;; ---- Sampled mode (sampling=3) ----
+
+;; Entry block: sampling decision computed once
+; SAMPLED-LABEL: define {{.*}} @test_kernel
+; SAMPLED: entry:
+; SAMPLED: %pgo.sampled = call i32 @__llvm_profile_sampling_gpu(i32 3)
+; SAMPLED: %pgo.matched = icmp ne i32 %pgo.sampled, 0
+; SAMPLED: br i1 %pgo.matched, label %po_then, label %po_cont
+
+;; Second instrumentation point reuses same sampling decision
+; SAMPLED: br i1 %pgo.matched, label %po_then{{[0-9]+}}, label %po_cont{{[0-9]+}}
+
+;; Both instrumentation points call the library function
+; SAMPLED: call void @__llvm_profile_instrument_gpu(
+; SAMPLED: call void @__llvm_profile_instrument_gpu(
+
+;; ---- No-sampling mode (sampling=0) ----
+
+;; No sampling guard — direct call
+; NOSAMPLE-LABEL: define {{.*}} @test_kernel
+; NOSAMPLE: entry:
+; NOSAMPLE-NOT: @__llvm_profile_sampling_gpu
+; NOSAMPLE: call void @__llvm_profile_instrument_gpu(
diff --git a/llvm/test/Instrumentation/InstrProfiling/amdgpu-uniform-counters.ll b/llvm/test/Instrumentation/InstrProfiling/amdgpu-uniform-counters.ll
new file mode 100644
index 0000000000000..ffdd41043bdbe
--- /dev/null
+++ b/llvm/test/Instrumentation/InstrProfiling/amdgpu-uniform-counters.ll
@@ -0,0 +1,22 @@
+;; Test that AMDGPU targets generate uniform counter arrays alongside regular
+;; counters. The uniform counter is passed to __llvm_profile_instrument_gpu which
+;; updates it when all lanes in the wave are active.
+
+; RUN: opt %s -mtriple=amdgcn-amd-amdhsa -passes=instrprof -S | FileCheck %s
+
+ at __hip_cuid_test123 = addrspace(1) global i8 0
+ at __profn_test_kernel = private constant [11 x i8] c"test_kernel"
+
+define amdgpu_kernel void @test_kernel() {
+ call void @llvm.instrprof.increment(ptr @__profn_test_kernel, i64 12345, i32 1, i32 0)
+ ret void
+}
+
+declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
+
+;; Per-function counter + uniform-counter globals (comdat)
+; CHECK: @__profc_test_kernel = linkonce_odr protected addrspace(1) global [1 x i64]
+; CHECK: @__llvm_prf_unifcnt_test_kernel = linkonce_odr protected addrspace(1) global [1 x i64]
+
+;; __llvm_profile_instrument_gpu receives counter and uniform-counter bases
+; CHECK: call void @__llvm_profile_instrument_gpu(ptr addrspacecast (ptr addrspace(1) @__profc_test_kernel to ptr), ptr addrspacecast (ptr addrspace(1) @__llvm_prf_unifcnt_test_kernel to ptr), i64 1)
diff --git a/llvm/test/Instrumentation/InstrProfiling/amdgpu-wave32.ll b/llvm/test/Instrumentation/InstrProfiling/amdgpu-wave32.ll
new file mode 100644
index 0000000000000..9cdf733da1c6f
--- /dev/null
+++ b/llvm/test/Instrumentation/InstrProfiling/amdgpu-wave32.ll
@@ -0,0 +1,35 @@
+;; Test that AMDGPU PGO instrumentation generates library calls for Wave32.
+;; Verifies sampling guard and __llvm_profile_instrument_gpu call with correct
+;; addrspace(1) counter pointers.
+
+; RUN: opt %s -mtriple=amdgcn-amd-amdhsa -passes=instrprof -S | FileCheck %s
+
+ at __hip_cuid_abcdef456 = addrspace(1) global i8 0
+ at __profn_kernel_w32 = private constant [10 x i8] c"kernel_w32"
+
+define amdgpu_kernel void @kernel_w32() #0 {
+ call void @llvm.instrprof.increment(ptr @__profn_kernel_w32, i64 12345, i32 1, i32 0)
+ ret void
+}
+
+declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
+
+attributes #0 = { "target-cpu"="gfx1100" }
+
+;; Per-function comdat counters + uniform counters
+; CHECK: @__profc_kernel_w32 = linkonce_odr protected addrspace(1) global [1 x i64]
+; CHECK: @__llvm_prf_unifcnt_kernel_w32 = linkonce_odr protected addrspace(1) global [1 x i64]
+; CHECK: @__profd_kernel_w32 = linkonce_odr protected addrspace(1) global { {{.*}} i16 0, i32 0 }
+
+;; Check wave size stored via intrinsic
+; CHECK: %wavesize.i16 = trunc i32 %{{.*}} to i16
+; CHECK: store i16 %wavesize.i16, ptr addrspace(1) getelementptr inbounds {{.*}} @__profd_kernel_w32
+
+;; Check sampling guard (default sampling=3)
+; CHECK: %pgo.sampled = call i32 @__llvm_profile_sampling_gpu(i32 3)
+; CHECK: %pgo.matched = icmp ne i32 %pgo.sampled, 0
+; CHECK: br i1 %pgo.matched, label %po_then, label %po_cont
+
+;; Check library call
+; CHECK: po_then:
+; CHECK: call void @__llvm_profile_instrument_gpu(ptr addrspacecast (ptr addrspace(1) @__profc_kernel_w32 to ptr), ptr addrspacecast (ptr addrspace(1) @__llvm_prf_unifcnt_kernel_w32 to ptr), i64 1)
diff --git a/llvm/test/Instrumentation/InstrProfiling/amdgpu-wave64.ll b/llvm/test/Instrumentation/InstrProfiling/amdgpu-wave64.ll
new file mode 100644
index 0000000000000..0a776ba5da687
--- /dev/null
+++ b/llvm/test/Instrumentation/InstrProfiling/amdgpu-wave64.ll
@@ -0,0 +1,34 @@
+;; Test that AMDGPU PGO instrumentation generates library calls for Wave64.
+;; Wave64 targets (e.g., gfx908) should embed wave size 64 in profile data.
+
+; RUN: opt %s -mtriple=amdgcn-amd-amdhsa -passes=instrprof -S | FileCheck %s
+
+ at __hip_cuid_abcdef123 = addrspace(1) global i8 0
+ at __profn_kernel_w64 = private constant [10 x i8] c"kernel_w64"
+
+define amdgpu_kernel void @kernel_w64() #0 {
+ call void @llvm.instrprof.increment(ptr @__profn_kernel_w64, i64 12345, i32 1, i32 0)
+ ret void
+}
+
+declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
+
+attributes #0 = { "target-cpu"="gfx908" }
+
+;; Per-function comdat counters
+; CHECK: @__profc_kernel_w64 = linkonce_odr protected addrspace(1) global [1 x i64]
+; CHECK: @__llvm_prf_unifcnt_kernel_w64 = linkonce_odr protected addrspace(1) global [1 x i64]
+; CHECK: @__profd_kernel_w64 = linkonce_odr protected addrspace(1) global { {{.*}} i16 0, i32 0 }
+
+;; Check wave size stored via intrinsic
+; CHECK: %wavesize.i16 = trunc i32 %{{.*}} to i16
+; CHECK: store i16 %wavesize.i16, ptr addrspace(1) getelementptr inbounds {{.*}} @__profd_kernel_w64
+
+;; Check sampling guard
+; CHECK: %pgo.sampled = call i32 @__llvm_profile_sampling_gpu(i32 3)
+; CHECK: %pgo.matched = icmp ne i32 %pgo.sampled, 0
+; CHECK: br i1 %pgo.matched, label %po_then, label %po_cont
+
+;; Check library call
+; CHECK: po_then:
+; CHECK: call void @__llvm_profile_instrument_gpu(ptr addrspacecast (ptr addrspace(1) @__profc_kernel_w64 to ptr), ptr addrspacecast (ptr addrspace(1) @__llvm_prf_unifcnt_kernel_w64 to ptr), i64 1)
diff --git a/llvm/test/Instrumentation/InstrProfiling/coverage.ll b/llvm/test/Instrumentation/InstrProfiling/coverage.ll
index 08cbcaa962b76..695a8829fdf75 100644
--- a/llvm/test/Instrumentation/InstrProfiling/coverage.ll
+++ b/llvm/test/Instrumentation/InstrProfiling/coverage.ll
@@ -5,12 +5,12 @@ target triple = "aarch64-unknown-linux-gnu"
@__profn_foo = private constant [3 x i8] c"foo"
; CHECK: @__profc_foo = private global [1 x i8] c"\FF", section "__llvm_prf_cnts", comdat, align 1
-; CHECK: @__profd_foo = private global { i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i32 } { i64 {{.*}}, i64 {{.*}}, i64 sub (i64 ptrtoint (ptr @__profc_foo to i64)
-; BINARY: @__profd_foo = private global { i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i32 } { i64 {{.*}}, i64 {{.*}}, i64 ptrtoint (ptr @__profc_foo to i64),
+; CHECK: @__profd_foo = private global { i64, i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i16, i32 } { i64 {{.*}}, i64 {{.*}}, i64 sub (i64 ptrtoint (ptr @__profc_foo to i64)
+; BINARY: @__profd_foo = private global { i64, i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i16, i32 } { i64 {{.*}}, i64 {{.*}}, i64 ptrtoint (ptr @__profc_foo to i64),
@__profn_bar = private constant [3 x i8] c"bar"
; CHECK: @__profc_bar = private global [1 x i8] c"\FF", section "__llvm_prf_cnts", comdat, align 1
-; CHECK: @__profd_bar = private global { i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i32 } { i64 {{.*}}, i64 {{.*}}, i64 sub (i64 ptrtoint (ptr @__profc_bar to i64)
-; BINARY: @__profd_bar = private global { i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i32 } { i64 {{.*}}, i64 {{.*}}, i64 ptrtoint (ptr @__profc_bar to i64),
+; CHECK: @__profd_bar = private global { i64, i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i16, i32 } { i64 {{.*}}, i64 {{.*}}, i64 sub (i64 ptrtoint (ptr @__profc_bar to i64)
+; BINARY: @__profd_bar = private global { i64, i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i16, i32 } { i64 {{.*}}, i64 {{.*}}, i64 ptrtoint (ptr @__profc_bar to i64),
; CHECK: @__llvm_prf_nm = {{.*}} section "__llvm_prf_names"
; BINARY: @__llvm_prf_nm ={{.*}} section "__llvm_covnames"
diff --git a/llvm/test/Instrumentation/InstrProfiling/inline-data-var-create.ll b/llvm/test/Instrumentation/InstrProfiling/inline-data-var-create.ll
index 456103164378e..8644c320f867b 100644
--- a/llvm/test/Instrumentation/InstrProfiling/inline-data-var-create.ll
+++ b/llvm/test/Instrumentation/InstrProfiling/inline-data-var-create.ll
@@ -7,17 +7,18 @@
target triple = "x86_64-unknown-linux-gnu"
-; INLINEFIRST: @__profd_foo = private global{{.*}}zeroinitializer, i32 21
-; INLINEFIRST: @__profd_bar = private global{{.*}}zeroinitializer, i32 23
-; INLINEFIRST: @__profd_foobar = private global{{.*}}zeroinitializer, i32 99
-
-; INLINEAFTER: @__profd_foobar = private global{{.*}}zeroinitializer, i32 99
-; INLINEAFTER: @__profd_foo = private global{{.*}}zeroinitializer, i32 21
-; INLINEAFTER: @__profd_bar = private global{{.*}}zeroinitializer, i32 23
-
-; NOINLINE: @__profd_foobar = private global{{.*}}zeroinitializer, i32 99
-; NOINLINE: @__profd_foo = private global{{.*}}zeroinitializer, i32 21
-; NOINLINE: @__profd_bar = private global{{.*}}zeroinitializer, i32 23
+;; Note: struct layout is { ..., i32, [3 x i16], i16, i32 } where the i16 is OffloadDeviceWaveSize
+; INLINEFIRST: @__profd_foo = private global{{.*}}i16 0, i32 21
+; INLINEFIRST: @__profd_bar = private global{{.*}}i16 0, i32 23
+; INLINEFIRST: @__profd_foobar = private global{{.*}}i16 0, i32 99
+
+; INLINEAFTER: @__profd_foobar = private global{{.*}}i16 0, i32 99
+; INLINEAFTER: @__profd_foo = private global{{.*}}i16 0, i32 21
+; INLINEAFTER: @__profd_bar = private global{{.*}}i16 0, i32 23
+
+; NOINLINE: @__profd_foobar = private global{{.*}}i16 0, i32 99
+; NOINLINE: @__profd_foo = private global{{.*}}i16 0, i32 21
+; NOINLINE: @__profd_bar = private global{{.*}}i16 0, i32 23
declare void @llvm.instrprof.increment(ptr %0, i64 %1, i32 %2, i32 %3)
declare void @llvm.instrprof.mcdc.parameters(ptr %0, i64 %1, i32 %2)
diff --git a/llvm/test/Instrumentation/InstrProfiling/platform.ll b/llvm/test/Instrumentation/InstrProfiling/platform.ll
index 9c76a5caf2a51..5f7eb47d2b611 100644
--- a/llvm/test/Instrumentation/InstrProfiling/platform.ll
+++ b/llvm/test/Instrumentation/InstrProfiling/platform.ll
@@ -9,26 +9,35 @@
; RUN: opt < %s -mtriple=x86_64-pc-windows -passes=instrprof -S | FileCheck %s -check-prefix=WINDOWS
; RUN: opt < %s -mtriple=powerpc64-ibm-aix-xcoff -passes=instrprof -S | FileCheck %s -check-prefix=AIX
; RUN: opt < %s -mtriple=arm-elf -passes=instrprof -S | FileCheck %s -check-prefix=BAREMETAL
+; RUN: opt < %s -mtriple=amdgcn-amd-amdhsa -passes=instrprof -S | FileCheck %s -check-prefix=AMDGPU
@__profn_foo = private constant [3 x i8] c"foo"
; MACHO-NOT: __profn_foo
; ELF-NOT: __profn_foo
; WINDOWS-NOT: __profn_foo
; AIX-NOT: __profn_foo
+; AMDGPU-NOT: __profn_foo
; MACHO: @__profc_foo = private global [1 x i64] zeroinitializer, section "__DATA,__llvm_prf_cnts", align 8
; ELF: @__profc_foo = private global [1 x i64] zeroinitializer, section "__llvm_prf_cnts", comdat, align 8
; WINDOWS: @__profc_foo = private global [1 x i64] zeroinitializer, section ".lprfc$M", align 8
; AIX: @__profc_foo = private global [1 x i64] zeroinitializer, section "__llvm_prf_cnts", align 8
+;; AMDGPU without CUID: COMDAT counters + uniform counter slot (HIP-style GPU profiling).
+; AMDGPU: $__llvm_prf_unifcnt_foo = comdat any
+; AMDGPU: @__profc_foo = linkonce_odr protected addrspace(1) global [{{[0-9]+}} x i64] zeroinitializer, section "__llvm_prf_cnts", comdat, align 8
+; AMDGPU: @__llvm_prf_unifcnt_foo = linkonce_odr protected addrspace(1) global [{{[0-9]+}} x i64] zeroinitializer, section "__llvm_prf_ucnts", comdat, align 8
; MACHO: @__profd_foo = private {{.*}}, section "__DATA,__llvm_prf_data,regular,live_support", align 8
; ELF: @__profd_foo = private {{.*}}, section "__llvm_prf_data", comdat($__profc_foo), align 8
; WINDOWS: @__profd_foo = private global {{.*}}, section ".lprfd$M", align 8
; AIX: @__profd_foo = private {{.*}}, section "__llvm_prf_data", align 8
+;; AMDGPU without CUID uses per-function data (not alias)
+; AMDGPU: @__profd_foo = linkonce_odr protected addrspace(1) global {{.*}}, section "__llvm_prf_data", comdat($__profc_foo), align 8
; ELF: @__llvm_prf_nm = private constant [{{.*}} x i8] c"{{.*}}", section "{{.*}}__llvm_prf_names"{{.*}}, align 1
; WINDOWS: @__llvm_prf_nm = private constant [{{.*}} x i8] c"{{.*}}", section "{{.*}}lprfn$M", align 1
; AIX: @__llvm_prf_nm = private constant [{{.*}} x i8] c"{{.*}}", section "{{.*}}__llvm_prf_names", align 1
+; AMDGPU: @__llvm_prf_nm = private addrspace(1) constant [{{.*}} x i8] c"{{.*}}", section "__llvm_prf_names", align 1
define void @foo() {
call void @llvm.instrprof.increment(ptr @__profn_foo, i64 0, i32 1, i32 0)
@@ -37,6 +46,9 @@ define void @foo() {
declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
+;; AMDGPU without CUID uses standard per-function allocation (for OpenMP compatibility)
+;; Start/stop symbols behavior is platform-specific
+
;; Emit registration functions for platforms that don't find the
;; symbols by their sections.
@@ -48,6 +60,7 @@ declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
; WINDOWS-NOT: define internal void @__llvm_profile_register_functions
; AIX-NOT: define internal void @__llvm_profile_register_functions
; BAREMETAL-NOT: define internal void @__llvm_profile_register_functions
+; AMDGPU-NOT: define internal void @__llvm_profile_register_functions
;; PR38340: When dynamic registration is used, we had a bug where we'd register
;; something that's not a __profd_* variable.
@@ -60,3 +73,4 @@ declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
; WINDOWS-NOT: define internal void @__llvm_profile_init
; AIX-NOT: define internal void @__llvm_profile_init
; BAREMETAL-NOT: define internal void @__llvm_profile_init
+; AMDGPU-NOT: define internal void @__llvm_profile_init
diff --git a/llvm/test/Transforms/PGOProfile/Inputs/thinlto_indirect_call_promotion.profraw b/llvm/test/Transforms/PGOProfile/Inputs/thinlto_indirect_call_promotion.profraw
index 3daa98f937b691880ffff203c9426bfacddf749d..596ba3236108eedb5cf7546418d99f9d24f2c7d7 100644
GIT binary patch
delta 207
zcmZ3$a)X7lu_!ISs37M*_e4(ni4GiKB*7ODqoTHZ<^TT at H*Vc at 37eQ^#m4{z1}wKh
zBFnVo7W{_*m<)_|fT(z;rIvt0MF3a at 1H*+SAhRaxF)GO^;83~%taNf3<17V;xeN at 9
IC at w+p0h4(_hX4Qo
delta 174
zcmcb?vVeuNu_!ISs37M**F;YHi3tLH8J8YEkzV!x|K>Vr=_jER>#W!UR)JWP`52`p
zJ`fOF at gD*hpdbMzt0h-}CRG5IVqln|Cj&Adj!{X>0!_&Qi0*levlJi(F)%Pf4Te~Y
G-~#}Nbw4lw
diff --git a/llvm/test/Transforms/PGOProfile/comdat_internal.ll b/llvm/test/Transforms/PGOProfile/comdat_internal.ll
index 1bad0db1b4762..c948ad84392a0 100644
--- a/llvm/test/Transforms/PGOProfile/comdat_internal.ll
+++ b/llvm/test/Transforms/PGOProfile/comdat_internal.ll
@@ -13,9 +13,9 @@ $foo = comdat any
; CHECK: @__llvm_profile_raw_version = hidden constant i64 {{[0-9]+}}, comdat
; CHECK-NOT: __profn__stdin__foo
; CHECK: @__profc__stdin__foo.[[#FOO_HASH]] = private global [1 x i64] zeroinitializer, section "__llvm_prf_cnts", comdat, align 8
-; CHECK: @__profd__stdin__foo.[[#FOO_HASH]] = private global { i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i32 } { i64 {{.*}}, i64 [[#FOO_HASH]], i64 sub (i64 ptrtoint (ptr @__profc__stdin__foo.742261418966908927 to i64), i64 ptrtoint (ptr @__profd__stdin__foo.742261418966908927 to i64)), i64 0, ptr null
+; CHECK: @__profd__stdin__foo.[[#FOO_HASH]] = private global { i64, i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i16, i32 } { i64 {{.*}}, i64 [[#FOO_HASH]], i64 sub (i64 ptrtoint (ptr @__profc__stdin__foo.742261418966908927 to i64), i64 ptrtoint (ptr @__profd__stdin__foo.742261418966908927 to i64)), i64 0, i64 0, ptr null
; CHECK-NOT: @foo
-; CHECK-SAME: , ptr null, i32 1, [3 x i16] zeroinitializer, i32 0 }, section "__llvm_prf_data", comdat($__profc__stdin__foo.[[#FOO_HASH]]), align 8
+; CHECK-SAME: , ptr null, i32 1, [3 x i16] zeroinitializer, i16 0, i32 0 }, section "__llvm_prf_data", comdat($__profc__stdin__foo.[[#FOO_HASH]]), align 8
; CHECK: @__llvm_prf_nm
; CHECK: @llvm.compiler.used
diff --git a/llvm/test/Transforms/PGOProfile/instrprof_burst_sampling_fast.ll b/llvm/test/Transforms/PGOProfile/instrprof_burst_sampling_fast.ll
index 56d8364d8f543..7261cc24b0dcf 100644
--- a/llvm/test/Transforms/PGOProfile/instrprof_burst_sampling_fast.ll
+++ b/llvm/test/Transforms/PGOProfile/instrprof_burst_sampling_fast.ll
@@ -14,7 +14,7 @@ $__llvm_profile_raw_version = comdat any
; SAMPLE-VAR: @__llvm_profile_sampling = thread_local global i16 0, comdat
; SAMPLE-VAR: @__profc_f = private global [1 x i64] zeroinitializer, section "__llvm_prf_cnts", comdat, align 8
-; SAMPLE-VAR: @__profd_f = private global { i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i32 } { i64 -3706093650706652785, i64 12884901887, i64 sub (i64 ptrtoint (ptr @__profc_f to i64), i64 ptrtoint (ptr @__profd_f to i64)), i64 0, ptr @f.local, ptr null, i32 1, [3 x i16] zeroinitializer, i32 0 }, section "__llvm_prf_data", comdat($__profc_f), align 8
+; SAMPLE-VAR: @__profd_f = private global { i64, i64, i64, i64, i64, ptr, ptr, i32, [3 x i16], i16, i32 } { i64 -3706093650706652785, i64 12884901887, i64 sub (i64 ptrtoint (ptr @__profc_f to i64), i64 ptrtoint (ptr @__profd_f to i64)), i64 0, i64 0, ptr @f.local, ptr null, i32 1, [3 x i16] zeroinitializer, i16 0, i32 0 }, section "__llvm_prf_data", comdat($__profc_f), align 8
; SAMPLE-VAR: @__llvm_prf_nm = private constant {{.*}}, section "__llvm_prf_names", align 1
; SAMPLE-VAR: @llvm.compiler.used = appending global [2 x ptr] [ptr @__llvm_profile_sampling, ptr @__profd_f], section "llvm.metadata"
; SAMPLE-VAR: @llvm.used = appending global [1 x ptr] [ptr @__llvm_prf_nm], section "llvm.metadata"
diff --git a/llvm/test/Transforms/PGOProfile/vtable_profile.ll b/llvm/test/Transforms/PGOProfile/vtable_profile.ll
index aae1e2d8b4e49..0c554db05cfb4 100644
--- a/llvm/test/Transforms/PGOProfile/vtable_profile.ll
+++ b/llvm/test/Transforms/PGOProfile/vtable_profile.ll
@@ -49,7 +49,7 @@ target triple = "x86_64-unknown-linux-gnu"
@llvm.compiler.used = appending global [1 x ptr] [ptr @_ZTV5Base1], section "llvm.metadata"
; GEN: __llvm_profile_raw_version = comdat any
-; GEN: __llvm_profile_raw_version = hidden constant i64 72057594037927946, comdat
+; GEN: __llvm_profile_raw_version = hidden constant i64 72057594037927947, comdat
; GEN: __profn__Z4funci = private constant [8 x i8] c"_Z4funci"
; LOWER: $__profvt__ZTV7Derived = comdat nodeduplicate
diff --git a/llvm/test/tools/llvm-profdata/Inputs/c-general.profraw b/llvm/test/tools/llvm-profdata/Inputs/c-general.profraw
index a3e884343942ebc70ba95ab4ee006630b6816d80..dec90151a79352bf898c8c9d1d2d23eaa4bf5eed 100644
GIT binary patch
delta 197
zcmeys|3ZMXu_!ISs37M*_e4(ni4GiKB*Ev<^}llG%K!gm%^2F;!Y0QtOHFJr0Lg8T
zXJQeUI8h$N{vgl9ZUA9VHei%zVt{c97!??2KqWQ+CD>j-IUj%$8j~+Fg0xRIV3KFz
n34rkmm=qWbV4Mv=DGmo1?*mXu0>&|5R$#mVF>Z4Ka{~(iCWJse
delta 160
zcmaDM at PVJRu_!ISs37M**F;YHi3tLH4>VUWO0D|;f9~Aye@{av)_G05AkX9?G4Z4P
zWC=!jCUpl8D*?#5-~nQ-U{qk-0b+ar3T+dZypTzLvIUa at V+5Ghz$DMKrUT480TkQ-
nW(hFMGgW{ECteVkoWU%>m;n}Fz^uS%0Afv?z_IxQa|H_kt<yS_
diff --git a/llvm/test/tools/llvm-profdata/Inputs/compressed.profraw b/llvm/test/tools/llvm-profdata/Inputs/compressed.profraw
index e3f77e870d4d20828119348e70eb44e6d39e0ec0..778e80fce2691a35d9654748763b0aa8c233ec13 100644
GIT binary patch
delta 197
zcmX at Wze9ktu_!ISs37M*_e4(ni4GiKB*EwKYmREqs{j9Gk0sqI4xJpwEH$yg03^3T
zo{2?Z;zW56`-40ay8(nf*?>`=i2=qbU{qk70hQPQlwf-S<$M52XiUDy2+}^;fJvT-
nCjiDPU{YW#fN?eer8pd5ybnMr2^hzKS%L8e#JJ4`%mFL_d2B$e
delta 159
zcmdlXaDbn)u_!ISs37M**F;YHi3tL{4=Vq*Em-yc|6E_D%J_-3-V-m#Gr34i{3t(J
zf>EAH-2uc(0J1K4fLJRS6&QDb7$1N_+XN;rWRjn3!KA<#0cJHY$uq6#05eYj1vh|M
m0?hJE6=1=M7X&6}FbgndfW;RuD=->>SQ963Z2rLP!2$q9VL7b;
diff --git a/llvm/test/tools/llvm-profdata/binary-ids-padding.test b/llvm/test/tools/llvm-profdata/binary-ids-padding.test
index cc3e6c38e6907..f20161f9886e3 100644
--- a/llvm/test/tools/llvm-profdata/binary-ids-padding.test
+++ b/llvm/test/tools/llvm-profdata/binary-ids-padding.test
@@ -18,7 +18,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\201rforpl\377' > %t.profraw
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t.profraw
// There will be 2 20-byte binary IDs, so the total Binary IDs size will be 64 bytes.
// 2 * 8 binary ID sizes
// + 2 * 20 binary IDs (of size 20)
@@ -32,6 +32,9 @@ RUN: printf '\3\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\20\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\4\0\1\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
@@ -64,12 +67,14 @@ RUN: printf '\0\0\4\0\1\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\067\265\035\031\112\165\023\344' >> %t.profraw
RUN: printf '\02\0\0\0\0\0\0\0' >> %t.profraw
-RUN: printf '\310\377\3\0\1\0\0\0' >> %t.profraw
+RUN: printf '\270\377\3\0\1\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
@@ -81,8 +86,8 @@ RUN: printf '\067\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\101\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\7\0foo\1bar\0\0\0\0\0\0\0' >> %t.profraw
-// RUN: llvm-profdata show --binary-ids %t.profraw | FileCheck %s
-// CHECK: Instrumentation level: Front-end
-// CHECK: Binary IDs:
-// CHECK-NEXT: 0001020304050607000102030405060700010203
-// CHECK-NEXT: 0101010101010101020202020202020203030303
+RUN: llvm-profdata show --binary-ids %t.profraw | FileCheck %s
+CHECK: Instrumentation level: Front-end
+CHECK: Binary IDs:
+CHECK-NEXT: 0001020304050607000102030405060700010203
+CHECK-NEXT: 0101010101010101020202020202020203030303
diff --git a/llvm/test/tools/llvm-profdata/large-binary-id-size.test b/llvm/test/tools/llvm-profdata/large-binary-id-size.test
index b62bdad4ddb21..4ffcc89c421ee 100644
--- a/llvm/test/tools/llvm-profdata/large-binary-id-size.test
+++ b/llvm/test/tools/llvm-profdata/large-binary-id-size.test
@@ -2,7 +2,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\201rforpl\377' > %t.profraw
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\40\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
@@ -17,6 +17,9 @@ RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
// Check for a corrupted size being too large past the end of the file.
RUN: printf '\7\7\7\7\7\7\7\7' >> %t.profraw
@@ -24,5 +27,5 @@ RUN: printf '\0\1\2\3\4\5\6\7' >> %t.profraw
RUN: printf '\0\1\2\3\4\5\6\7' >> %t.profraw
RUN: printf '\0\1\2\3\0\0\0\0' >> %t.profraw
-// RUN: not llvm-profdata show --binary-ids %t.profraw 2>&1 | FileCheck %s
-// CHECK: malformed instrumentation profile data: not enough data to read binary id data
+RUN: not llvm-profdata show --binary-ids %t.profraw 2>&1 | FileCheck %s
+CHECK: malformed instrumentation profile data: not enough data to read binary id data
diff --git a/llvm/test/tools/llvm-profdata/malformed-not-space-for-another-header.test b/llvm/test/tools/llvm-profdata/malformed-not-space-for-another-header.test
index 705e5efaf5875..4f9876e4d8fb2 100644
--- a/llvm/test/tools/llvm-profdata/malformed-not-space-for-another-header.test
+++ b/llvm/test/tools/llvm-profdata/malformed-not-space-for-another-header.test
@@ -18,7 +18,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\201rforpl\377' > %t.profraw
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
@@ -26,6 +26,9 @@ RUN: printf '\1\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\10\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\4\0\1\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
@@ -46,10 +49,12 @@ RUN: printf '\254\275\030\333\114\302\370\134' >> %t.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\4\0\1\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\4\0\3\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\023\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\3\0foo\0\0\0' >> %t.profraw
diff --git a/llvm/test/tools/llvm-profdata/malformed-num-counters-zero.test b/llvm/test/tools/llvm-profdata/malformed-num-counters-zero.test
index 157c13b926a7e..e7c2fa7877c1a 100644
--- a/llvm/test/tools/llvm-profdata/malformed-num-counters-zero.test
+++ b/llvm/test/tools/llvm-profdata/malformed-num-counters-zero.test
@@ -18,7 +18,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\201rforpl\377' > %t.profraw
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
@@ -26,6 +26,9 @@ RUN: printf '\1\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\10\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\4\0\1\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
@@ -56,6 +59,7 @@ RUN: cp %t.profraw %t-good.profraw
// Make NumCounters = 0 so that we get "number of counters is zero" error message
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\023\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\3\0foo\0\0\0' >> %t.profraw
@@ -64,12 +68,14 @@ RUN: not llvm-profdata show %t.profraw 2>&1 | FileCheck %s --check-prefix=ZERO
ZERO: malformed instrumentation profile data: number of counters is zero
// Test a counter value greater than 2^56.
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bad.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t-bad.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bad.profraw
// Counter value is 72057594037927937
RUN: printf '\1\0\0\0\0\0\0\1' >> %t-bad.profraw
RUN: printf '\3\0foo\0\0\0' >> %t-bad.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-good.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t-good.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-good.profraw
// Counter value is 72057594037927937
diff --git a/llvm/test/tools/llvm-profdata/malformed-ptr-to-counter-array.test b/llvm/test/tools/llvm-profdata/malformed-ptr-to-counter-array.test
index 83cf76f68fb63..ee4b63986e43f 100644
--- a/llvm/test/tools/llvm-profdata/malformed-ptr-to-counter-array.test
+++ b/llvm/test/tools/llvm-profdata/malformed-ptr-to-counter-array.test
@@ -18,7 +18,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\201rforpl\377' > %t.profraw
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
@@ -26,6 +26,9 @@ RUN: printf '\2\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\10\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\6\0\1\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
@@ -51,11 +54,12 @@ RUN: printf '\02\0\0\0\0\0\0\0' >> %t.profraw
// the profile reader should error out.
RUN: printf '\11\0\6\0\1\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
-
+RUN: printf '\0\0\6\0\2\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\02\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
// Counter Section
diff --git a/llvm/test/tools/llvm-profdata/misaligned-binary-ids-size.test b/llvm/test/tools/llvm-profdata/misaligned-binary-ids-size.test
index 0f20a1b0b369c..b43b2408869c2 100644
--- a/llvm/test/tools/llvm-profdata/misaligned-binary-ids-size.test
+++ b/llvm/test/tools/llvm-profdata/misaligned-binary-ids-size.test
@@ -2,7 +2,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\201rforpl\377' > %t.profraw
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t.profraw
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t.profraw
// We should fail on this because the binary IDs is not a multiple of 8 bytes.
RUN: printf '\77\0\0\0\0\0\0\0' >> %t.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t.profraw
diff --git a/llvm/test/tools/llvm-profdata/mismatched-raw-profile-header.test b/llvm/test/tools/llvm-profdata/mismatched-raw-profile-header.test
index dfa163f1f3439..8d71914a99995 100644
--- a/llvm/test/tools/llvm-profdata/mismatched-raw-profile-header.test
+++ b/llvm/test/tools/llvm-profdata/mismatched-raw-profile-header.test
@@ -11,6 +11,9 @@ RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\20' >> %t
RUN: printf '\0\0\0\1\0\4\0\0' >> %t
RUN: printf '\0\0\0\2\0\4\0\0' >> %t
diff --git a/llvm/test/tools/llvm-profdata/profile-version.test b/llvm/test/tools/llvm-profdata/profile-version.test
index e811699ac63ed..7f627cb89c7d0 100644
--- a/llvm/test/tools/llvm-profdata/profile-version.test
+++ b/llvm/test/tools/llvm-profdata/profile-version.test
@@ -2,7 +2,7 @@ Test the profile version.
RUN: llvm-profdata merge -o %t.profdata %p/Inputs/basic.proftext
RUN: llvm-profdata show --profile-version %t.profdata | FileCheck %s
-CHECK: Profile version: 13
+CHECK: Profile version: 14
RUN: llvm-profdata merge -o %t.prev.profdata %p/Inputs/basic.proftext --write-prev-version
RUN: llvm-profdata show --profile-version %t.prev.profdata | FileCheck %s --check-prefix=PREV
diff --git a/llvm/test/tools/llvm-profdata/raw-32-bits-be.test b/llvm/test/tools/llvm-profdata/raw-32-bits-be.test
index 6b3dc96b36270..4f24eb624a2a7 100644
--- a/llvm/test/tools/llvm-profdata/raw-32-bits-be.test
+++ b/llvm/test/tools/llvm-profdata/raw-32-bits-be.test
@@ -3,7 +3,7 @@
UNSUPPORTED: system-windows
// Header
RUN: printf '\377lprofR\201' > %t
-RUN: printf '\0\0\0\0\0\0\0\12' >> %t
+RUN: printf '\0\0\0\0\0\0\0\13' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\2' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
@@ -11,6 +11,9 @@ RUN: printf '\0\0\0\0\0\0\0\3' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\4' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\20' >> %t
RUN: printf '\0\0\0\0\1\0\0\0' >> %t
RUN: printf '\0\0\0\0\3\0\0\0' >> %t
@@ -22,22 +25,28 @@ RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\134\370\302\114\333\030\275\254' >> %t
RUN: printf '\0\0\0\0\0\0\0\1' >> %t
RUN: printf '\1\0\0\0' >> %t
+RUN: printf '\0\0\0\0' >> %t
RUN: printf '\3\0\0\0' >> %t
RUN: printf '\0\0\0\0' >> %t
RUN: printf '\0\0\0\0' >> %t
RUN: printf '\0\0\0\1' >> %t
-RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0' >> %t
RUN: printf '\0\0\0\3' >> %t
+RUN: printf '\0\0\0\0' >> %t
RUN: printf '\344\023\165\112\031\035\265\067' >> %t
RUN: printf '\0\0\0\0\0\0\0\2' >> %t
-RUN: printf '\0\377\377\330' >> %t
-RUN: printf '\2\377\377\323' >> %t
+RUN: printf '\0\377\377\320' >> %t
+RUN: printf '\0\0\0\0' >> %t
+RUN: printf '\2\377\377\313' >> %t
RUN: printf '\0\0\0\0' >> %t
RUN: printf '\0\0\0\0' >> %t
RUN: printf '\0\0\0\2' >> %t
-RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0' >> %t
RUN: printf '\0\0\0\1' >> %t
+RUN: printf '\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\023' >> %t
RUN: printf '\0\0\0\0\0\0\0\067' >> %t
diff --git a/llvm/test/tools/llvm-profdata/raw-32-bits-le.test b/llvm/test/tools/llvm-profdata/raw-32-bits-le.test
index 95625565f5c0c..96db225b16b78 100644
--- a/llvm/test/tools/llvm-profdata/raw-32-bits-le.test
+++ b/llvm/test/tools/llvm-profdata/raw-32-bits-le.test
@@ -2,7 +2,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\201Rforpl\377' > %t
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\2\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
@@ -10,6 +10,9 @@ RUN: printf '\3\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\4\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\20\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\1\0\0\0\0' >> %t
RUN: printf '\0\0\0\3\0\0\0\0' >> %t
@@ -21,22 +24,28 @@ RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\254\275\030\333\114\302\370\134' >> %t
RUN: printf '\1\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\1' >> %t
+RUN: printf '\0\0\0\0' >> %t
RUN: printf '\0\0\0\3' >> %t
RUN: printf '\0\0\0\0' >> %t
RUN: printf '\0\0\0\0' >> %t
RUN: printf '\1\0\0\0' >> %t
-RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0' >> %t
RUN: printf '\3\0\0\0' >> %t
+RUN: printf '\0\0\0\0' >> %t
RUN: printf '\067\265\035\031\112\165\023\344' >> %t
RUN: printf '\02\0\0\0\0\0\0\0' >> %t
-RUN: printf '\330\377\377\0' >> %t
-RUN: printf '\323\377\377\2' >> %t
+RUN: printf '\320\377\377\0' >> %t
+RUN: printf '\0\0\0\0' >> %t
+RUN: printf '\313\377\377\2' >> %t
RUN: printf '\0\0\0\0' >> %t
RUN: printf '\0\0\0\0' >> %t
RUN: printf '\2\0\0\0' >> %t
-RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0' >> %t
RUN: printf '\1\0\0\0' >> %t
+RUN: printf '\0\0\0\0' >> %t
RUN: printf '\023\0\0\0\0\0\0\0' >> %t
RUN: printf '\067\0\0\0\0\0\0\0' >> %t
diff --git a/llvm/test/tools/llvm-profdata/raw-64-bits-be.test b/llvm/test/tools/llvm-profdata/raw-64-bits-be.test
index 5316ef5a3e559..8d29323ff472a 100644
--- a/llvm/test/tools/llvm-profdata/raw-64-bits-be.test
+++ b/llvm/test/tools/llvm-profdata/raw-64-bits-be.test
@@ -2,7 +2,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\377lprofr\201' > %t
-RUN: printf '\0\0\0\0\0\0\0\12' >> %t
+RUN: printf '\0\0\0\0\0\0\0\13' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\2' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
@@ -10,6 +10,9 @@ RUN: printf '\0\0\0\0\0\0\0\3' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\4' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\20' >> %t
RUN: printf '\0\0\0\1\0\4\0\0' >> %t
RUN: printf '\0\0\0\3\0\4\0\0' >> %t
@@ -21,22 +24,22 @@ RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\134\370\302\114\333\030\275\254' >> %t
RUN: printf '\0\0\0\0\0\0\0\1' >> %t
RUN: printf '\0\0\0\1\0\4\0\0' >> %t
-RUN: printf '\0\0\0\3\0\4\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\3\0\4\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
-RUN: printf '\0\0\0\1' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
-RUN: printf '\0\0\0\3' >> %t
+RUN: printf '\0\0\0\1\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\3' >> %t
RUN: printf '\344\023\165\112\031\035\265\067' >> %t
-RUN: printf '\0\0\0\0\0\0\0\02' >> %t
-RUN: printf '\0\0\0\1\0\3\377\310' >> %t
-RUN: printf '\0\0\0\3\0\3\377\303' >> %t
+RUN: printf '\0\0\0\0\0\0\0\2' >> %t
+RUN: printf '\0\0\0\1\0\3\377\300' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\3\0\3\377\273' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
-RUN: printf '\0\0\0\02' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
-RUN: printf '\0\0\0\1' >> %t
+RUN: printf '\0\0\0\2\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\1' >> %t
RUN: printf '\0\0\0\0\0\0\0\023' >> %t
RUN: printf '\0\0\0\0\0\0\0\067' >> %t
diff --git a/llvm/test/tools/llvm-profdata/raw-64-bits-le.test b/llvm/test/tools/llvm-profdata/raw-64-bits-le.test
index 58f4da8cf0386..d90c18cd31b12 100644
--- a/llvm/test/tools/llvm-profdata/raw-64-bits-le.test
+++ b/llvm/test/tools/llvm-profdata/raw-64-bits-le.test
@@ -2,7 +2,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\201rforpl\377' > %t
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\2\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
@@ -10,6 +10,9 @@ RUN: printf '\3\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\4\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\20\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\4\0\1\0\0\0' >> %t
RUN: printf '\0\0\4\0\3\0\0\0' >> %t
@@ -21,22 +24,22 @@ RUN: printf '\0\0\0\0\0\0\0\0' >> %t
RUN: printf '\254\275\030\333\114\302\370\134' >> %t
RUN: printf '\1\0\0\0\0\0\0\0' >> %t
RUN: printf '\0\0\4\0\1\0\0\0' >> %t
-RUN: printf '\0\0\4\0\3\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\4\0\3\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
-RUN: printf '\1\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
-RUN: printf '\3\0\0\0' >> %t
+RUN: printf '\1\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\3\0\0\0' >> %t
RUN: printf '\067\265\035\031\112\165\023\344' >> %t
RUN: printf '\02\0\0\0\0\0\0\0' >> %t
-RUN: printf '\310\377\3\0\1\0\0\0' >> %t
-RUN: printf '\303\377\3\0\3\0\0\0' >> %t
+RUN: printf '\300\377\3\0\1\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
+RUN: printf '\273\377\3\0\3\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
-RUN: printf '\02\0\0\0' >> %t
RUN: printf '\0\0\0\0\0\0\0\0' >> %t
-RUN: printf '\1\0\0\0' >> %t
+RUN: printf '\02\0\0\0\0\0\0\0' >> %t
+RUN: printf '\0\0\0\0\1\0\0\0' >> %t
RUN: printf '\023\0\0\0\0\0\0\0' >> %t
RUN: printf '\067\0\0\0\0\0\0\0' >> %t
diff --git a/llvm/test/tools/llvm-profdata/raw-two-profiles.test b/llvm/test/tools/llvm-profdata/raw-two-profiles.test
index 47cc6fa4fd7fe..e14e70e747a9a 100644
--- a/llvm/test/tools/llvm-profdata/raw-two-profiles.test
+++ b/llvm/test/tools/llvm-profdata/raw-two-profiles.test
@@ -2,7 +2,7 @@
// whenever \12 (LF) is in the input string.
UNSUPPORTED: system-windows
RUN: printf '\201rforpl\377' > %t-foo.profraw
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t-foo.profraw
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
@@ -10,6 +10,9 @@ RUN: printf '\1\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\10\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\4\0\1\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
@@ -24,6 +27,7 @@ RUN: printf '\0\0\4\0\1\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-foo.profraw
@@ -31,7 +35,7 @@ RUN: printf '\023\0\0\0\0\0\0\0' >> %t-foo.profraw
RUN: printf '\3\0foo\0\0\0' >> %t-foo.profraw
RUN: printf '\201rforpl\377' > %t-bar.profraw
-RUN: printf '\12\0\0\0\0\0\0\0' >> %t-bar.profraw
+RUN: printf '\13\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\1\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
@@ -39,6 +43,9 @@ RUN: printf '\2\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\10\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\6\0\1\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
@@ -53,6 +60,7 @@ RUN: printf '\0\0\6\0\1\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
+RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\02\0\0\0\0\0\0\0' >> %t-bar.profraw
RUN: printf '\0\0\0\0\0\0\0\0' >> %t-bar.profraw
diff --git a/llvm/tools/llvm-profdata/llvm-profdata.cpp b/llvm/tools/llvm-profdata/llvm-profdata.cpp
index 6bf3179696e4a..67396f409a93a 100644
--- a/llvm/tools/llvm-profdata/llvm-profdata.cpp
+++ b/llvm/tools/llvm-profdata/llvm-profdata.cpp
@@ -31,6 +31,7 @@
#include "llvm/Support/BalancedPartitioning.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/Discriminator.h"
+#include "llvm/Support/Endian.h"
#include "llvm/Support/Errc.h"
#include "llvm/Support/FileSystem.h"
#include "llvm/Support/Format.h"
@@ -825,6 +826,22 @@ loadInput(const WeightedFile &Input, SymbolRemapper *Remapper,
I.Name = (*Remapper)(I.Name);
const StringRef FuncName = I.Name;
bool Reported = false;
+
+ if (!I.UniformCounts.empty()) {
+ size_t NumBlocks = I.Counts.size();
+ I.UniformityBits.resize((NumBlocks + 7) / 8, 0xFF);
+
+ for (size_t BlockIdx = 0; BlockIdx < NumBlocks; ++BlockIdx) {
+ uint64_t TotalCount = I.Counts[BlockIdx];
+ uint64_t UniformCount = I.UniformCounts[BlockIdx];
+
+ bool IsUniform =
+ (TotalCount == 0) || ((double)UniformCount / TotalCount >= 0.9);
+ if (!IsUniform)
+ I.UniformityBits[BlockIdx / 8] &= ~(1 << (BlockIdx % 8));
+ }
+ }
+
WC->Writer.addRecord(std::move(I), Input.Weight, [&](Error E) {
if (Reported) {
consumeError(std::move(E));
@@ -2979,6 +2996,16 @@ static int showInstrProfile(ShowFormat SFormat, raw_fd_ostream &OS) {
OS << (I == Start ? "" : ", ") << Func.Counts[I];
}
OS << "]\n";
+
+ // Show uniformity bits if present
+ if (!Func.UniformityBits.empty()) {
+ OS << " Block uniformity: [";
+ for (size_t I = Start, E = Func.Counts.size(); I < E; ++I) {
+ bool IsUniform = Func.isBlockUniform(I);
+ OS << (I == Start ? "" : ", ") << (IsUniform ? "U" : "D");
+ }
+ OS << "]\n";
+ }
}
if (ShowIndirectCallTargets) {
diff --git a/offload/plugins-nextgen/common/include/GlobalHandler.h b/offload/plugins-nextgen/common/include/GlobalHandler.h
index fc8d6fe384754..8f333fcf1056e 100644
--- a/offload/plugins-nextgen/common/include/GlobalHandler.h
+++ b/offload/plugins-nextgen/common/include/GlobalHandler.h
@@ -72,11 +72,15 @@ struct __llvm_profile_gpu_sections {
};
extern "C" {
-extern int __attribute__((weak)) __llvm_write_custom_profile(
- const char *Target, const __llvm_profile_data *DataBegin,
- const __llvm_profile_data *DataEnd, const char *CountersBegin,
- const char *CountersEnd, const char *NamesBegin, const char *NamesEnd,
- const uint64_t *VersionOverride);
+extern int __attribute__((weak))
+__llvm_write_custom_profile(const char *Target,
+ const __llvm_profile_data *DataBegin,
+ const __llvm_profile_data *DataEnd,
+ const char *CountersBegin, const char *CountersEnd,
+ const char *UniformCountersBegin,
+ const char *UniformCountersEnd,
+ const char *NamesBegin, const char *NamesEnd,
+ const uint64_t *VersionOverride);
}
/// PGO profiling data extracted from a GPU device via __llvm_profile_sections.
struct GPUProfGlobals {
diff --git a/offload/plugins-nextgen/common/src/GlobalHandler.cpp b/offload/plugins-nextgen/common/src/GlobalHandler.cpp
index 9216834b1e15e..9b7f9bfec3163 100644
--- a/offload/plugins-nextgen/common/src/GlobalHandler.cpp
+++ b/offload/plugins-nextgen/common/src/GlobalHandler.cpp
@@ -295,7 +295,9 @@ Error GPUProfGlobals::write() const {
reinterpret_cast<const __llvm_profile_data *>(DataBegin),
reinterpret_cast<const __llvm_profile_data *>(DataBegin +
DataSection.size()),
- CountersBegin, CountersBegin + CountersSection.size(), NamesBegin,
+ CountersBegin, CountersBegin + CountersSection.size(),
+ /*UniformCountersBegin=*/nullptr,
+ /*UniformCountersEnd=*/nullptr, NamesBegin,
NamesBegin + NamesSection.size(), &Version);
if (Result != 0)
return Plugin::error(ErrorCode::HOST_IO,
>From f612d4136e110c8f5afb18f4799d9de38a6d2350 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Thu, 5 Feb 2026 11:52:20 -0500
Subject: [PATCH 2/2] [CodeGen] Add BlockUniformityProfile for PGO-guided spill
placement
Add BlockUniformityProfile analysis that reads block uniformity metadata
from PGO profiles to guide register allocation decisions.
When PGO data indicates a basic block executes divergently (partial wave
on AMDGPU), SpillPlacement flattens its frequency to the function entry
frequency. This prevents PGO from biasing spill decisions toward blocks
that appear hot in profiles but are actually divergent at runtime.
Key components:
- PGOInstrumentation: emits block.uniformity.profile metadata during the
PGO-use phase, keyed by a fixed metadata kind (MD_block_uniformity_profile)
- BlockUniformityProfile: Reads !block-uniformity-profile metadata
- BlockUniformityProfileProxy: Analysis pass for new pass manager
- BlockUniformityProfilePrinterPass: Debug printer pass
- Integration with SpillPlacement to flatten divergent block frequencies
The metadata is attached during PGO-use phase by PGOInstrumentation.cpp
and uses a fixed metadata kind (MD_block_uniformity_profile) for efficient
lookup.
---
.../llvm/CodeGen/BlockUniformityProfile.h | 79 ++++++++++++
llvm/include/llvm/CodeGen/SpillPlacement.h | 4 +-
llvm/include/llvm/IR/FixedMetadataKinds.def | 1 +
llvm/include/llvm/Passes/CodeGenPassBuilder.h | 1 +
.../llvm/Passes/MachinePassRegistry.def | 4 +
llvm/lib/CodeGen/BlockUniformityProfile.cpp | 121 ++++++++++++++++++
llvm/lib/CodeGen/CMakeLists.txt | 1 +
llvm/lib/CodeGen/SpillPlacement.cpp | 24 +++-
llvm/lib/Passes/PassBuilder.cpp | 1 +
.../Instrumentation/PGOInstrumentation.cpp | 40 ++++++
.../AMDGPU/block-uniformity-profile.ll | 79 ++++++++++++
11 files changed, 349 insertions(+), 6 deletions(-)
create mode 100644 llvm/include/llvm/CodeGen/BlockUniformityProfile.h
create mode 100644 llvm/lib/CodeGen/BlockUniformityProfile.cpp
create mode 100644 llvm/test/CodeGen/AMDGPU/block-uniformity-profile.ll
diff --git a/llvm/include/llvm/CodeGen/BlockUniformityProfile.h b/llvm/include/llvm/CodeGen/BlockUniformityProfile.h
new file mode 100644
index 0000000000000..7d4382494e9e7
--- /dev/null
+++ b/llvm/include/llvm/CodeGen/BlockUniformityProfile.h
@@ -0,0 +1,79 @@
+//===- BlockUniformityProfile.h - Block uniformity from PGO -*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Provide per-(Machine)basic-block uniformity information from PGO profiles.
+//
+// The source of truth is IR metadata attached during PGO use:
+// - Metadata name: "block-uniformity-profile"
+// - Payload: i1 (true = uniform, false = divergent)
+//
+// This is intentionally target-agnostic: any backend that produces
+// uniformity bits in the profile can attach the same metadata and reuse this
+// proxy in codegen.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CODEGEN_BLOCKUNIFORMITYPROFILE_H
+#define LLVM_CODEGEN_BLOCKUNIFORMITYPROFILE_H
+
+#include "llvm/ADT/BitVector.h"
+#include "llvm/CodeGen/MachineFunctionAnalysis.h"
+#include "llvm/CodeGen/MachineFunctionAnalysisManager.h"
+#include "llvm/CodeGen/MachinePassManager.h"
+#include "llvm/Support/Compiler.h"
+
+namespace llvm {
+
+class MachineBasicBlock;
+class MachineFunction;
+class raw_ostream;
+
+class BlockUniformityProfile {
+public:
+ LLVM_ABI void compute(const MachineFunction &MF);
+
+ bool hasProfile() const { return HasProfile; }
+
+ // Returns true if the block is considered divergent. If profile exists for
+ // the function but a block has no explicit annotation, it is treated as
+ // divergent (conservative).
+ LLVM_ABI bool isDivergent(const MachineBasicBlock &MBB) const;
+
+ LLVM_ABI void print(raw_ostream &OS, const MachineFunction &MF) const;
+
+private:
+ bool HasProfile = false;
+ unsigned NumBlockIDs = 0;
+ BitVector DivergentBlocks;
+};
+
+class BlockUniformityProfileProxy
+ : public AnalysisInfoMixin<BlockUniformityProfileProxy> {
+ friend AnalysisInfoMixin<BlockUniformityProfileProxy>;
+ static AnalysisKey Key;
+
+public:
+ using Result = BlockUniformityProfile;
+ LLVM_ABI Result run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &MFAM);
+};
+
+class BlockUniformityProfilePrinterPass
+ : public PassInfoMixin<BlockUniformityProfilePrinterPass> {
+ raw_ostream &OS;
+
+public:
+ explicit BlockUniformityProfilePrinterPass(raw_ostream &OS) : OS(OS) {}
+ PreservedAnalyses run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &MFAM);
+ static bool isRequired() { return true; }
+};
+
+} // end namespace llvm
+
+#endif // LLVM_CODEGEN_BLOCKUNIFORMITYPROFILE_H
diff --git a/llvm/include/llvm/CodeGen/SpillPlacement.h b/llvm/include/llvm/CodeGen/SpillPlacement.h
index bc206075d9ba4..857aa37f0b301 100644
--- a/llvm/include/llvm/CodeGen/SpillPlacement.h
+++ b/llvm/include/llvm/CodeGen/SpillPlacement.h
@@ -36,6 +36,7 @@
namespace llvm {
class BitVector;
+class BlockUniformityProfile;
class EdgeBundles;
class MachineBlockFrequencyInfo;
class MachineFunction;
@@ -169,7 +170,8 @@ class SpillPlacement {
LLVM_ABI void releaseMemory();
void run(MachineFunction &MF, EdgeBundles *Bundles,
- MachineBlockFrequencyInfo *MBFI);
+ MachineBlockFrequencyInfo *MBFI,
+ const BlockUniformityProfile *Profile);
void activate(unsigned n);
void setThreshold(BlockFrequency Entry);
diff --git a/llvm/include/llvm/IR/FixedMetadataKinds.def b/llvm/include/llvm/IR/FixedMetadataKinds.def
index 552c1f60f2d38..6271633ca6d66 100644
--- a/llvm/include/llvm/IR/FixedMetadataKinds.def
+++ b/llvm/include/llvm/IR/FixedMetadataKinds.def
@@ -64,3 +64,4 @@ LLVM_FIXED_MD_KIND(MD_inline_history, "inline_history", 49)
LLVM_FIXED_MD_KIND(MD_elf_section_properties, "elf_section_properties", 50)
LLVM_FIXED_MD_KIND(MD_unique_id, "guid", 51)
LLVM_FIXED_MD_KIND(MD_mem_cache_hint, "mem.cache_hint", 52)
+LLVM_FIXED_MD_KIND(MD_block_uniformity_profile, "block.uniformity.profile", 53)
diff --git a/llvm/include/llvm/Passes/CodeGenPassBuilder.h b/llvm/include/llvm/Passes/CodeGenPassBuilder.h
index 9641ac7313c69..6afca2f6a8558 100644
--- a/llvm/include/llvm/Passes/CodeGenPassBuilder.h
+++ b/llvm/include/llvm/Passes/CodeGenPassBuilder.h
@@ -25,6 +25,7 @@
#include "llvm/Analysis/TypeBasedAliasAnalysis.h"
#include "llvm/CodeGen/AsmPrinter.h"
#include "llvm/CodeGen/AsmPrinterAnalysis.h"
+#include "llvm/CodeGen/BlockUniformityProfile.h"
#include "llvm/CodeGen/BranchFoldingPass.h"
#include "llvm/CodeGen/CodeGenPrepare.h"
#include "llvm/CodeGen/DeadMachineInstructionElim.h"
diff --git a/llvm/include/llvm/Passes/MachinePassRegistry.def b/llvm/include/llvm/Passes/MachinePassRegistry.def
index a297c08fac1b2..2a1f228f0675e 100644
--- a/llvm/include/llvm/Passes/MachinePassRegistry.def
+++ b/llvm/include/llvm/Passes/MachinePassRegistry.def
@@ -63,6 +63,8 @@ LOOP_PASS("loop-term-fold", LoopTermFoldPass())
// LiveVariables can be removed completely, and LiveIntervals can be directly
// computed. (We still either need to regenerate kill flags after regalloc, or
// preferably fix the scavenger to not depend on them).
+MACHINE_FUNCTION_ANALYSIS("block-uniformity-profile",
+ BlockUniformityProfileProxy())
MACHINE_FUNCTION_ANALYSIS("edge-bundles", EdgeBundlesAnalysis())
MACHINE_FUNCTION_ANALYSIS("gisel-cse-analysis", GISelCSEAnalysis(TM))
MACHINE_FUNCTION_ANALYSIS("gisel-value-tracking", GISelValueTrackingAnalysis())
@@ -140,6 +142,8 @@ MACHINE_FUNCTION_PASS("postra-machine-sink", PostRAMachineSinkingPass())
MACHINE_FUNCTION_PASS("postmisched", PostMachineSchedulerPass(TM))
MACHINE_FUNCTION_PASS("post-ra-pseudos", ExpandPostRAPseudosPass())
MACHINE_FUNCTION_PASS("print", PrintMIRPass())
+MACHINE_FUNCTION_PASS("print<block-uniformity-profile>",
+ BlockUniformityProfilePrinterPass(errs()))
MACHINE_FUNCTION_PASS("print<gisel-value-tracking>", GISelValueTrackingPrinterPass(errs()))
MACHINE_FUNCTION_PASS("print<livedebugvars>", LiveDebugVariablesPrinterPass(errs()))
MACHINE_FUNCTION_PASS("print<live-intervals>", LiveIntervalsPrinterPass(errs()))
diff --git a/llvm/lib/CodeGen/BlockUniformityProfile.cpp b/llvm/lib/CodeGen/BlockUniformityProfile.cpp
new file mode 100644
index 0000000000000..fff9843ce59eb
--- /dev/null
+++ b/llvm/lib/CodeGen/BlockUniformityProfile.cpp
@@ -0,0 +1,121 @@
+//===- BlockUniformityProfile.cpp - Block uniformity from PGO -----------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "llvm/CodeGen/BlockUniformityProfile.h"
+#include "llvm/CodeGen/MachineBasicBlock.h"
+#include "llvm/CodeGen/MachineFunction.h"
+#include "llvm/IR/BasicBlock.h"
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/Instruction.h"
+#include "llvm/IR/LLVMContext.h"
+#include "llvm/IR/Metadata.h"
+#include "llvm/Support/raw_ostream.h"
+#include <optional>
+
+using namespace llvm;
+
+static std::optional<bool> getIRBlockUniformity(const BasicBlock &BB) {
+ const Instruction *TI = BB.getTerminator();
+ if (!TI)
+ return std::nullopt;
+
+ MDNode *MD = TI->getMetadata(LLVMContext::MD_block_uniformity_profile);
+ if (!MD)
+ return std::nullopt;
+
+ // Metadata format: !{i1 IsUniform} - structural validity assumed (verifier).
+ // Returns true if uniform (not divergent).
+ return mdconst::extract<ConstantInt>(MD->getOperand(0))->isOne();
+}
+
+void BlockUniformityProfile::compute(const MachineFunction &MF) {
+ HasProfile = false;
+ NumBlockIDs = MF.getNumBlockIDs();
+ DivergentBlocks.clear();
+ DivergentBlocks.resize(NumBlockIDs);
+
+ // First determine whether any uniformity profile exists for this function.
+ for (const MachineBasicBlock &MBB : MF) {
+ const BasicBlock *BB = MBB.getBasicBlock();
+ if (!BB)
+ continue;
+ if (getIRBlockUniformity(*BB).has_value()) {
+ HasProfile = true;
+ break;
+ }
+ }
+
+ if (!HasProfile)
+ return;
+
+ // Conservative behavior: if profile exists for the function but we
+ // cannot classify a particular (Machine)basic block, treat it as divergent.
+ for (const MachineBasicBlock &MBB : MF) {
+ const unsigned Num = MBB.getNumber();
+ bool IsDivergent = true;
+ if (const BasicBlock *BB = MBB.getBasicBlock()) {
+ if (auto U = getIRBlockUniformity(*BB))
+ IsDivergent = !*U; // Metadata stores IsUniform, we want IsDivergent
+ }
+ if (Num < DivergentBlocks.size() && IsDivergent)
+ DivergentBlocks.set(Num);
+ }
+}
+
+void BlockUniformityProfile::print(raw_ostream &OS,
+ const MachineFunction &MF) const {
+ OS << "BlockUniformityProfile for function: ";
+ MF.getFunction().printAsOperand(OS, /*PrintType=*/false);
+ OS << '\n';
+ OS << "HasProfile: " << (HasProfile ? "true" : "false") << '\n';
+ if (!HasProfile)
+ return;
+
+ for (const MachineBasicBlock &MBB : MF) {
+ const BasicBlock *BB = MBB.getBasicBlock();
+ if (!BB)
+ continue;
+ OS << " " << printMBBReference(MBB);
+ if (BB->hasName())
+ OS << " (%" << BB->getName() << ")";
+ if (auto U = getIRBlockUniformity(*BB)) {
+ OS << ": " << (*U ? "uniform" : "divergent") << '\n';
+ continue;
+ }
+ OS << ": no PGO annotation (treated divergent for spill placement)\n";
+ }
+}
+
+bool BlockUniformityProfile::isDivergent(const MachineBasicBlock &MBB) const {
+ if (!HasProfile)
+ return false;
+ assert(MBB.getParent()->getNumBlockIDs() == NumBlockIDs &&
+ "MachineFunction was modified without invalidating "
+ "BlockUniformityProfile");
+ const unsigned Num = MBB.getNumber();
+ assert(Num < DivergentBlocks.size() && "Block number out of range");
+ return DivergentBlocks.test(Num);
+}
+
+AnalysisKey BlockUniformityProfileProxy::Key;
+
+BlockUniformityProfileProxy::Result
+BlockUniformityProfileProxy::run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &) {
+ BlockUniformityProfile Profile;
+ Profile.compute(MF);
+ return Profile;
+}
+
+PreservedAnalyses
+BlockUniformityProfilePrinterPass::run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &MFAM) {
+ auto &Profile = MFAM.getResult<BlockUniformityProfileProxy>(MF);
+ Profile.print(OS, MF);
+ return PreservedAnalyses::all();
+}
diff --git a/llvm/lib/CodeGen/CMakeLists.txt b/llvm/lib/CodeGen/CMakeLists.txt
index c572128b023c1..1305e12693fbd 100644
--- a/llvm/lib/CodeGen/CMakeLists.txt
+++ b/llvm/lib/CodeGen/CMakeLists.txt
@@ -36,6 +36,7 @@ add_llvm_component_library(LLVMCodeGen
BasicBlockPathCloning.cpp
BasicBlockSectionsProfileReader.cpp
BasicBlockMatchingAndInference.cpp
+ BlockUniformityProfile.cpp
CalcSpillWeights.cpp
CallingConvLower.cpp
CFGuardLongjmp.cpp
diff --git a/llvm/lib/CodeGen/SpillPlacement.cpp b/llvm/lib/CodeGen/SpillPlacement.cpp
index 55a96a22a00ec..4ea05ac2b2987 100644
--- a/llvm/lib/CodeGen/SpillPlacement.cpp
+++ b/llvm/lib/CodeGen/SpillPlacement.cpp
@@ -28,11 +28,14 @@
#include "llvm/CodeGen/SpillPlacement.h"
#include "llvm/ADT/BitVector.h"
+#include "llvm/CodeGen/BlockUniformityProfile.h"
#include "llvm/CodeGen/EdgeBundles.h"
#include "llvm/CodeGen/MachineBasicBlock.h"
#include "llvm/CodeGen/MachineBlockFrequencyInfo.h"
#include "llvm/CodeGen/MachineFunction.h"
#include "llvm/CodeGen/Passes.h"
+#include "llvm/CodeGen/TargetSubtargetInfo.h"
+#include "llvm/IR/Function.h"
#include "llvm/InitializePasses.h"
#include "llvm/Pass.h"
#include <algorithm>
@@ -193,7 +196,9 @@ bool SpillPlacementWrapperLegacy::runOnMachineFunction(MachineFunction &MF) {
auto *Bundles = &getAnalysis<EdgeBundlesWrapperLegacy>().getEdgeBundles();
auto *MBFI = &getAnalysis<MachineBlockFrequencyInfoWrapperPass>().getMBFI();
- Impl.run(MF, Bundles, MBFI);
+ BlockUniformityProfile Profile;
+ Profile.compute(MF);
+ Impl.run(MF, Bundles, MBFI, &Profile);
return false;
}
@@ -204,8 +209,9 @@ SpillPlacementAnalysis::run(MachineFunction &MF,
MachineFunctionAnalysisManager &MFAM) {
auto *Bundles = &MFAM.getResult<EdgeBundlesAnalysis>(MF);
auto *MBFI = &MFAM.getResult<MachineBlockFrequencyAnalysis>(MF);
+ auto &Profile = MFAM.getResult<BlockUniformityProfileProxy>(MF);
SpillPlacement Impl;
- Impl.run(MF, Bundles, MBFI);
+ Impl.run(MF, Bundles, MBFI, &Profile);
return Impl;
}
@@ -217,7 +223,8 @@ bool SpillPlacementAnalysis::Result::invalidate(
return true;
// Check dependencies.
return Inv.invalidate<EdgeBundlesAnalysis>(MF, PA) ||
- Inv.invalidate<MachineBlockFrequencyAnalysis>(MF, PA);
+ Inv.invalidate<MachineBlockFrequencyAnalysis>(MF, PA) ||
+ Inv.invalidate<BlockUniformityProfileProxy>(MF, PA);
}
SpillPlacement::SpillPlacement() = default;
@@ -230,7 +237,8 @@ void SpillPlacement::releaseMemory() {
}
void SpillPlacement::run(MachineFunction &mf, EdgeBundles *Bundles,
- MachineBlockFrequencyInfo *MBFI) {
+ MachineBlockFrequencyInfo *MBFI,
+ const BlockUniformityProfile *Profile) {
MF = &mf;
this->bundles = Bundles;
this->MBFI = MBFI;
@@ -240,12 +248,18 @@ void SpillPlacement::run(MachineFunction &mf, EdgeBundles *Bundles,
TodoList.clear();
TodoList.setUniverse(bundles->getNumBundles());
+ const bool HasProfile = Profile && Profile->hasProfile();
+
// Compute total ingoing and outgoing block frequencies for all bundles.
BlockFrequencies.resize(mf.getNumBlockIDs());
setThreshold(MBFI->getEntryFreq());
for (auto &I : mf) {
unsigned Num = I.getNumber();
- BlockFrequencies[Num] = MBFI->getBlockFreq(&I);
+ if (HasProfile && Profile->isDivergent(I)) {
+ BlockFrequencies[Num] = MBFI->getEntryFreq();
+ } else {
+ BlockFrequencies[Num] = MBFI->getBlockFreq(&I);
+ }
}
}
diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp
index f3110558f146d..330f3185877ea 100644
--- a/llvm/lib/Passes/PassBuilder.cpp
+++ b/llvm/lib/Passes/PassBuilder.cpp
@@ -81,6 +81,7 @@
#include "llvm/CodeGen/AssignmentTrackingAnalysis.h"
#include "llvm/CodeGen/AtomicExpand.h"
#include "llvm/CodeGen/BasicBlockSectionsProfileReader.h"
+#include "llvm/CodeGen/BlockUniformityProfile.h"
#include "llvm/CodeGen/BranchFoldingPass.h"
#include "llvm/CodeGen/BranchRelaxation.h"
#include "llvm/CodeGen/BreakFalseDeps.h"
diff --git a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
index ba0d5187f3c08..b30121a87d483 100644
--- a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
+++ b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
@@ -1206,6 +1206,9 @@ class PGOUseFunc {
// Annotate the irreducible loop header weights.
void annotateIrrLoopHeaderWeights();
+ // Annotate per-block uniformity info for offload profiling.
+ void setBlockUniformityAttribute();
+
// The hotness of the function from the profile count.
enum FuncFreqAttr { FFA_Normal, FFA_Cold, FFA_Hot };
@@ -1784,6 +1787,42 @@ void PGOUseFunc::annotateIrrLoopHeaderWeights() {
}
}
+void PGOUseFunc::setBlockUniformityAttribute() {
+ if (ProfileRecord.UniformityBits.empty())
+ return;
+
+ // Annotate uniformity on each instrumented IR basic block so later codegen
+ // passes (MachineFunction) can consume it without relying on fragile block
+ // numbering heuristics.
+ //
+ // Metadata kind: LLVMContext::MD_block_uniformity_profile
+ // Payload: i1 (true = uniform, false = divergent)
+
+ std::vector<BasicBlock *> InstrumentBBs;
+ FuncInfo.getInstrumentBBs(InstrumentBBs);
+
+ LLVMContext &Ctx = F.getContext();
+ Type *Int1Ty = Type::getInt1Ty(Ctx);
+
+ for (size_t I = 0, E = InstrumentBBs.size(); I < E; ++I) {
+ BasicBlock *BB = InstrumentBBs[I];
+ if (!BB || !BB->getTerminator())
+ continue;
+ bool IsUniform = ProfileRecord.isBlockUniform(I);
+ auto *MD = MDNode::get(
+ Ctx, ConstantAsMetadata::get(ConstantInt::get(Int1Ty, IsUniform)));
+ BB->getTerminator()->setMetadata(LLVMContext::MD_block_uniformity_profile,
+ MD);
+ }
+
+ LLVM_DEBUG({
+ dbgs() << "PGO: Set block uniformity profile for " << F.getName() << ": ";
+ for (size_t I = 0, E = InstrumentBBs.size(); I < E; ++I)
+ dbgs() << (ProfileRecord.isBlockUniform(I) ? 'U' : 'D');
+ dbgs() << "\n";
+ });
+}
+
void SelectInstVisitor::instrumentOneSelectInst(SelectInst &SI) {
Module *M = F.getParent();
IRBuilder<> Builder(&SI);
@@ -2295,6 +2334,7 @@ static bool annotateAllFunctions(
Func.setBranchWeights();
Func.annotateValueSites();
Func.annotateIrrLoopHeaderWeights();
+ Func.setBlockUniformityAttribute();
PGOUseFunc::FuncFreqAttr FreqAttr = Func.getFuncFreqAttr();
if (FreqAttr == PGOUseFunc::FFA_Cold)
ColdFunctions.push_back(&F);
diff --git a/llvm/test/CodeGen/AMDGPU/block-uniformity-profile.ll b/llvm/test/CodeGen/AMDGPU/block-uniformity-profile.ll
new file mode 100644
index 0000000000000..5a8cff00d84c7
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/block-uniformity-profile.ll
@@ -0,0 +1,79 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -O0 -stop-after=finalize-isel -o - %s | \
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -passes='print<block-uniformity-profile>' -x mir -filetype=null 2>&1 | FileCheck %s
+
+; Test that BlockUniformityProfileProxy correctly reads block-uniformity-profile
+; metadata from IR basic blocks and classifies machine blocks.
+;
+; This metadata is attached during PGO-use phase to indicate whether a basic block
+; was executed uniformly (all lanes together) or divergently (partial wave).
+;
+; The analysis is consumed by SpillPlacement to flatten block frequencies for
+; divergent blocks, preventing PGO from causing regressions on divergent code paths.
+
+; CHECK-LABEL: BlockUniformityProfile for function: @uniform_blocks
+; CHECK-NEXT: HasProfile: true
+; CHECK: %bb.{{[0-9]+}} (%entry): uniform
+define amdgpu_kernel void @uniform_blocks(ptr addrspace(1) %out) #0 {
+entry:
+ store i32 1, ptr addrspace(1) %out, align 4
+ ret void, !block-uniformity-profile !0
+}
+
+; CHECK-LABEL: BlockUniformityProfile for function: @divergent_blocks
+; CHECK-NEXT: HasProfile: true
+; CHECK-DAG: %bb.{{[0-9]+}} (%if.then): divergent
+; CHECK-DAG: %bb.{{[0-9]+}} (%if.else): uniform
+define amdgpu_kernel void @divergent_blocks(ptr addrspace(1) %out, i32 %tid) #0 {
+entry:
+ %cmp = icmp eq i32 %tid, 0
+ br i1 %cmp, label %if.then, label %if.else
+
+if.then:
+ store i32 1, ptr addrspace(1) %out, align 4
+ ret void, !block-uniformity-profile !1
+
+if.else:
+ store i32 2, ptr addrspace(1) %out, align 4
+ ret void, !block-uniformity-profile !0
+}
+
+; CHECK-LABEL: BlockUniformityProfile for function: @missing_metadata
+; CHECK-NEXT: HasProfile: true
+; CHECK-DAG: %bb.{{[0-9]+}} (%if.then): no PGO annotation (treated divergent for spill placement)
+; CHECK-DAG: %bb.{{[0-9]+}} (%if.else): uniform
+define amdgpu_kernel void @missing_metadata(ptr addrspace(1) %out, i32 %cond) #0 {
+entry:
+ %cmp = icmp sgt i32 %cond, 0
+ br i1 %cmp, label %if.then, label %if.else
+
+if.then:
+ store i32 1, ptr addrspace(1) %out, align 4
+ ret void
+
+if.else:
+ store i32 2, ptr addrspace(1) %out, align 4
+ ret void, !block-uniformity-profile !0
+}
+
+; CHECK-LABEL: BlockUniformityProfile for function: @no_divergence_metadata
+; CHECK-NEXT: HasProfile: false
+define amdgpu_kernel void @no_divergence_metadata(ptr addrspace(1) %out, i32 %cond) #0 {
+entry:
+ ; No uniformity metadata - analysis should report hasProfile() = false
+ %cmp = icmp sgt i32 %cond, 0
+ br i1 %cmp, label %if.then, label %if.else
+
+if.then:
+ store i32 1, ptr addrspace(1) %out, align 4
+ ret void
+
+if.else:
+ store i32 2, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
+
+; Metadata: i1 true = uniform, i1 false = divergent
+!0 = !{i1 true} ; uniform
+!1 = !{i1 false} ; divergent
More information about the cfe-commits
mailing list