[clang] [compiler-rt] [llvm] [PGO][AMDGPU] Add uniformity-aware offload profile format and instrumentation (PR #190708)
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Mon Jun 1 06:54:48 PDT 2026
https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/190708
>From 56edfebc8f80feacfcdae621a9816a4407754d3c 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] [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/FixedMetadataKinds.def | 1 +
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 | 71 +++++-
.../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 +-
53 files changed, 1076 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/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/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..b30121a87d483 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 {
@@ -1199,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 };
@@ -1305,11 +1315,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 +1341,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
@@ -1762,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);
@@ -2273,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);
@@ -2410,11 +2472,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,
More information about the cfe-commits
mailing list