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()); 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 +} 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..8da516778 --- /dev/null +++ b/tests/runtime/TestFixVolatileLoadLoweringSPIRV.bash @@ -0,0 +1,134 @@ +#!/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 + # 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 + 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"