From de40b4417b18d042a025b379a3385434ec976788 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Fri, 28 Aug 2026 19:50:27 +0300 Subject: [PATCH 1/4] tests: add reproducer for volatile global accesses reaching SPIR-V as plain loads A `volatile` load or store of a 32 or 64 bit value in global memory is emitted as OpLoad / OpStore with the Volatile memory operand, which IGC serves from L1 like any other access. CUDA compiles the same source to ld.volatile / st.volatile, which the PTX ISA defines as relaxed memory operations at system scope, and Kokkos::volatile_load relies on that in the UnorderedMap insert list walk: on Aurora PVC hip.UnorderedMap_insert counts 901 keys for 900 because a losing inserter reads a stale key. TestFixVolatileLoadLoweringSPIRV.bash compiles TestFixVolatileLoadLowering.hip with --save-temps and checks the lowered device bitcode (and the SPIR-V module when the Khronos translator produced it) for relaxed atomic forms of the 32 and 64 bit global accesses, and that 16 bit and work-group local ones stay as they are. TestFixKokkosUnorderedMapInsert.hip is the insert idiom itself (claim, publish with __threadfence, relaxed CAS, volatile list walk) checked on the host for duplicate keys and orphaned slots; it passes on the CPU OpenCL runtime before and after and is meant for PVC. --- tests/runtime/CMakeLists.txt | 6 + .../TestFixKokkosUnorderedMapInsert.hip | 299 ++++++++++++++++++ tests/runtime/TestFixVolatileLoadLowering.hip | 150 +++++++++ .../TestFixVolatileLoadLoweringSPIRV.bash | 126 ++++++++ 4 files changed, 581 insertions(+) create mode 100644 tests/runtime/TestFixKokkosUnorderedMapInsert.hip create mode 100644 tests/runtime/TestFixVolatileLoadLowering.hip create mode 100755 tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash diff --git a/tests/runtime/CMakeLists.txt b/tests/runtime/CMakeLists.txt index 6daa519ec..1d0a062d8 100644 --- a/tests/runtime/CMakeLists.txt +++ b/tests/runtime/CMakeLists.txt @@ -113,6 +113,12 @@ if(SPIRV_DIS AND (NOT USE_NEW_OFFLOAD_DRIVER OR CLANG_OFFLOAD_BUNDLER)) add_shell_test(TestBoolKernelParamSPIRV.bash) endif() +# Volatile 32 and 64 bit global accesses must reach the SPIR-V producer as +# relaxed device-scope atomics (Kokkos::volatile_load on PVC reads stale L1 +# data otherwise). Inspects the lowered device bitcode of +# TestFixVolatileLoadLowering.hip, and the SPIR-V module when spirv-dis exists. +add_shell_test(TestFixVolatileLoadLoweringSPIRV.bash) + add_shell_test(../run_testenvvars.sh) set_tests_properties(run_testenvvars PROPERTIES TIMEOUT 60) diff --git a/tests/runtime/TestFixKokkosUnorderedMapInsert.hip b/tests/runtime/TestFixKokkosUnorderedMapInsert.hip new file mode 100644 index 000000000..82d7da013 --- /dev/null +++ b/tests/runtime/TestFixKokkosUnorderedMapInsert.hip @@ -0,0 +1,299 @@ +// Reproduces the Kokkos::UnorderedMap::insert claim-then-link idiom that +// miscounts on Aurora PVC (Kokkos_ContainersUnitTest_HIP hip.UnorderedMap_insert +// reports 901 distinct keys for 900). Mirrors containers/src/Kokkos_UnorderedMap.hpp +// insert() and containers/src/Kokkos_Bitset.hpp (Kokkos develop 93b924764): +// +// claim : Bitset::set(i), a relaxed device-scope atomic fetch_or +// publish: keys[i] = key; memory_fence() (= __threadfence()); relaxed +// device-scope CAS of the slot index into the hash list head / link +// read : list walk through Kokkos::volatile_load of the link and of the key +// undo : Bitset::reset(i), a relaxed device-scope atomic fetch_and +// +// A losing inserter that reads a stale key from the winner's slot appends its +// own slot behind it, and the same key ends up twice in one list. Every run +// walks the lists on the host and fails on a duplicate key, an orphaned slot +// (bit set, never linked) or a bad count. The volatile loads are what the +// lowering in HipLowerVolatileAccessesPass turns into relaxed device-scope +// atomic loads; on the CPU OpenCL runtime this passes before and after. + +#include + +#include +#include +#include +#include +#include + +#define INVALID 0xFFFFFFFFu + +#define CHECK(Cmd) \ + do { \ + hipError_t Err = (Cmd); \ + if (Err != hipSuccess) { \ + printf("FAILED: %s at %s:%d\n", hipGetErrorString(Err), __FILE__, \ + __LINE__); \ + exit(1); \ + } \ + } while (0) + +static constexpr unsigned Reps = 40; +static constexpr unsigned NumInserts = 90000; +static constexpr unsigned NumKeys = 900; +static constexpr unsigned Dups = NumInserts / NumKeys; +static constexpr unsigned NumNodes = 100000; +static constexpr unsigned BlockSize = 256; + +enum { S_SUCCESS = 0, S_EXISTING = 1, S_FAILED = 2, S_FREE_FAIL = 3, S_COUNT = 4 }; + +struct Table { + unsigned *Bits; // Capacity / 32 words + unsigned *Heads; // NumLists + unsigned *Next; // Capacity + unsigned *Keys; // Capacity + unsigned *Stats; // S_COUNT +}; + +// Kokkos::pod_hash = MurmurHash3_x86_32(&key, 4, seed 0) +__host__ __device__ inline uint32_t rotl32(uint32_t X, int R) { + return (X << R) | (X >> (32 - R)); +} +__host__ __device__ inline uint32_t murmur3U32(uint32_t Key, uint32_t Seed) { + uint32_t H = Seed; + uint32_t K = Key * 0xcc9e2d51u; + K = rotl32(K, 15) * 0x1b873593u; + H ^= K; + H = rotl32(H, 13) * 5u + 0xe6546b64u; + H ^= 4u; + H ^= H >> 16; + H *= 0x85ebca6bu; + H ^= H >> 13; + H *= 0xc2b2ae35u; + H ^= H >> 16; + return H; +} + +// Kokkos::UnorderedMap sizing (Kokkos_UnorderedMap.hpp, Kokkos_UnorderedMap_impl.cpp) +static uint32_t calculateCapacity(uint32_t Hint) { + return ((static_cast(7ull * Hint / 6u) + 127u) / 128u) * 128u; +} +static uint32_t findHashSize(uint32_t Size) { + static const uint32_t Primes[] = { + 85703, 90749, 95783, 100823, 105871, 110909, 115963, 120997, + 126031, 141157, 151237, 161323, 171401, 181499, 191579, 201653}; + for (uint32_t P : Primes) + if (Size <= P) + return P; + return Primes[sizeof(Primes) / sizeof(Primes[0]) - 1]; +} + +// Kokkos::volatile_load +__device__ __forceinline__ unsigned vload(unsigned *P) { + return *reinterpret_cast(P); +} + +// Kokkos::Bitset::set: true if this thread flipped the bit 0 -> 1 +__device__ __forceinline__ bool bitSet(unsigned *Bits, unsigned I) { + unsigned Mask = 1u << (I & 31); + return !(__hip_atomic_fetch_or(&Bits[I >> 5], Mask, __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT) & + Mask); +} +// Kokkos::Bitset::reset: true if the bit was set +__device__ __forceinline__ bool bitReset(unsigned *Bits, unsigned I) { + unsigned Mask = 1u << (I & 31); + return __hip_atomic_fetch_and(&Bits[I >> 5], ~Mask, __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT) & + Mask; +} +// Kokkos::Bitset::find_any_unset_near, forward scan, forward hint move. +__device__ __forceinline__ bool findAnyUnsetNear(unsigned *Bits, + unsigned NumBlocks, + unsigned LastBlockMask, + unsigned *Hint) { + unsigned BlockIdx = *Hint >> 5; + unsigned Offset = *Hint & 31; + unsigned Block = __hip_atomic_load(&Bits[BlockIdx], __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT); + Block = (!LastBlockMask || BlockIdx < NumBlocks - 1) ? ~Block + : (~Block & LastBlockMask); + if (Block == 0u) { + unsigned NB = BlockIdx + 1; + if (NB >= NumBlocks) + NB = 0; + *Hint = NB * 32u + Offset; + return false; + } + unsigned Rot = Offset ? ((Block >> Offset) | (Block << (32 - Offset))) : Block; + *Hint = (BlockIdx << 5) + ((__builtin_ctz(Rot) + Offset) & 31); + return true; +} + +// Kokkos::UnorderedMap::insert, unbounded, key == value, NoOp insert op. +__device__ void mapInsert(Table T, unsigned Capacity, unsigned NumLists, + unsigned Key) { + const unsigned NumBlocks = (Capacity + 31) / 32; + const unsigned LastBlockMask = + (Capacity & 31) ? ((1u << (Capacity & 31)) - 1u) : 0u; + const unsigned HashList = murmur3U32(Key, 0) % NumLists; + unsigned *CurrPtr = &T.Heads[HashList]; + unsigned NewIndex = INVALID; + unsigned IndexHint = static_cast( + (static_cast(HashList) * Capacity) / NumLists); + unsigned FindAttempts = 0; + + bool NotDone = true; + while (NotDone) { + // Need volatile_load as other threads may be appending. + unsigned Curr = vload(CurrPtr); + while (Curr != INVALID && vload(&T.Keys[Curr]) != Key) { + IndexHint = Curr; + CurrPtr = &T.Next[Curr]; + Curr = vload(CurrPtr); + } + + if (Curr != INVALID) { + if (NewIndex != INVALID && !bitReset(T.Bits, NewIndex)) + atomicAdd(&T.Stats[S_FREE_FAIL], 1u); + atomicAdd(&T.Stats[S_EXISTING], 1u); + NotDone = false; + } else { + if (NewIndex == INVALID) { + bool Found = + findAnyUnsetNear(T.Bits, NumBlocks, LastBlockMask, &IndexHint); + if (!Found && ++FindAttempts >= NumBlocks) { + atomicAdd(&T.Stats[S_FAILED], 1u); + NotDone = false; + } else if (bitSet(T.Bits, IndexHint)) { + NewIndex = IndexHint; + T.Keys[NewIndex] = Key; // plain store, as in Kokkos + __threadfence(); // Kokkos::memory_fence() + } + } + unsigned Expected = INVALID; + if (NewIndex != INVALID) { + __hip_atomic_compare_exchange_strong(CurrPtr, &Expected, NewIndex, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT); + if (Curr == Expected) { + atomicAdd(&T.Stats[S_SUCCESS], 1u); + NotDone = false; + } + } + } + } +} + +__global__ void insertKernel(Table T, unsigned Capacity, unsigned NumLists, + int Near) { + unsigned I = blockIdx.x * blockDim.x + threadIdx.x; + if (I >= NumInserts) + return; + const unsigned Key = Near ? I / Dups : I % NumKeys; + mapInsert(T, Capacity, NumLists, Key); +} + +int main() { + const unsigned Capacity = calculateCapacity(NumNodes); + const unsigned NumLists = findHashSize(Capacity); + const unsigned NumBlocks = (Capacity + 31) / 32; + printf("inserts=%u keys=%u capacity=%u hash_lists=%u reps=%u\n", NumInserts, + NumKeys, Capacity, NumLists, Reps); + + Table T; + CHECK(hipMalloc(&T.Bits, NumBlocks * sizeof(unsigned))); + CHECK(hipMalloc(&T.Heads, NumLists * sizeof(unsigned))); + CHECK(hipMalloc(&T.Next, Capacity * sizeof(unsigned))); + CHECK(hipMalloc(&T.Keys, Capacity * sizeof(unsigned))); + CHECK(hipMalloc(&T.Stats, S_COUNT * sizeof(unsigned))); + + std::vector Bits(NumBlocks), Heads(NumLists), Next(Capacity), + Keys(Capacity), Stats(S_COUNT); + std::vector FirstSlot(NumKeys); + std::vector Reached(Capacity); + + unsigned BadRuns = 0, TotalSuccess = 0, TotalExisting = 0; + for (unsigned R = 0; R < Reps; ++R) { + const int Near = R & 1; // the Kokkos test alternates near / not near + CHECK(hipMemset(T.Bits, 0, NumBlocks * sizeof(unsigned))); + CHECK(hipMemset(T.Heads, 0xFF, NumLists * sizeof(unsigned))); + CHECK(hipMemset(T.Next, 0xFF, Capacity * sizeof(unsigned))); + CHECK(hipMemset(T.Keys, 0, Capacity * sizeof(unsigned))); + CHECK(hipMemset(T.Stats, 0, S_COUNT * sizeof(unsigned))); + insertKernel<<<(NumInserts + BlockSize - 1) / BlockSize, BlockSize>>>( + T, Capacity, NumLists, Near); + CHECK(hipGetLastError()); + CHECK(hipDeviceSynchronize()); + CHECK(hipMemcpy(Bits.data(), T.Bits, NumBlocks * sizeof(unsigned), + hipMemcpyDeviceToHost)); + CHECK(hipMemcpy(Heads.data(), T.Heads, NumLists * sizeof(unsigned), + hipMemcpyDeviceToHost)); + CHECK(hipMemcpy(Next.data(), T.Next, Capacity * sizeof(unsigned), + hipMemcpyDeviceToHost)); + CHECK(hipMemcpy(Keys.data(), T.Keys, Capacity * sizeof(unsigned), + hipMemcpyDeviceToHost)); + CHECK(hipMemcpy(Stats.data(), T.Stats, S_COUNT * sizeof(unsigned), + hipMemcpyDeviceToHost)); + + // Walk every list: count linked slots, duplicate keys, orphaned bits. + unsigned SetBits = 0, Linked = 0, DupKeys = 0, Orphans = 0, Broken = 0; + for (unsigned W = 0; W < NumBlocks; ++W) + SetBits += __builtin_popcount(Bits[W]); + std::fill(Reached.begin(), Reached.end(), 0); + std::fill(FirstSlot.begin(), FirstSlot.end(), INVALID); + for (unsigned L = 0; L < NumLists; ++L) { + unsigned Curr = Heads[L]; + unsigned Steps = 0; + while (Curr != INVALID) { + if (Curr >= Capacity || Reached[Curr] || ++Steps > Capacity || + Keys[Curr] >= NumKeys || murmur3U32(Keys[Curr], 0) % NumLists != L) { + Broken++; + break; + } + Reached[Curr] = 1; + Linked++; + unsigned K = Keys[Curr]; + if (FirstSlot[K] != INVALID) { + DupKeys++; + printf("rep %u: duplicate key %u in list %u: slots %u and %u\n", R, + K, L, FirstSlot[K], Curr); + } else { + FirstSlot[K] = Curr; + } + Curr = Next[Curr]; + } + } + for (unsigned I = 0; I < Capacity; ++I) + if (((Bits[I >> 5] >> (I & 31)) & 1u) && !Reached[I]) + Orphans++; + TotalSuccess += Stats[S_SUCCESS]; + TotalExisting += Stats[S_EXISTING]; + bool Ok = SetBits == NumKeys && Linked == NumKeys && DupKeys == 0 && + Orphans == 0 && Broken == 0 && Stats[S_FAILED] == 0 && + Stats[S_FREE_FAIL] == 0 && + Stats[S_SUCCESS] + Stats[S_EXISTING] == NumInserts; + if (!Ok) { + BadRuns++; + printf("rep %u near=%d: bits=%u linked=%u dup_keys=%u orphans=%u " + "broken=%u success=%u existing=%u failed=%u free_fail=%u\n", + R, Near, SetBits, Linked, DupKeys, Orphans, Broken, + Stats[S_SUCCESS], Stats[S_EXISTING], Stats[S_FAILED], + Stats[S_FREE_FAIL]); + } + } + + CHECK(hipFree(T.Bits)); + CHECK(hipFree(T.Heads)); + CHECK(hipFree(T.Next)); + CHECK(hipFree(T.Keys)); + CHECK(hipFree(T.Stats)); + + printf("%u runs: %u new keys, %u existing keys (expected %u and %u)\n", Reps, + TotalSuccess, TotalExisting, Reps * NumKeys, + Reps * (NumInserts - NumKeys)); + if (BadRuns) { + printf("FAILED: %u of %u runs miscounted\n", BadRuns, Reps); + return 1; + } + printf("PASSED\n"); + return 0; +} diff --git a/tests/runtime/TestFixVolatileLoadLowering.hip b/tests/runtime/TestFixVolatileLoadLowering.hip new file mode 100644 index 000000000..4d5393cf4 --- /dev/null +++ b/tests/runtime/TestFixVolatileLoadLowering.hip @@ -0,0 +1,150 @@ +// Reproducer: volatile loads and stores of 32 and 64 bit values in global +// memory reach the SPIR-V producer as plain `load volatile` / `store volatile` +// and end up as OpLoad / OpStore with the Volatile memory operand, which IGC +// treats as ordinary cached accesses. CUDA compiles the same source to +// ld.volatile / st.volatile, which the PTX ISA (8.4.2 "volatile Operation") +// defines as relaxed memory operations at system scope, and code written +// against that (Kokkos::volatile_load in the UnorderedMap insert list walk) +// reads stale L1 data on PVC. chipStar must lower them to relaxed +// device-scope atomic loads and stores, which the SPIR-V producers emit as +// OpAtomicLoad / OpAtomicStore. +// +// This executable checks that the accesses still compute the right values. +// TestFixVolatileLoadLoweringSPIRV.bash compiles this file with --save-temps and +// inspects the lowered device bitcode (and the SPIR-V module) for the atomic +// forms; that is the check that fails without the lowering. + +#include + +#include +#include +#include + +static constexpr int NumBlocks = 4; +static constexpr int BlockSize = 256; +static constexpr int N = NumBlocks * BlockSize; + +#define CHECK(Cmd) \ + do { \ + hipError_t Err = (Cmd); \ + if (Err != hipSuccess) { \ + printf("FAILED: %s at %s:%d\n", hipGetErrorString(Err), __FILE__, \ + __LINE__); \ + exit(1); \ + } \ + } while (0) + +// The accesses the lowering must rewrite: 32 and 64 bit volatile loads and +// stores through global memory pointers. +__global__ void volatileAccess(const unsigned *In32, + const unsigned long long *In64, unsigned *Out32, + unsigned long long *Out64) { + int Tid = blockIdx.x * blockDim.x + threadIdx.x; + const volatile unsigned *VIn32 = In32; + const volatile unsigned long long *VIn64 = In64; + volatile unsigned *VOut32 = Out32; + volatile unsigned long long *VOut64 = Out64; + unsigned A = VIn32[Tid]; + unsigned long long B = VIn64[Tid]; + VOut32[Tid] = A + 1u; + VOut64[Tid] = B + 1ull; +} + +// Accesses the lowering must leave as they are: 16 bit values (OpenCL SPIR-V +// consumers implement 32 and 64 bit atomics only) and work-group local memory +// (never shared across cores, so no L1 staleness to defeat). +__global__ void volatileLeftAlone(const unsigned short *In16, + unsigned short *Out16, unsigned *OutLocal) { + __shared__ unsigned Scratch[BlockSize]; + int Tid = threadIdx.x; + const volatile unsigned short *VIn16 = In16; + volatile unsigned short *VOut16 = Out16; + volatile unsigned *VScratch = Scratch; + VOut16[Tid] = VIn16[Tid] + 1; + VScratch[Tid] = Tid; + __syncthreads(); + OutLocal[Tid] = VScratch[(Tid + 1) % BlockSize]; +} + +int main() { + std::vector HIn32(N), HOut32(N), HOutLocal(BlockSize); + std::vector HIn64(N), HOut64(N); + std::vector HIn16(BlockSize), HOut16(BlockSize); + for (int I = 0; I < N; ++I) { + HIn32[I] = 0x10000u + I; + HIn64[I] = 0x100000000ull + I; + } + for (int I = 0; I < BlockSize; ++I) + HIn16[I] = 0x100 + I; + + unsigned *In32, *Out32, *OutLocal; + unsigned long long *In64, *Out64; + unsigned short *In16, *Out16; + CHECK(hipMalloc(&In32, N * sizeof(unsigned))); + CHECK(hipMalloc(&Out32, N * sizeof(unsigned))); + CHECK(hipMalloc(&In64, N * sizeof(unsigned long long))); + CHECK(hipMalloc(&Out64, N * sizeof(unsigned long long))); + CHECK(hipMalloc(&In16, BlockSize * sizeof(unsigned short))); + CHECK(hipMalloc(&Out16, BlockSize * sizeof(unsigned short))); + CHECK(hipMalloc(&OutLocal, BlockSize * sizeof(unsigned))); + CHECK(hipMemcpy(In32, HIn32.data(), N * sizeof(unsigned), + hipMemcpyHostToDevice)); + CHECK(hipMemcpy(In64, HIn64.data(), N * sizeof(unsigned long long), + hipMemcpyHostToDevice)); + CHECK(hipMemcpy(In16, HIn16.data(), BlockSize * sizeof(unsigned short), + hipMemcpyHostToDevice)); + + volatileAccess<<>>(In32, In64, Out32, Out64); + CHECK(hipGetLastError()); + volatileLeftAlone<<<1, BlockSize>>>(In16, Out16, OutLocal); + CHECK(hipGetLastError()); + CHECK(hipDeviceSynchronize()); + + CHECK(hipMemcpy(HOut32.data(), Out32, N * sizeof(unsigned), + hipMemcpyDeviceToHost)); + CHECK(hipMemcpy(HOut64.data(), Out64, N * sizeof(unsigned long long), + hipMemcpyDeviceToHost)); + CHECK(hipMemcpy(HOut16.data(), Out16, BlockSize * sizeof(unsigned short), + hipMemcpyDeviceToHost)); + CHECK(hipMemcpy(HOutLocal.data(), OutLocal, BlockSize * sizeof(unsigned), + hipMemcpyDeviceToHost)); + + int Errors = 0; + for (int I = 0; I < N; ++I) { + if (HOut32[I] != HIn32[I] + 1u) { + if (Errors++ < 5) + printf("Out32[%d] = %u, expected %u\n", I, HOut32[I], HIn32[I] + 1u); + } + if (HOut64[I] != HIn64[I] + 1ull) { + if (Errors++ < 5) + printf("Out64[%d] = %llu, expected %llu\n", I, HOut64[I], + HIn64[I] + 1ull); + } + } + for (int I = 0; I < BlockSize; ++I) { + if (HOut16[I] != (unsigned short)(HIn16[I] + 1)) { + if (Errors++ < 5) + printf("Out16[%d] = %u, expected %u\n", I, HOut16[I], HIn16[I] + 1); + } + if (HOutLocal[I] != (unsigned)((I + 1) % BlockSize)) { + if (Errors++ < 5) + printf("OutLocal[%d] = %u, expected %u\n", I, HOutLocal[I], + (I + 1) % BlockSize); + } + } + + CHECK(hipFree(In32)); + CHECK(hipFree(Out32)); + CHECK(hipFree(In64)); + CHECK(hipFree(Out64)); + CHECK(hipFree(In16)); + CHECK(hipFree(Out16)); + CHECK(hipFree(OutLocal)); + + if (Errors) { + printf("FAILED: %d mismatches\n", Errors); + return 1; + } + printf("PASSED\n"); + return 0; +} diff --git a/tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash b/tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash new file mode 100755 index 000000000..4810880da --- /dev/null +++ b/tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash @@ -0,0 +1,126 @@ +#!/bin/bash +# Volatile loads and stores of 32 and 64 bit values in global memory must reach +# the SPIR-V producer as relaxed atomic accesses, not as plain volatile ones. +# +# CUDA lowers a volatile global access to ld.volatile / st.volatile, which the +# PTX ISA (8.4.2 "volatile Operation") defines as a relaxed memory operation at +# system scope; code written against that (Kokkos::volatile_load in the +# UnorderedMap insert list walk) relies on the load bypassing a core's L1. In +# SPIR-V a `load volatile` is an OpLoad with the Volatile memory operand, which +# IGC serves from L1 like any other load, so on PVC the walk reads stale data. +# The fix lowers such accesses to `load atomic ... syncscope("device") +# monotonic`, which the SPIR-V producers emit as OpAtomicLoad with Relaxed +# semantics at Device scope. +# +# Compiles TestFixVolatileLoadLowering.hip with --save-temps and inspects the +# lowered device bitcode, the SPIR-V producer's input: the 32 and 64 bit global +# accesses must be atomic, the 16 bit and work-group local ones must not. When +# the module was produced by the Khronos translator and spirv-dis is available, +# the SPIR-V module is checked for OpAtomicLoad / OpAtomicStore as well. +set -u + +HIPCC="@CMAKE_BINARY_DIR@/bin/hipcc" +LLVM_DIS="@CLANG_ROOT_PATH_BIN@/llvm-dis" +SPIRV_DIS="@SPIRV_DIS@" +SRC="@CMAKE_CURRENT_SOURCE_DIR@/TestFixVolatileLoadLowering.hip" +OUT="@CMAKE_CURRENT_BINARY_DIR@/@TEST_NAME@.d" + +if [ ! -x "${LLVM_DIS}" ]; then + echo "HIP_SKIP_THIS_TEST: llvm-dis not found at ${LLVM_DIS}" + exit 0 +fi + +rm -rf "${OUT}" +mkdir -p "${OUT}" +cd "${OUT}" + +# -O2 so the accesses are not hidden behind allocas; --save-temps keeps the +# lowered device bitcode (*-lower.bc with either offload driver) and, with the +# old driver, the SPIR-V module (*.out). +"${HIPCC}" -O2 --save-temps=cwd -c "${SRC}" -o TestFixVolatileLoadLowering.o +BC=$(ls "${OUT}"/*-lower.bc 2>/dev/null | head -1) +if [ -z "${BC}" ]; then + echo "FAIL: no lowered device bitcode (*-lower.bc) produced by hipcc" + exit 1 +fi +"${LLVM_DIS}" "${BC}" -o lowered.ll + +# The body of one kernel, from its define line to the closing brace. +kernel_body() { + sed -n "/^define .*@_Z[0-9]*$1/,/^}/p" lowered.ll +} + +STATUS=0 +fail() { + echo "FAIL: $1" + STATUS=1 +} + +ACCESS=$(kernel_body volatileAccess) +if [ -z "${ACCESS}" ]; then + echo "FAIL: kernel volatileAccess not found in lowered.ll" + exit 1 +fi +for PATTERN in 'load atomic (volatile )?i32,' 'load atomic (volatile )?i64,' \ + 'store atomic (volatile )?i32 ' 'store atomic (volatile )?i64 '; do + if ! echo "${ACCESS}" | grep -qE "${PATTERN}"; then + fail "volatileAccess has no '${PATTERN}' after the pass pipeline" + fi +done +PLAIN=$(echo "${ACCESS}" | grep -E 'load volatile (i32|i64),|store volatile (i32|i64) ' || true) +if [ -n "${PLAIN}" ]; then + fail "volatileAccess still has non-atomic volatile global accesses:" + echo "${PLAIN}" +fi + +LEFT=$(kernel_body volatileLeftAlone) +if [ -z "${LEFT}" ]; then + echo "FAIL: kernel volatileLeftAlone not found in lowered.ll" + exit 1 +fi +if ! echo "${LEFT}" | grep -qE 'load volatile i16,'; then + fail "the 16 bit volatile load in volatileLeftAlone was not left alone" +fi +# The __shared__ array is accessed through a generic pointer over an +# addrspace(3) object; the accesses stay volatile and non-atomic either way. +if ! echo "${LEFT}" | grep -qE '(load|store) volatile i32'; then + fail "the work-group local volatile accesses in volatileLeftAlone were not left alone" +fi +if echo "${LEFT}" | grep -qE 'atomic'; then + fail "volatileLeftAlone gained atomic accesses:" + echo "${LEFT}" | grep -E 'atomic' +fi + +# SPIR-V level check. The LLVM SPIR-V backend selects every load as OpLoad +# regardless of its atomic ordering (SPIRVInstructionSelector::selectLoad), so +# the module check is only meaningful for the Khronos translator. +SPV=$(ls "${OUT}"/*.out 2>/dev/null | head -1) +if [ -n "${SPV}" ] && [ -n "${SPIRV_DIS}" ] && [ -x "${SPIRV_DIS}" ]; then + "${SPIRV_DIS}" "${SPV}" > module.spvasm + if grep -q "Generator: Khronos LLVM/SPIR-V Translator" module.spvasm; then + KID=$(grep -E 'OpEntryPoint Kernel %[0-9]+ "_Z[0-9]+volatileAccess' module.spvasm | + sed -E 's/.*Kernel (%[0-9]+) .*/\1/') + FUNC=$(sed -n "/^ *${KID} = OpFunction /,/OpFunctionEnd/p" module.spvasm) + LOADS=$(echo "${FUNC}" | grep -c 'OpAtomicLoad' || true) + STORES=$(echo "${FUNC}" | grep -c 'OpAtomicStore' || true) + if [ "${LOADS}" -lt 2 ] || [ "${STORES}" -lt 2 ]; then + fail "SPIR-V volatileAccess has ${LOADS} OpAtomicLoad and ${STORES} OpAtomicStore, expected at least 2 each" + fi + if echo "${FUNC}" | grep -qE 'Op(Load|Store) .*Volatile'; then + fail "SPIR-V volatileAccess still has Volatile OpLoad / OpStore:" + echo "${FUNC}" | grep -E 'Op(Load|Store) .*Volatile' + fi + echo "SPIR-V module checked (Khronos translator)" + else + echo "NOTE: SPIR-V module not produced by the Khronos translator; module check skipped" + grep -m1 "Generator" module.spvasm || true + fi +else + echo "NOTE: no SPIR-V module or spirv-dis; module check skipped" +fi + +if [ "${STATUS}" -ne 0 ]; then + echo "See ${OUT}/lowered.ll" + exit 1 +fi +echo "PASSED" From f42845d068f35b1161252ddfdc88cce8bebfbb3e Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Fri, 28 Aug 2026 19:55:23 +0300 Subject: [PATCH 2/4] llvm_passes: lower volatile global accesses to relaxed device-scope atomics CUDA compiles a volatile access to global memory into PTX ld.volatile / st.volatile, which the PTX ISA (8.4.2 "volatile Operation") defines as "equivalent to a relaxed memory operation with system-scope": a strong access that bypasses the core's L1 and observes other cores' relaxed atomics. HIP code relies on that; Kokkos::volatile_load walks the UnorderedMap insert list with it after the writer's __threadfence() plus relaxed CAS. SPIR-V's OpLoad with the Volatile memory operand only means the access cannot be eliminated, duplicated or combined, and IGC serves it from L1, so on Aurora PVC hip.UnorderedMap_insert reports 901 keys for 900. HipLowerVolatileAccessesPass rewrites every volatile load and store of a 32 or 64 bit integer, float or pointer through a global or generic pointer into `load atomic volatile ... syncscope("device") monotonic` (floats and pointers through the same-width integer), which both SPIR-V producers emit as OpAtomicLoad / OpAtomicStore with Relaxed semantics at Device scope. The explicit "device" syncscope matters: the translator at LLVM 17 hardcodes Device for atomic loads, newer ones map the default scope to CrossDevice. 8 and 16 bit values, private, constant and work-group local objects, under-aligned accesses and accesses that are already atomic are left as they are. --- llvm_passes/CMakeLists.txt | 1 + llvm_passes/HipLowerVolatileAccesses.cpp | 211 +++++++++++++++++++++++ llvm_passes/HipLowerVolatileAccesses.h | 32 ++++ llvm_passes/HipPasses.cpp | 15 ++ 4 files changed, 259 insertions(+) create mode 100644 llvm_passes/HipLowerVolatileAccesses.cpp create mode 100644 llvm_passes/HipLowerVolatileAccesses.h diff --git a/llvm_passes/CMakeLists.txt b/llvm_passes/CMakeLists.txt index 4aa4657c7..604ed0c01 100644 --- a/llvm_passes/CMakeLists.txt +++ b/llvm_passes/CMakeLists.txt @@ -120,6 +120,7 @@ add_library(LLVMHipPasses MODULE HipLowerOverflowIntrinsics.cpp HipLowerRoundIntrinsics.cpp HipLowerSwitch.cpp + HipLowerVolatileAccesses.cpp HipLowerZeroLengthArrays.cpp HipPasses.cpp HipPrintf.cpp diff --git a/llvm_passes/HipLowerVolatileAccesses.cpp b/llvm_passes/HipLowerVolatileAccesses.cpp new file mode 100644 index 000000000..84d318315 --- /dev/null +++ b/llvm_passes/HipLowerVolatileAccesses.cpp @@ -0,0 +1,211 @@ +//===- HipLowerVolatileAccesses.cpp ---------------------------------------===// +// +// Part of the chipStar Project, under the Apache License v2.0 with LLVM +// Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// A `volatile` access in HIP source means what it means in CUDA: clang lowers +// it to PTX ld.volatile / st.volatile, whose semantics the PTX ISA (8.4.2, +// "volatile Operation") defines as "equivalent to a relaxed memory operation +// with system-scope", a strong access that bypasses the core's L1 and observes +// other cores' relaxed atomics. HIP code written for that (Kokkos' +// volatile_load in the UnorderedMap insert list walk) polls memory another +// work-item publishes with __threadfence() followed by a relaxed CAS. +// +// SPIR-V has no such access. OpLoad with the Volatile memory operand only says +// the access "cannot be eliminated, duplicated, or combined with other +// accesses", and IGC serves it from L1 like any other load, so the poll reads +// stale data. The SPIR-V operation with the intended semantics is OpAtomicLoad +// / OpAtomicStore with Relaxed semantics at Device scope, which is what both +// SPIR-V producers emit for `load atomic ... syncscope("device") monotonic`. +// +// So every volatile load or store of a 32 or 64 bit integer, float or pointer +// through a global (addrspace 1) or generic (addrspace 4) pointer becomes the +// corresponding monotonic device-scope atomic. It stays volatile so nothing +// downstream merges or drops it. Floats and pointers go through the integer of +// the same width with a bitcast / ptrtoint / inttoptr around the access: +// OpAtomicLoad's result type must be an integer or float scalar, and the +// integer forms are the ones every OpenCL SPIR-V consumer implements (64 bit +// ones under Int64Atomics). +// +// Left as they are, and why: +// - 8 and 16 bit values, which OpenCL SPIR-V consumers have no atomics for +// (see HipLowerSubwordAtomics.cpp), and wider or vector values, which have +// no atomic form at all. +// - accesses whose pointer comes from a private (addrspace 0), constant +// (addrspace 2) or work-group local (addrspace 3) object: private memory +// is never shared, local memory never crosses a core, and atomics on the +// Function storage class have no defined behaviour. +// - under-aligned accesses: LLVM requires atomic loads and stores to be +// naturally aligned. Clang never emits a volatile access that is not, so +// these are reported and left volatile. +// - accesses that are already atomic. +// +// (c) 2026 chipStar developers +//===----------------------------------------------------------------------===// + +#include "HipLowerVolatileAccesses.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "PassPluginCompat.h" + +#define PASS_NAME "hip-lower-volatile-accesses" +#define DEBUG_TYPE PASS_NAME + +using namespace llvm; + +namespace { + +// SPIR-V address spaces as clang numbers them for spirv64. +constexpr unsigned PrivateAS = 0; +constexpr unsigned GlobalAS = 1; +constexpr unsigned ConstantAS = 2; +constexpr unsigned LocalAS = 3; +constexpr unsigned GenericAS = 4; + +/// Whether the access may reach memory another core can also access. Kernel +/// argument pointers arrive as generic pointers (clang launders them through +/// ptrtoint / inttoptr, which InferAddressSpaces does not see through), so a +/// generic pointer is assumed to be global unless the object it is derived +/// from says otherwise. +bool mayBeShared(Value *Ptr) { + unsigned AS = Ptr->getType()->getPointerAddressSpace(); + if (AS != GlobalAS && AS != GenericAS) + return false; + unsigned ObjAS = + getUnderlyingObject(Ptr)->getType()->getPointerAddressSpace(); + return ObjAS != PrivateAS && ObjAS != ConstantAS && ObjAS != LocalAS; +} + +/// The integer the access is performed on, or null when the value has no +/// 32 or 64 bit atomic form. +Type *atomicIntType(Type *ValTy, const DataLayout &DL) { + if (!ValTy->isIntegerTy() && !ValTy->isFloatingPointTy() && + !ValTy->isPointerTy()) + return nullptr; + uint64_t Bits = DL.getTypeStoreSizeInBits(ValTy); + if (Bits != 32 && Bits != 64) + return nullptr; + if (ValTy->isIntegerTy() && ValTy->getIntegerBitWidth() != Bits) + return nullptr; // i33 and friends round up to 64. + return Type::getIntNTy(ValTy->getContext(), Bits); +} + +bool isNaturallyAligned(Align A, Type *IntTy) { + return A.value() * 8 >= IntTy->getIntegerBitWidth(); +} + +void lowerLoad(LoadInst *LI, Type *IntTy, SyncScope::ID SSID) { + if (LI->getType() == IntTy) { + LI->setAtomic(AtomicOrdering::Monotonic, SSID); + return; + } + IRBuilder<> B(LI); + LoadInst *Bits = B.CreateAlignedLoad(IntTy, LI->getPointerOperand(), + LI->getAlign(), /*isVolatile=*/true, + LI->getName() + ".bits"); + Bits->setAtomic(AtomicOrdering::Monotonic, SSID); + Bits->setDebugLoc(LI->getDebugLoc()); + Value *V = LI->getType()->isPointerTy() ? B.CreateIntToPtr(Bits, LI->getType()) + : B.CreateBitCast(Bits, LI->getType()); + V->takeName(LI); + LI->replaceAllUsesWith(V); + LI->eraseFromParent(); +} + +void lowerStore(StoreInst *SI, Type *IntTy, SyncScope::ID SSID) { + Value *V = SI->getValueOperand(); + if (V->getType() == IntTy) { + SI->setAtomic(AtomicOrdering::Monotonic, SSID); + return; + } + IRBuilder<> B(SI); + Value *Bits = V->getType()->isPointerTy() ? B.CreatePtrToInt(V, IntTy) + : B.CreateBitCast(V, IntTy); + StoreInst *Raw = B.CreateAlignedStore(Bits, SI->getPointerOperand(), + SI->getAlign(), /*isVolatile=*/true); + Raw->setAtomic(AtomicOrdering::Monotonic, SSID); + Raw->setDebugLoc(SI->getDebugLoc()); + SI->eraseFromParent(); +} + +bool lowerVolatileAccesses(Function &F) { + const DataLayout &DL = F.getParent()->getDataLayout(); + SyncScope::ID DeviceSSID = F.getContext().getOrInsertSyncScopeID("device"); + SmallVector, 16> WorkList; + for (auto &BB : F) + for (auto &I : BB) { + Value *Ptr = nullptr; + Type *ValTy = nullptr; + Align A; + if (auto *LI = dyn_cast(&I)) { + if (!LI->isVolatile() || LI->isAtomic()) + continue; + Ptr = LI->getPointerOperand(); + ValTy = LI->getType(); + A = LI->getAlign(); + } else if (auto *SI = dyn_cast(&I)) { + if (!SI->isVolatile() || SI->isAtomic()) + continue; + Ptr = SI->getPointerOperand(); + ValTy = SI->getValueOperand()->getType(); + A = SI->getAlign(); + } else { + continue; + } + Type *IntTy = atomicIntType(ValTy, DL); + if (!IntTy || !mayBeShared(Ptr)) + continue; + if (!isNaturallyAligned(A, IntTy)) { + errs() << "warning: HipLowerVolatileAccesses: leaving under-aligned " + "volatile access alone in " + << F.getName() << ": " << I << "\n"; + continue; + } + WorkList.emplace_back(&I, IntTy); + } + + for (auto &[I, IntTy] : WorkList) { + if (auto *LI = dyn_cast(I)) + lowerLoad(LI, IntTy, DeviceSSID); + else + lowerStore(cast(I), IntTy, DeviceSSID); + } + return !WorkList.empty(); +} + +} // namespace + +PreservedAnalyses HipLowerVolatileAccessesPass::run(Function &F, + FunctionAnalysisManager &AM) { + return lowerVolatileAccesses(F) ? PreservedAnalyses::none() + : PreservedAnalyses::all(); +} + +#ifndef CHIP_COMBINED_PASS_PLUGIN +extern "C" ::llvm::PassPluginLibraryInfo LLVM_ATTRIBUTE_WEAK +llvmGetPassPluginInfo() { + return {LLVM_PLUGIN_API_VERSION, PASS_NAME, LLVM_VERSION_STRING, + [](PassBuilder &PB) { + PB.registerPipelineParsingCallback( + [](StringRef Name, FunctionPassManager &FPM, + ArrayRef) { + if (Name == PASS_NAME) { + FPM.addPass(HipLowerVolatileAccessesPass()); + return true; + } + return false; + }); + }}; +} +#endif // CHIP_COMBINED_PASS_PLUGIN diff --git a/llvm_passes/HipLowerVolatileAccesses.h b/llvm_passes/HipLowerVolatileAccesses.h new file mode 100644 index 000000000..ca32c5645 --- /dev/null +++ b/llvm_passes/HipLowerVolatileAccesses.h @@ -0,0 +1,32 @@ +//===- HipLowerVolatileAccesses.h -----------------------------------------===// +// +// Part of the chipStar Project, under the Apache License v2.0 with LLVM +// Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// Rewrites volatile 32 and 64 bit loads and stores through global and generic +// pointers into relaxed device-scope atomic ones, because that is the SPIR-V +// access with the semantics CUDA gives a volatile global access (PTX +// ld.volatile / st.volatile), while OpLoad / OpStore with the Volatile memory +// operand is served from L1 like any other access. +// +// (c) 2026 chipStar developers +//===----------------------------------------------------------------------===// + +#ifndef LLVM_PASSES_HIP_LOWER_VOLATILE_ACCESSES_H +#define LLVM_PASSES_HIP_LOWER_VOLATILE_ACCESSES_H + +#include + +using namespace llvm; + +class HipLowerVolatileAccessesPass + : public PassInfoMixin { +public: + PreservedAnalyses run(Function &F, FunctionAnalysisManager &AM); + static bool isRequired() { return true; } +}; + +#endif diff --git a/llvm_passes/HipPasses.cpp b/llvm_passes/HipPasses.cpp index 3dede76ad..1ef544628 100644 --- a/llvm_passes/HipPasses.cpp +++ b/llvm_passes/HipPasses.cpp @@ -32,6 +32,7 @@ #include "HipLowerMemset.h" #include "HipLowerFPAtomicMinMax.h" #include "HipLowerRoundIntrinsics.h" +#include "HipLowerVolatileAccesses.h" #include "HipIGBADetector.h" #include "HipPromoteInts.h" #include "HipLowerOverflowIntrinsics.h" @@ -210,6 +211,13 @@ static void addFullLinkTimePasses(ModulePassManager &MPM) { addPassWithVerification(MPM, HipIGBADetectorPass(), "HipIGBADetectorPass"); + // A volatile global access carries CUDA's ld.volatile / st.volatile meaning + // (a relaxed system-scope access that bypasses L1) and SPIR-V's Volatile + // memory operand does not, so rewrite them into relaxed device-scope + // atomics. Runs after the IGBA detector, which sees pointer-typed volatile + // loads as pointer loads before this launders them through an integer. + addPassWithVerification(MPM, createModuleToFunctionPassAdaptor(HipLowerVolatileAccessesPass()), "HipLowerVolatileAccessesPass"); + // Fix InvalidBitWidth errors due to non-standard integer types addPassWithVerification(MPM, HipPromoteIntsPass(), "HipPromoteIntsPass"); @@ -273,6 +281,13 @@ llvmGetPassPluginInfo() { MPM.addPass(HipLowerOverflowIntrinsicsPass()); return true; } + // Register the volatile access lowering as standalone, + // which makes it directly testable with opt. + if (Name == "hip-lower-volatile-accesses") { + MPM.addPass(createModuleToFunctionPassAdaptor( + HipLowerVolatileAccessesPass())); + return true; + } // Register SPIR-V function reorder pass as standalone if (Name == "hip-spirv-function-reorder") { MPM.addPass(HipSpirvFunctionReorderPass()); From 87b47eda842700cb2741a94cf53f8a44facbd0c3 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Fri, 28 Aug 2026 19:57:47 +0300 Subject: [PATCH 3/4] tests: cover every rule of the volatile access lowering with an opt fixture volatile-accesses.ll holds one kernel of accesses HipLowerVolatileAccessesPass must rewrite (32 and 64 bit integers, floats and pointers through global and generic pointers) and one of accesses it must leave alone (8 and 16 bit, vectors, work-group local, private, constant, under-aligned, already atomic). The runner drives the standalone hip-lower-volatile-accesses pass with opt, counts both sets, expects the under-aligned warnings, and translates the result to SPIR-V through spirv-val. --- tests/compiler/CMakeLists.txt | 3 + .../compiler/volatileAccesses/CMakeLists.txt | 58 ++++++++++ .../run_volatile_accesses_pass.bash | 102 ++++++++++++++++++ .../volatileAccesses/volatile-accesses.ll | 92 ++++++++++++++++ 4 files changed, 255 insertions(+) create mode 100644 tests/compiler/volatileAccesses/CMakeLists.txt create mode 100755 tests/compiler/volatileAccesses/run_volatile_accesses_pass.bash create mode 100644 tests/compiler/volatileAccesses/volatile-accesses.ll diff --git a/tests/compiler/CMakeLists.txt b/tests/compiler/CMakeLists.txt index c7dd72585..a42896d00 100644 --- a/tests/compiler/CMakeLists.txt +++ b/tests/compiler/CMakeLists.txt @@ -197,3 +197,6 @@ add_subdirectory(promoteInt) # Add the IR verification tests add_subdirectory(irVerification) + +# Add the volatile global access lowering pass tests +add_subdirectory(volatileAccesses) diff --git a/tests/compiler/volatileAccesses/CMakeLists.txt b/tests/compiler/volatileAccesses/CMakeLists.txt new file mode 100644 index 000000000..ff8588f81 --- /dev/null +++ b/tests/compiler/volatileAccesses/CMakeLists.txt @@ -0,0 +1,58 @@ +#============================================================================= +# Copyright (c) 2026 chipStar developers +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. +# +#============================================================================= + +configure_file(run_volatile_accesses_pass.bash + ${CMAKE_CURRENT_BINARY_DIR}/run_volatile_accesses_pass.bash @ONLY) + +# LLVM_SPIRV carries the sentinel "NOT_NEEDED" when chipStar targets LLVM's +# integrated SPIR-V backend; this test still translates to SPIR-V to validate +# the pass output, so fall back to a translator next to the other LLVM tools. +set(VOLATILE_ACCESSES_LLVM_SPIRV "${LLVM_SPIRV}") +if(NOT VOLATILE_ACCESSES_LLVM_SPIRV OR VOLATILE_ACCESSES_LLVM_SPIRV STREQUAL "NOT_NEEDED") + set(VOLATILE_ACCESSES_LLVM_SPIRV "${LLVM_TOOLS_BINARY_DIR}/llvm-spirv") +endif() +find_program(SPIRV_VAL spirv-val) + +# Every volatile access shape the pass has a rule for: 32 and 64 bit integer, +# float and pointer accesses through global and generic pointers (rewritten) +# and the 8 / 16 bit, vector, local, private, constant, under-aligned and +# already atomic ones (left alone). +list(APPEND TEST_IR_FILES ${CMAKE_CURRENT_SOURCE_DIR}/volatile-accesses.ll) + +foreach(IR_FILE ${TEST_IR_FILES}) + get_filename_component(FILENAME ${IR_FILE} NAME) + get_filename_component(BASENAME ${IR_FILE} NAME_WE) + configure_file(${IR_FILE} ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) + + add_test( + NAME hipVolatileAccesses-${BASENAME} + COMMAND env + LLVM_OPT=${LLVM_TOOLS_BINARY_DIR}/opt + LLVM_DIS=${LLVM_TOOLS_BINARY_DIR}/llvm-dis + LLVM_SPIRV=${VOLATILE_ACCESSES_LLVM_SPIRV} + SPIRV_VAL=${SPIRV_VAL} + HIP_SPV_PASSES_LIB=${CMAKE_BINARY_DIR}/lib/libLLVMHipSpvPasses.so + ${CMAKE_CURRENT_BINARY_DIR}/run_volatile_accesses_pass.bash + ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} + ) +endforeach() diff --git a/tests/compiler/volatileAccesses/run_volatile_accesses_pass.bash b/tests/compiler/volatileAccesses/run_volatile_accesses_pass.bash new file mode 100755 index 000000000..0d8841a33 --- /dev/null +++ b/tests/compiler/volatileAccesses/run_volatile_accesses_pass.bash @@ -0,0 +1,102 @@ +#!/bin/bash +# Check that HipLowerVolatileAccessesPass turns the volatile 32 and 64 bit +# global and generic accesses of a module into relaxed device-scope atomics, +# leaves every other volatile access as it is, and that the result still +# translates to valid SPIR-V. +# +# Usage: run_volatile_accesses_pass.bash +# +# The input has two kernels: @rewritten holds only accesses the pass must +# rewrite, @left_alone only accesses it must not touch. Every load and store +# in @rewritten must come out `atomic volatile ... syncscope("device") +# monotonic` on i32 or i64 (floats and pointers are laundered through the +# integer), and @left_alone must contain no `syncscope("device") monotonic` +# access at all. + +set -e + +if [ $# -ne 1 ]; then + echo "Usage: $0 " + exit 1 +fi + +INPUT_FILE="$1" +BASE_NAME=$(basename "${INPUT_FILE}" .ll) +OUTPUT_BC="${BASE_NAME}.lowered.bc" +OUTPUT_LL="${BASE_NAME}.lowered.ll" +OUTPUT_SPV="${BASE_NAME}.lowered.spv" +SPIRV_OPTS="--spirv-max-version=1.2 --spirv-ext=-all,+SPV_INTEL_function_pointers,+SPV_INTEL_subgroups" + +# CHIP_VERIFY_MODE=off: the in-pass IR->SPIR-V re-verification defaults to on +# in Debug builds and is redundant here; the translation below is the check. +CHIP_VERIFY_MODE=off "${LLVM_OPT}" -load-pass-plugin "${HIP_SPV_PASSES_LIB}" \ + -passes=hip-lower-volatile-accesses "${INPUT_FILE}" -o "${OUTPUT_BC}" \ + 2> "${BASE_NAME}.stderr" +"${LLVM_DIS}" "${OUTPUT_BC}" -o "${OUTPUT_LL}" + +kernel_body() { + sed -n "/^define .*@$1(/,/^}/p" "${OUTPUT_LL}" +} + +REWRITTEN=$(kernel_body rewritten) +LEFT=$(kernel_body left_alone) +if [ -z "${REWRITTEN}" ] || [ -z "${LEFT}" ]; then + echo "ERROR: kernels @rewritten / @left_alone not found after the pass" + exit 1 +fi + +# Volatile accesses in @rewritten that are not device-scope monotonic atomics +# on i32 / i64: the 12 volatile accesses of the input must all have become +# `load atomic volatile i32|i64 ... syncscope("device") monotonic` or the +# store equivalent. +STALE=$(echo "${REWRITTEN}" | grep -E '(load|store) volatile' || true) +if [ -n "${STALE}" ]; then + echo "ERROR: volatile access(es) in @rewritten survived the pass:" + echo "${STALE}" + exit 1 +fi +ATOMIC=$(echo "${REWRITTEN}" | grep -c -E '(load|store) atomic volatile i(32|64)[, ].*syncscope\("device"\) monotonic' || true) +if [ "${ATOMIC}" -ne 12 ]; then + echo "ERROR: expected 12 device-scope monotonic i32 / i64 atomics in @rewritten, found ${ATOMIC}" + echo "See ${OUTPUT_LL} for details" + exit 1 +fi +# Floats and pointers must be laundered through the integer, not loaded as such. +if echo "${REWRITTEN}" | grep -q -E 'atomic volatile (float|double|ptr)'; then + echo "ERROR: float / pointer typed atomic access in @rewritten:" + echo "${REWRITTEN}" | grep -E 'atomic volatile (float|double|ptr)' + exit 1 +fi + +# Nothing in @left_alone may have been rewritten. Its 15 volatile accesses go +# in as is (two of them atomic already, with their own scope and ordering). +if echo "${LEFT}" | grep -q 'syncscope("device") monotonic'; then + echo "ERROR: access(es) in @left_alone were rewritten:" + echo "${LEFT}" | grep 'syncscope("device") monotonic' + exit 1 +fi +KEPT=$(echo "${LEFT}" | grep -c -E '(load|store) (atomic )?volatile' || true) +if [ "${KEPT}" -ne 15 ]; then + echo "ERROR: expected the 15 volatile accesses of @left_alone to survive, found ${KEPT}" + echo "See ${OUTPUT_LL} for details" + exit 1 +fi +# The two under-aligned accesses are reported, not silently skipped. +WARNED=$(grep -c "leaving under-aligned volatile access alone" "${BASE_NAME}.stderr" || true) +if [ "${WARNED}" -ne 2 ]; then + echo "ERROR: expected 2 under-aligned access warnings, got ${WARNED}:" + cat "${BASE_NAME}.stderr" + exit 1 +fi + +# The rewrite is only useful if the result is valid SPIR-V. +"${LLVM_SPIRV}" "${OUTPUT_BC}" ${SPIRV_OPTS} -o "${OUTPUT_SPV}" +if [ -n "${SPIRV_VAL}" ] && [ -x "${SPIRV_VAL}" ]; then + "${SPIRV_VAL}" "${OUTPUT_SPV}" + VALIDATED="spirv-val ok" +else + VALIDATED="spirv-val not available" +fi + +echo "rewritten=${ATOMIC} left alone=${KEPT} warnings=${WARNED}, SPIR-V ok, ${VALIDATED}" +exit 0 diff --git a/tests/compiler/volatileAccesses/volatile-accesses.ll b/tests/compiler/volatileAccesses/volatile-accesses.ll new file mode 100644 index 000000000..912e2aa3b --- /dev/null +++ b/tests/compiler/volatileAccesses/volatile-accesses.ll @@ -0,0 +1,92 @@ +; Every volatile access shape HipLowerVolatileAccessesPass has a rule for. The +; 32 and 64 bit integer, float and pointer accesses through global and generic +; pointers must become relaxed device-scope atomics; everything else must come +; out exactly as it went in. + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-G1" +target triple = "spirv64" + +@lmem = internal addrspace(3) global [64 x i32] undef, align 4 +@cmem = internal addrspace(2) constant [4 x i32] zeroinitializer, align 4 + +define spir_kernel void @rewritten(ptr addrspace(1) %P32, ptr addrspace(1) %P64, + ptr addrspace(1) %PF, ptr addrspace(1) %PD, + ptr addrspace(1) %PP, + i32 %V32, i64 %V64, float %VF, double %VD, + ptr addrspace(1) %VP) { +entry: + %G32 = addrspacecast ptr addrspace(1) %P32 to ptr addrspace(4) + %GF = addrspacecast ptr addrspace(1) %PF to ptr addrspace(4) + ; integers, global and generic + %ld32 = load volatile i32, ptr addrspace(1) %P32, align 4 + %ld64 = load volatile i64, ptr addrspace(1) %P64, align 8 + %ld32g = load volatile i32, ptr addrspace(4) %G32, align 4 + store volatile i32 %V32, ptr addrspace(1) %P32, align 4 + store volatile i64 %V64, ptr addrspace(1) %P64, align 8 + store volatile i32 %V32, ptr addrspace(4) %G32, align 4 + ; floats go through the same-width integer + %ldf = load volatile float, ptr addrspace(4) %GF, align 4 + %ldd = load volatile double, ptr addrspace(1) %PD, align 8 + store volatile float %VF, ptr addrspace(4) %GF, align 4 + store volatile double %VD, ptr addrspace(1) %PD, align 8 + ; pointers go through i64 + %ldp = load volatile ptr addrspace(1), ptr addrspace(1) %PP, align 8 + store volatile ptr addrspace(1) %VP, ptr addrspace(1) %PP, align 8 + ; keep every loaded value alive + %f2i = bitcast float %ldf to i32 + %d2i = bitcast double %ldd to i64 + %p2i = ptrtoint ptr addrspace(1) %ldp to i64 + %s1 = add i32 %ld32, %ld32g + %s2 = add i32 %s1, %f2i + %s3 = add i64 %ld64, %d2i + %s4 = add i64 %s3, %p2i + store i32 %s2, ptr addrspace(1) %P32, align 4 + store i64 %s4, ptr addrspace(1) %P64, align 8 + ret void +} + +define spir_kernel void @left_alone(ptr addrspace(1) %P8, ptr addrspace(1) %P16, + ptr addrspace(1) %P32, ptr addrspace(1) %PV, + i8 %V8, i16 %V16, i32 %V32, <2 x i32> %VV) { +entry: + ; 8 and 16 bit values + %ld8 = load volatile i8, ptr addrspace(1) %P8, align 1 + %ld16 = load volatile i16, ptr addrspace(1) %P16, align 2 + store volatile i8 %V8, ptr addrspace(1) %P8, align 1 + store volatile i16 %V16, ptr addrspace(1) %P16, align 2 + ; vectors + %ldv = load volatile <2 x i32>, ptr addrspace(1) %PV, align 8 + store volatile <2 x i32> %VV, ptr addrspace(1) %PV, align 8 + ; work-group local memory reached through a generic pointer + %L = getelementptr inbounds [64 x i32], ptr addrspace(3) @lmem, i64 0, i64 5 + %LG = addrspacecast ptr addrspace(3) %L to ptr addrspace(4) + %ldl = load volatile i32, ptr addrspace(4) %LG, align 4 + store volatile i32 %V32, ptr addrspace(4) %LG, align 4 + ; private memory reached through a generic pointer + %A = alloca i32, align 4 + %AG = addrspacecast ptr %A to ptr addrspace(4) + %lda = load volatile i32, ptr addrspace(4) %AG, align 4 + store volatile i32 %V32, ptr addrspace(4) %AG, align 4 + ; constant memory + %C = getelementptr inbounds [4 x i32], ptr addrspace(2) @cmem, i64 0, i64 1 + %ldc = load volatile i32, ptr addrspace(2) %C, align 4 + ; under-aligned + %ldu = load volatile i32, ptr addrspace(1) %P32, align 2 + store volatile i32 %V32, ptr addrspace(1) %P32, align 2 + ; already atomic + %lda2 = load atomic volatile i32, ptr addrspace(1) %P32 syncscope("workgroup") acquire, align 4 + store atomic volatile i32 %V32, ptr addrspace(1) %P32 seq_cst, align 4 + ; keep every loaded value alive + %e8 = zext i8 %ld8 to i32 + %e16 = zext i16 %ld16 to i32 + %v0 = extractelement <2 x i32> %ldv, i32 0 + %s1 = add i32 %e8, %e16 + %s2 = add i32 %s1, %v0 + %s3 = add i32 %s2, %ldl + %s4 = add i32 %s3, %lda + %s5 = add i32 %s4, %ldc + %s6 = add i32 %s5, %ldu + %s7 = add i32 %s6, %lda2 + store i32 %s7, ptr addrspace(1) %P32, align 4 + ret void +} From aa9356c1a2907a476e4013b25825e15f388154ab Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Sat, 29 Aug 2026 15:25:45 +0300 Subject: [PATCH 4/4] tests: follow the entry point wrapper when checking the volatile access SPIR-V The Khronos translator built from LLVM 21 keeps an OpName on the kernel body and emits the OpEntryPoint as a separate wrapper function whose only instruction is an OpFunctionCall to it, so the entry point id in the disassembly is neither numeric nor the function that holds the accesses. TestFixVolatileLoadLoweringSPIRV looked up the body by a numeric entry point id and counted 0 OpAtomicLoad / 0 OpAtomicStore on the pocl 7.1 (llvm-21) CI lane. Accept any id and follow a wrapper's OpFunctionCall to the callee before counting. --- tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash b/tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash index 4810880da..8da516778 100755 --- a/tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash +++ b/tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash @@ -98,9 +98,17 @@ SPV=$(ls "${OUT}"/*.out 2>/dev/null | head -1) if [ -n "${SPV}" ] && [ -n "${SPIRV_DIS}" ] && [ -x "${SPIRV_DIS}" ]; then "${SPIRV_DIS}" "${SPV}" > module.spvasm if grep -q "Generator: Khronos LLVM/SPIR-V Translator" module.spvasm; then - KID=$(grep -E 'OpEntryPoint Kernel %[0-9]+ "_Z[0-9]+volatileAccess' module.spvasm | - sed -E 's/.*Kernel (%[0-9]+) .*/\1/') + # The entry point id is a number or, when the translator kept an OpName, + # the mangled name. Translators from LLVM 21 on emit the entry point as + # a wrapper whose only instruction is an OpFunctionCall to the kernel + # body, so a wrapper is followed to its callee before inspecting the body. + KID=$(grep -E 'OpEntryPoint Kernel %[^ ]+ "_Z[0-9]+volatileAccess' module.spvasm | + sed -E 's/.*Kernel (%[^ ]+) .*/\1/') FUNC=$(sed -n "/^ *${KID} = OpFunction /,/OpFunctionEnd/p" module.spvasm) + CALLEE=$(echo "${FUNC}" | grep -oE 'OpFunctionCall %[^ ]+ %[^ ]+' | awk '{print $3}' | head -1) + if [ -n "${CALLEE}" ]; then + FUNC=$(sed -n "/^ *${CALLEE} = OpFunction /,/OpFunctionEnd/p" module.spvasm) + fi LOADS=$(echo "${FUNC}" | grep -c 'OpAtomicLoad' || true) STORES=$(echo "${FUNC}" | grep -c 'OpAtomicStore' || true) if [ "${LOADS}" -lt 2 ] || [ "${STORES}" -lt 2 ]; then