Script 'mail_helper' called by obssrc Hello community, here is the log from the commit of package spirv-llvm-translator for openSUSE:Factory checked in at 2025-06-10 08:56:18 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Comparing /work/SRC/openSUSE:Factory/spirv-llvm-translator (Old) and /work/SRC/openSUSE:Factory/.spirv-llvm-translator.new.19631 (New) ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
Package is "spirv-llvm-translator" Tue Jun 10 08:56:18 2025 rev:21 rq:1283886 version:20.1.3 Changes: -------- --- /work/SRC/openSUSE:Factory/spirv-llvm-translator/spirv-llvm-translator.changes 2025-05-06 16:38:18.612345854 +0200 +++ /work/SRC/openSUSE:Factory/.spirv-llvm-translator.new.19631/spirv-llvm-translator.changes 2025-06-10 08:56:39.390015477 +0200 @@ -1,0 +2,15 @@ +Sat Jun 7 15:55:24 UTC 2025 - Aaron Puchert <aaronpuch...@alice-dsl.net> + +- Update to version 20.1.3. + * Implement SPV_KHR_bfloat16 extension. + * Remove `Aligned 0` from tests. + * Rounding modes on int to int conversions are valid OpenCL C + builtin functions. + * Add reverse translation test for integer convert with explicit + rounding. + * Make OCLUtil.h compatible with C++20 standard. + * Missing addExtension in SPIRVWriter.cpp. + * Use native for the system separator for source path string in + debug info. + +------------------------------------------------------------------- Old: ---- SPIRV-LLVM-Translator-20.1.2.tar.gz New: ---- SPIRV-LLVM-Translator-20.1.3.tar.gz ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Other differences: ------------------ ++++++ spirv-llvm-translator.spec ++++++ --- /var/tmp/diff_new_pack.QLGXBa/_old 2025-06-10 08:56:40.042042410 +0200 +++ /var/tmp/diff_new_pack.QLGXBa/_new 2025-06-10 08:56:40.042042410 +0200 @@ -23,7 +23,7 @@ %define sover 20 Name: spirv-llvm-translator -Version: 20.1.2 +Version: 20.1.3 Release: 0 Summary: LLVM/SPIR-V Bi-Directional Translator library License: BSD-3-Clause ++++++ SPIRV-LLVM-Translator-20.1.2.tar.gz -> SPIRV-LLVM-Translator-20.1.3.tar.gz ++++++ diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/include/LLVMSPIRVExtensions.inc new/SPIRV-LLVM-Translator-20.1.3/include/LLVMSPIRVExtensions.inc --- old/SPIRV-LLVM-Translator-20.1.2/include/LLVMSPIRVExtensions.inc 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/include/LLVMSPIRVExtensions.inc 2025-05-15 10:45:53.000000000 +0200 @@ -77,3 +77,4 @@ EXT(SPV_INTEL_bindless_images) EXT(SPV_INTEL_2d_block_io) EXT(SPV_INTEL_subgroup_matrix_multiply_accumulate) +EXT(SPV_KHR_bfloat16) diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/Mangler/ManglingUtils.cpp new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/Mangler/ManglingUtils.cpp --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/Mangler/ManglingUtils.cpp 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/Mangler/ManglingUtils.cpp 2025-05-15 10:45:53.000000000 +0200 @@ -28,6 +28,7 @@ "half", "float", "double", + "__bf16", "void", "...", "image1d_ro_t", @@ -105,6 +106,7 @@ "Dh", // HALF "f", // FLOAT "d", // DOUBLE + "u6__bf16", // __BF16 "v", // VOID "z", // VarArg "14ocl_image1d_ro", // PRIMITIVE_IMAGE1D_RO_T @@ -197,6 +199,7 @@ SPIR12, // HALF SPIR12, // FLOAT SPIR12, // DOUBLE + SPIR12, // __BF16 SPIR12, // VOID SPIR12, // VarArg SPIR12, // PRIMITIVE_IMAGE1D_RO_T diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/Mangler/ParameterType.h new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/Mangler/ParameterType.h --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/Mangler/ParameterType.h 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/Mangler/ParameterType.h 2025-05-15 10:45:53.000000000 +0200 @@ -45,6 +45,7 @@ PRIMITIVE_HALF, PRIMITIVE_FLOAT, PRIMITIVE_DOUBLE, + PRIMITIVE_BFLOAT, PRIMITIVE_VOID, PRIMITIVE_VAR_ARG, PRIMITIVE_STRUCT_FIRST, diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/OCLToSPIRV.cpp new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/OCLToSPIRV.cpp --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/OCLToSPIRV.cpp 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/OCLToSPIRV.cpp 2025-05-15 10:45:53.000000000 +0200 @@ -800,9 +800,6 @@ OC = OpFConvert; } - if (!Rounding.empty() && (isa<IntegerType>(SrcTy) && IsTargetInt)) - return; - assert(CI->getCalledFunction() && "Unexpected indirect call"); mutateCallInst( CI, getSPIRVFuncName(OC, "_R" + DestTy + VecSize + Sat + Rounding)); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/OCLUtil.h new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/OCLUtil.h --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/OCLUtil.h 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/OCLUtil.h 2025-05-15 10:45:53.000000000 +0200 @@ -98,11 +98,19 @@ // To avoid any inconsistence here, constants are explicitly initialized with // the corresponding constants from 'std::memory_order' enum. enum OCLMemOrderKind { +#if __cplusplus >= 202002L + OCLMO_relaxed = std::memory_order_relaxed, + OCLMO_acquire = std::memory_order_acquire, + OCLMO_release = std::memory_order_release, + OCLMO_acq_rel = std::memory_order_acq_rel, + OCLMO_seq_cst = std::memory_order_seq_cst +#else OCLMO_relaxed = std::memory_order::memory_order_relaxed, OCLMO_acquire = std::memory_order::memory_order_acquire, OCLMO_release = std::memory_order::memory_order_release, OCLMO_acq_rel = std::memory_order::memory_order_acq_rel, OCLMO_seq_cst = std::memory_order::memory_order_seq_cst +#endif }; enum IntelFPGAMemoryAccessesVal { @@ -453,11 +461,12 @@ template <typename T> std::string getFullPath(const T *Scope) { if (!Scope) return std::string(); - std::string Filename = Scope->getFilename().str(); - if (sys::path::is_absolute(Filename)) - return Filename; + StringRef Filename = Scope->getFilename(); + auto Style = sys::path::Style::native; + if (sys::path::is_absolute(Filename, Style)) + return Filename.str(); SmallString<16> DirName = Scope->getDirectory(); - sys::path::append(DirName, sys::path::Style::posix, Filename); + sys::path::append(DirName, Style, Filename.str()); return DirName.str().str(); } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/SPIRVReader.cpp new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/SPIRVReader.cpp --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/SPIRVReader.cpp 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/SPIRVReader.cpp 2025-05-15 10:45:53.000000000 +0200 @@ -314,6 +314,8 @@ Type *SPIRVToLLVM::transFPType(SPIRVType *T) { switch (T->getFloatBitWidth()) { case 16: + if (T->isTypeFloat(16, FPEncodingBFloat16KHR)) + return Type::getBFloatTy(*Context); return Type::getHalfTy(*Context); case 32: return Type::getFloatTy(*Context); @@ -1476,7 +1478,9 @@ const llvm::fltSemantics *FS = nullptr; switch (BT->getFloatBitWidth()) { case 16: - FS = &APFloat::IEEEhalf(); + FS = + (BT->isTypeFloat(16, FPEncodingBFloat16KHR) ? &APFloat::BFloat() + : &APFloat::IEEEhalf()); break; case 32: FS = &APFloat::IEEEsingle(); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/SPIRVUtil.cpp new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/SPIRVUtil.cpp --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/SPIRVUtil.cpp 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/SPIRVUtil.cpp 2025-05-15 10:45:53.000000000 +0200 @@ -1338,6 +1338,8 @@ return SPIR::RefParamType(new SPIR::PrimitiveType(SPIR::PRIMITIVE_FLOAT)); if (Ty->isDoubleTy()) return SPIR::RefParamType(new SPIR::PrimitiveType(SPIR::PRIMITIVE_DOUBLE)); + if (Ty->isBFloatTy()) + return SPIR::RefParamType(new SPIR::PrimitiveType(SPIR::PRIMITIVE_BFLOAT)); if (auto *VecTy = dyn_cast<FixedVectorType>(Ty)) { return SPIR::RefParamType(new SPIR::VectorType( transTypeDesc(VecTy->getElementType(), Info), VecTy->getNumElements())); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/SPIRVWriter.cpp new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/SPIRVWriter.cpp --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/SPIRVWriter.cpp 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/SPIRVWriter.cpp 2025-05-15 10:45:53.000000000 +0200 @@ -397,6 +397,16 @@ } } + if (T->isBFloatTy()) { + BM->getErrorLog().checkError( + BM->isAllowedToUseExtension(ExtensionID::SPV_KHR_bfloat16), + SPIRVEC_RequiresExtension, + "SPV_KHR_bfloat16\n" + "NOTE: LLVM module contains bfloat type, translation of which " + "requires this extension"); + return mapType(T, BM->addFloatType(16, FPEncodingBFloat16KHR)); + } + if (T->isFloatingPointTy()) return mapType(T, BM->addFloatType(T->getPrimitiveSizeInBits())); @@ -3118,10 +3128,12 @@ if (FMF.allowContract()) { M |= FPFastMathModeAllowContractFastINTELMask; BM->addCapability(CapabilityFPFastMathModeINTEL); + BM->addExtension(ExtensionID::SPV_INTEL_fp_fast_math_mode); } if (FMF.allowReassoc()) { M |= FPFastMathModeAllowReassocINTELMask; BM->addCapability(CapabilityFPFastMathModeINTEL); + BM->addExtension(ExtensionID::SPV_INTEL_fp_fast_math_mode); } } } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVEnum.h new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVEnum.h --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVEnum.h 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVEnum.h 2025-05-15 10:45:53.000000000 +0200 @@ -223,6 +223,9 @@ {CapabilityCooperativeMatrixKHR}); ADD_VEC_INIT(internal::CapabilityCooperativeMatrixOffsetInstructionsINTEL, {CapabilityCooperativeMatrixKHR}); + ADD_VEC_INIT(CapabilityBFloat16DotProductKHR, {CapabilityBFloat16TypeKHR}); + ADD_VEC_INIT(CapabilityBFloat16CooperativeMatrixKHR, + {CapabilityBFloat16TypeKHR, CapabilityCooperativeMatrixKHR}); } template <> inline void SPIRVMap<SPIRVExecutionModelKind, SPIRVCapVec>::init() { diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVInstruction.h new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVInstruction.h --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVInstruction.h 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVInstruction.h 2025-05-15 10:45:53.000000000 +0200 @@ -779,6 +779,18 @@ return VersionNumber::SPIRV_1_4; return VersionNumber::SPIRV_1_0; } + SPIRVCapVec getRequiredCapability() const override { + if (OpCode == OpDot) { + const SPIRVType *OpTy = getValueType(Ops[0]); + if (OpTy && OpTy->isTypeVector()) { + OpTy = OpTy->getVectorComponentType(); + if (OpTy && OpTy->isTypeFloat(16, FPEncodingBFloat16KHR)) { + return getVec(CapabilityBFloat16DotProductKHR); + } + } + } + return SPIRVInstruction::getRequiredCapability(); + } }; template <Op OC> diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVModule.cpp new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVModule.cpp --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVModule.cpp 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVModule.cpp 2025-05-15 10:45:53.000000000 +0200 @@ -253,7 +253,8 @@ template <class T> T *addType(T *Ty); SPIRVTypeArray *addArrayType(SPIRVType *, SPIRVValue *) override; SPIRVTypeBool *addBoolType() override; - SPIRVTypeFloat *addFloatType(unsigned BitWidth) override; + SPIRVTypeFloat *addFloatType(unsigned BitWidth, + unsigned FloatingPointEncoding) override; SPIRVTypeFunction *addFunctionType(SPIRVType *, const std::vector<SPIRVType *> &) override; SPIRVTypeInt *addIntegerType(unsigned BitWidth) override; @@ -580,7 +581,8 @@ SmallDenseMap<SPIRVStorageClassKind, SPIRVTypeUntypedPointerKHR *> UntypedPtrTyMap; SmallDenseMap<unsigned, SPIRVTypeInt *, 4> IntTypeMap; - SmallDenseMap<unsigned, SPIRVTypeFloat *, 4> FloatTypeMap; + SmallDenseMap<std::pair<unsigned, unsigned>, SPIRVTypeFloat *, 4> + FloatTypeMap; SmallDenseMap<std::pair<unsigned, SPIRVType *>, SPIRVTypePointer *, 4> PointerTypeMap; std::unordered_map<unsigned, SPIRVConstant *> LiteralMap; @@ -1010,12 +1012,14 @@ return addType(Ty); } -SPIRVTypeFloat *SPIRVModuleImpl::addFloatType(unsigned BitWidth) { - auto Loc = FloatTypeMap.find(BitWidth); +SPIRVTypeFloat *SPIRVModuleImpl::addFloatType(unsigned BitWidth, + unsigned FloatingPointEncoding) { + auto Desc = std::make_pair(BitWidth, FloatingPointEncoding); + auto Loc = FloatTypeMap.find(Desc); if (Loc != FloatTypeMap.end()) return Loc->second; - auto *Ty = new SPIRVTypeFloat(this, getId(), BitWidth); - FloatTypeMap[BitWidth] = Ty; + auto *Ty = new SPIRVTypeFloat(this, getId(), BitWidth, FloatingPointEncoding); + FloatTypeMap[Desc] = Ty; return addType(Ty); } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVModule.h new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVModule.h --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVModule.h 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVModule.h 2025-05-15 10:45:53.000000000 +0200 @@ -244,7 +244,7 @@ // Type creation functions virtual SPIRVTypeArray *addArrayType(SPIRVType *, SPIRVValue *) = 0; virtual SPIRVTypeBool *addBoolType() = 0; - virtual SPIRVTypeFloat *addFloatType(unsigned) = 0; + virtual SPIRVTypeFloat *addFloatType(unsigned, unsigned = FPEncodingMax) = 0; virtual SPIRVTypeFunction * addFunctionType(SPIRVType *, const std::vector<SPIRVType *> &) = 0; virtual SPIRVTypeImage *addImageType(SPIRVType *, diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h 2025-05-15 10:45:53.000000000 +0200 @@ -641,6 +641,9 @@ add(CapabilityFPGALatencyControlINTEL, "FPGALatencyControlINTEL"); add(CapabilityFPMaxErrorINTEL, "FPMaxErrorINTEL"); add(CapabilityRegisterLimitsINTEL, "RegisterLimitsINTEL"); + add(CapabilityBFloat16TypeKHR, "BFloat16TypeKHR"); + add(CapabilityBFloat16DotProductKHR, "BFloat16DotProductKHR"); + add(CapabilityBFloat16CooperativeMatrixKHR, "BFloat16CooperativeMatrixKHR"); add(CapabilityArithmeticFenceEXT, "ArithmeticFenceEXT"); add(CapabilitySubgroup2DBlockIOINTEL, "Subgroup2DBlockIOINTEL"); add(CapabilitySubgroup2DBlockTransformINTEL, "Subgroup2DBlockTransformINTEL"); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVType.cpp new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVType.cpp --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVType.cpp 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVType.cpp 2025-05-15 10:45:53.000000000 +0200 @@ -166,8 +166,15 @@ isTypeJointMatrixINTEL() || isTypeCooperativeMatrixKHR(); } -bool SPIRVType::isTypeFloat(unsigned Bits) const { - return isType<SPIRVTypeFloat>(this, Bits); +bool SPIRVType::isTypeFloat(unsigned Bits, + unsigned FloatingPointEncoding) const { + if (!isType<SPIRVTypeFloat>(this)) + return false; + if (Bits == 0) + return true; + const auto *ThisFloat = static_cast<const SPIRVTypeFloat *>(this); + return ThisFloat->getBitWidth() == Bits && + ThisFloat->getFloatingPointEncoding() == FloatingPointEncoding; } bool SPIRVType::isTypeOCLImage() const { diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVType.h new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVType.h --- old/SPIRV-LLVM-Translator-20.1.2/lib/SPIRV/libSPIRV/SPIRVType.h 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/lib/SPIRV/libSPIRV/SPIRVType.h 2025-05-15 10:45:53.000000000 +0200 @@ -84,7 +84,8 @@ bool isTypeEvent() const; bool isTypeDeviceEvent() const; bool isTypeReserveId() const; - bool isTypeFloat(unsigned Bits = 0) const; + bool isTypeFloat(unsigned Bits = 0, + unsigned FloatingPointEncoding = FPEncodingMax) const; bool isTypeImage() const; bool isTypeOCLImage() const; bool isTypePipe() const; @@ -204,16 +205,31 @@ public: static const Op OC = OpTypeFloat; // Complete constructor - SPIRVTypeFloat(SPIRVModule *M, SPIRVId TheId, unsigned TheBitWidth) - : SPIRVType(M, 3, OC, TheId), BitWidth(TheBitWidth) {} + SPIRVTypeFloat(SPIRVModule *M, SPIRVId TheId, unsigned TheBitWidth, + unsigned TheFloatingPointEncoding) + : SPIRVType(M, 3 + (TheFloatingPointEncoding != FPEncodingMax), OC, + TheId), + BitWidth(TheBitWidth), FloatingPointEncoding(TheFloatingPointEncoding) { + } // Incomplete constructor - SPIRVTypeFloat() : SPIRVType(OC), BitWidth(0) {} + SPIRVTypeFloat() + : SPIRVType(OC), BitWidth(0), FloatingPointEncoding(FPEncodingMax) {} unsigned getBitWidth() const { return BitWidth; } + unsigned getFloatingPointEncoding() const { return FloatingPointEncoding; } + + std::optional<ExtensionID> getRequiredExtension() const override { + if (isTypeFloat(16, FPEncodingBFloat16KHR)) + return ExtensionID::SPV_KHR_bfloat16; + return {}; + } + SPIRVCapVec getRequiredCapability() const override { SPIRVCapVec CV; - if (isTypeFloat(16)) { + if (isTypeFloat(16, FPEncodingBFloat16KHR)) { + CV.push_back(CapabilityBFloat16TypeKHR); + } else if (isTypeFloat(16)) { CV.push_back(CapabilityFloat16Buffer); auto Extensions = getModule()->getSourceExtension(); if (std::any_of(Extensions.begin(), Extensions.end(), @@ -225,14 +241,34 @@ } protected: - _SPIRV_DEF_ENCDEC2(Id, BitWidth) + void encode(spv_ostream &O) const override { + assert(WordCount == 3 || WordCount == 4); + auto Encoder = getEncoder(O); + Encoder << Id << BitWidth; + if (WordCount == 4) + Encoder << FloatingPointEncoding; + } + + void decode(std::istream &I) override { + assert(WordCount == 3 || WordCount == 4); + auto Decoder = getDecoder(I); + Decoder >> Id >> BitWidth; + if (WordCount == 4) + Decoder >> FloatingPointEncoding; + } + void validate() const override { SPIRVEntry::validate(); assert(BitWidth >= 16 && BitWidth <= 64 && "Invalid bit width"); + assert( + (FloatingPointEncoding == FPEncodingMax || + (BitWidth == 16 && FloatingPointEncoding == FPEncodingBFloat16KHR)) && + "Invalid floating point encoding"); } private: unsigned BitWidth; // Bit width + unsigned FloatingPointEncoding; }; template <Op TheOpCode = OpTypePointer, SPIRVWord WC = 3> @@ -1186,7 +1222,10 @@ return ExtensionID::SPV_KHR_cooperative_matrix; } SPIRVCapVec getRequiredCapability() const override { - return getVec(CapabilityCooperativeMatrixKHR); + auto CV = getVec(CapabilityCooperativeMatrixKHR); + if (CompType->isTypeFloat(16, FPEncodingBFloat16KHR)) + CV.push_back(CapabilityBFloat16CooperativeMatrixKHR); + return CV; } SPIRVType *getCompType() const { return CompType; } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/spirv-headers-tag.conf new/SPIRV-LLVM-Translator-20.1.3/spirv-headers-tag.conf --- old/SPIRV-LLVM-Translator-20.1.2/spirv-headers-tag.conf 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/spirv-headers-tag.conf 2025-05-15 10:45:53.000000000 +0200 @@ -1 +1 @@ -2b2e05e088841c63c0b6fd4c9fb380d8688738d3 +0e710677989b4326ac974fd80c5308191ed80965 diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/DebugInfo/LocalAddressSpace.ll new/SPIRV-LLVM-Translator-20.1.3/test/DebugInfo/LocalAddressSpace.ll --- old/SPIRV-LLVM-Translator-20.1.2/test/DebugInfo/LocalAddressSpace.ll 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/DebugInfo/LocalAddressSpace.ll 2025-05-15 10:45:53.000000000 +0200 @@ -22,7 +22,7 @@ ; CHECK: DW_TAG_variable ; CHECK-NEXT: DW_AT_name {{.*}} = "a") ; CHECK-NEXT: DW_AT_type {{.*}} "int") -; CHECK-NEXT: DW_AT_decl_file {{.*}} ("/work/tmp{{[/\\]}}tmp.cl") +; CHECK-NEXT: DW_AT_decl_file {{.*}} ("/work{{[/\\]}}tmp{{[/\\]}}tmp.cl") ; CHECK-NEXT: DW_AT_decl_line {{.*}} (2) ; CHECK-NEXT: DW_AT_location [DW_FORM_exprloc] (DW_OP_addr 0x0) diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/DebugInfo/RelativeSrcPath.ll new/SPIRV-LLVM-Translator-20.1.3/test/DebugInfo/RelativeSrcPath.ll --- old/SPIRV-LLVM-Translator-20.1.2/test/DebugInfo/RelativeSrcPath.ll 1970-01-01 01:00:00.000000000 +0100 +++ new/SPIRV-LLVM-Translator-20.1.3/test/DebugInfo/RelativeSrcPath.ll 2025-05-15 10:45:53.000000000 +0200 @@ -0,0 +1,61 @@ +; Source: +; __kernel void foo(__global int *a, __global int *b) { +; a[0] += b[0]; +; } + +; Command: +; clang -cc1 -triple spir -O0 -debug-info-kind=line-tables-only -emit-llvm -o RelativeSrcPath.ll RelativeSrcPath.cl + +; Directory: /tmp + +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s + +; ModuleID = 'RelativeSrcPath.cl' +source_filename = "RelativeSrcPath.cl" +target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir" + +; Function Attrs: convergent noinline nounwind optnone +define spir_kernel void @foo(ptr addrspace(1) %a, ptr addrspace(1) %b) #0 !dbg !8 !kernel_arg_addr_space !11 !kernel_arg_access_qual !12 !kernel_arg_type !13 !kernel_arg_base_type !13 !kernel_arg_type_qual !14 { +entry: + %a.addr = alloca ptr addrspace(1), align 4 + %b.addr = alloca ptr addrspace(1), align 4 + store ptr addrspace(1) %a, ptr %a.addr, align 4 + store ptr addrspace(1) %b, ptr %b.addr, align 4 + %0 = load ptr addrspace(1), ptr %b.addr, align 4, !dbg !15 + %1 = load i32, ptr addrspace(1) %0, align 4, !dbg !15 + %2 = load ptr addrspace(1), ptr %a.addr, align 4, !dbg !15 + %3 = load i32, ptr addrspace(1) %2, align 4, !dbg !15 + %add = add nsw i32 %3, %1, !dbg !15 + store i32 %add, ptr addrspace(1) %2, align 4, !dbg !15 + ret void, !dbg !16 +} + +attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!3, !4} +!opencl.ocl.version = !{!5} +!opencl.spir.version = !{!6} +!llvm.ident = !{!7} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang version 8.0.0 (cfe/trunk)", isOptimized: false, runtimeVersion: 0, emissionKind: LineTablesOnly, enums: !2, nameTableKind: None) +!1 = !DIFile(filename: "<stdin>", directory: "/tmp") +!2 = !{} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !{i32 1, !"wchar_size", i32 4} +!5 = !{i32 1, i32 0} +!6 = !{i32 1, i32 2} +!7 = !{!"clang version 8.0.0 (cfe/trunk)"} +!8 = distinct !DISubprogram(name: "foo", scope: !9, file: !9, line: 1, type: !10, isLocal: false, isDefinition: true, scopeLine: 1, flags: DIFlagPrototyped, isOptimized: false, unit: !0, retainedNodes: !2) +; CHECK: String [[ID:[0-9]+]] "/tmp{{[/\\]}}RelativeSrcPath.cl" +; CHECK: Line [[ID]] +!9 = !DIFile(filename: "RelativeSrcPath.cl", directory: "/tmp") +!10 = !DISubroutineType(types: !2) +!11 = !{i32 1, i32 1} +!12 = !{!"none", !"none"} +!13 = !{!"int*", !"int*"} +!14 = !{!"", !""} +!15 = !DILocation(line: 2, scope: !8) +!16 = !DILocation(line: 3, scope: !8) diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/OpLoopMergeNone.spt new/SPIRV-LLVM-Translator-20.1.3/test/OpLoopMergeNone.spt --- old/SPIRV-LLVM-Translator-20.1.2/test/OpLoopMergeNone.spt 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/OpLoopMergeNone.spt 2025-05-15 10:45:53.000000000 +0200 @@ -39,7 +39,7 @@ 2 Label 20 4 Variable 18 26 7 4 Variable 18 27 7 -6 Load 9 21 7 2 0 +4 Load 9 21 7 5 CompositeExtract 8 22 21 0 5 ShiftLeftLogical 8 23 22 11 5 ShiftRightArithmetic 8 24 23 11 diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/RelativeSrcPath.ll new/SPIRV-LLVM-Translator-20.1.3/test/RelativeSrcPath.ll --- old/SPIRV-LLVM-Translator-20.1.2/test/RelativeSrcPath.ll 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/RelativeSrcPath.ll 1970-01-01 01:00:00.000000000 +0100 @@ -1,61 +0,0 @@ -; Source: -; __kernel void foo(__global int *a, __global int *b) { -; a[0] += b[0]; -; } - -; Command: -; clang -cc1 -triple spir -O0 -debug-info-kind=line-tables-only -emit-llvm -o RelativeSrcPath.ll RelativeSrcPath.cl - -; Directory: /tmp - -; RUN: llvm-as %s -o %t.bc -; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s - -; ModuleID = 'RelativeSrcPath.cl' -source_filename = "RelativeSrcPath.cl" -target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" -target triple = "spir" - -; Function Attrs: convergent noinline nounwind optnone -define spir_kernel void @foo(ptr addrspace(1) %a, ptr addrspace(1) %b) #0 !dbg !8 !kernel_arg_addr_space !11 !kernel_arg_access_qual !12 !kernel_arg_type !13 !kernel_arg_base_type !13 !kernel_arg_type_qual !14 { -entry: - %a.addr = alloca ptr addrspace(1), align 4 - %b.addr = alloca ptr addrspace(1), align 4 - store ptr addrspace(1) %a, ptr %a.addr, align 4 - store ptr addrspace(1) %b, ptr %b.addr, align 4 - %0 = load ptr addrspace(1), ptr %b.addr, align 4, !dbg !15 - %1 = load i32, ptr addrspace(1) %0, align 4, !dbg !15 - %2 = load ptr addrspace(1), ptr %a.addr, align 4, !dbg !15 - %3 = load i32, ptr addrspace(1) %2, align 4, !dbg !15 - %add = add nsw i32 %3, %1, !dbg !15 - store i32 %add, ptr addrspace(1) %2, align 4, !dbg !15 - ret void, !dbg !16 -} - -attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } - -!llvm.dbg.cu = !{!0} -!llvm.module.flags = !{!3, !4} -!opencl.ocl.version = !{!5} -!opencl.spir.version = !{!6} -!llvm.ident = !{!7} - -!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang version 8.0.0 (cfe/trunk)", isOptimized: false, runtimeVersion: 0, emissionKind: LineTablesOnly, enums: !2, nameTableKind: None) -!1 = !DIFile(filename: "<stdin>", directory: "/tmp") -!2 = !{} -!3 = !{i32 2, !"Debug Info Version", i32 3} -!4 = !{i32 1, !"wchar_size", i32 4} -!5 = !{i32 1, i32 0} -!6 = !{i32 1, i32 2} -!7 = !{!"clang version 8.0.0 (cfe/trunk)"} -!8 = distinct !DISubprogram(name: "foo", scope: !9, file: !9, line: 1, type: !10, isLocal: false, isDefinition: true, scopeLine: 1, flags: DIFlagPrototyped, isOptimized: false, unit: !0, retainedNodes: !2) -; CHECK: String [[ID:[0-9]+]] "/tmp/RelativeSrcPath.cl" -; CHECK: Line [[ID]] -!9 = !DIFile(filename: "RelativeSrcPath.cl", directory: "/tmp") -!10 = !DISubroutineType(types: !2) -!11 = !{i32 1, i32 1} -!12 = !{!"none", !"none"} -!13 = !{!"int*", !"int*"} -!14 = !{!"", !""} -!15 = !DILocation(line: 2, scope: !8) -!16 = !DILocation(line: 3, scope: !8) diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/composite_construct_struct.spt new/SPIRV-LLVM-Translator-20.1.3/test/composite_construct_struct.spt --- old/SPIRV-LLVM-Translator-20.1.2/test/composite_construct_struct.spt 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/composite_construct_struct.spt 2025-05-15 10:45:53.000000000 +0200 @@ -37,7 +37,7 @@ 2 Label 21 5 CompositeConstruct 11 22 16 17 5 CompositeConstruct 12 23 20 22 -6 Load 5 24 3 2 0 +4 Load 5 24 3 5 CompositeExtract 4 25 24 0 5 ShiftLeftLogical 4 26 25 15 5 ShiftRightArithmetic 4 27 26 15 diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/composite_construct_vector.spt new/SPIRV-LLVM-Translator-20.1.3/test/composite_construct_vector.spt --- old/SPIRV-LLVM-Translator-20.1.2/test/composite_construct_vector.spt 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/composite_construct_vector.spt 2025-05-15 10:45:53.000000000 +0200 @@ -30,7 +30,7 @@ 3 FunctionParameter 10 2 2 Label 17 -6 Load 5 18 3 2 0 +4 Load 5 18 3 5 CompositeExtract 4 19 18 0 5 ShiftLeftLogical 4 20 19 12 5 ShiftRightArithmetic 4 21 20 12 diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/copy_object.spt new/SPIRV-LLVM-Translator-20.1.3/test/copy_object.spt --- old/SPIRV-LLVM-Translator-20.1.2/test/copy_object.spt 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/copy_object.spt 2025-05-15 10:45:53.000000000 +0200 @@ -27,7 +27,7 @@ 3 FunctionParameter 9 2 2 Label 13 -6 Load 5 14 3 2 0 +4 Load 5 14 3 5 CompositeExtract 4 15 14 0 5 ShiftLeftLogical 4 16 15 11 5 ShiftRightArithmetic 4 17 16 11 diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/extensions/KHR/SPV_KHR_bfloat16/bfloat16.ll new/SPIRV-LLVM-Translator-20.1.3/test/extensions/KHR/SPV_KHR_bfloat16/bfloat16.ll --- old/SPIRV-LLVM-Translator-20.1.2/test/extensions/KHR/SPV_KHR_bfloat16/bfloat16.ll 1970-01-01 01:00:00.000000000 +0100 +++ new/SPIRV-LLVM-Translator-20.1.3/test/extensions/KHR/SPV_KHR_bfloat16/bfloat16.ll 2025-05-15 10:45:53.000000000 +0200 @@ -0,0 +1,36 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_bfloat16 -o %t.spv +; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV + +; RUN: llvm-spirv -r %t.spv -o %t.rev.bc +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM + +; RUN: not llvm-spirv %t.bc 2>&1 | FileCheck %s --check-prefix=CHECK-ERROR + +; CHECK-ERROR: RequiresExtension: Feature requires the following SPIR-V extension: +; CHECK-ERROR-NEXT: SPV_KHR_bfloat16 +; CHECK-ERROR-NEXT: NOTE: LLVM module contains bfloat type, translation of which +; CHECK-ERROR-SAME: requires this extension + +source_filename = "bfloat16.cpp" +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" +target triple = "spirv64-unknown-unknown" + +; CHECK-SPIRV-DAG: Capability BFloat16TypeKHR +; CHECK-SPIRV-DAG: Extension "SPV_KHR_bfloat16" +; CHECK-SPIRV: 4 TypeFloat [[BFLOAT:[0-9]+]] 16 0 +; CHECK-SPIRV: 4 TypeVector [[#]] [[BFLOAT]] 2 + +; CHECK-LLVM: [[ADDR1:]] = alloca bfloat +; CHECK-LLVM: [[ADDR2:]] = alloca <2 x bfloat> +; CHECK-LLVM: [[DATA1:]] = load bfloat, ptr [[ADDR1]] +; CHECK-LLVM: [[DATA2:]] = load <2 x bfloat>, ptr [[ADDR2]] + +define spir_kernel void @test() { +entry: + %addr1 = alloca bfloat + %addr2 = alloca <2 x bfloat> + %data1 = load bfloat, ptr %addr1 + %data2 = load <2 x bfloat>, ptr %addr2 + ret void +} diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/extensions/KHR/SPV_KHR_bfloat16/bfloat16_dot.ll new/SPIRV-LLVM-Translator-20.1.3/test/extensions/KHR/SPV_KHR_bfloat16/bfloat16_dot.ll --- old/SPIRV-LLVM-Translator-20.1.2/test/extensions/KHR/SPV_KHR_bfloat16/bfloat16_dot.ll 1970-01-01 01:00:00.000000000 +0100 +++ new/SPIRV-LLVM-Translator-20.1.3/test/extensions/KHR/SPV_KHR_bfloat16/bfloat16_dot.ll 2025-05-15 10:45:53.000000000 +0200 @@ -0,0 +1,39 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_bfloat16 -o %t.spv +; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV + +; RUN: llvm-spirv -r %t.spv -o %t.rev.bc +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM + +source_filename = "bfloat16_dot.cpp" +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" +target triple = "spirv64-unknown-unknown" + +; CHECK-SPIRV-DAG: Capability BFloat16TypeKHR +; CHECK-SPIRV-DAG: Capability BFloat16DotProductKHR +; CHECK-SPIRV-DAG: Extension "SPV_KHR_bfloat16" +; CHECK-SPIRV: 4 TypeFloat [[BFLOAT:[0-9]+]] 16 0 +; CHECK-SPIRV: 4 TypeVector [[#]] [[BFLOAT]] 2 +; CHECK-SPIRV: Dot + +; CHECK-LLVM: %addrA = alloca <2 x bfloat> +; CHECK-LLVM: %addrB = alloca <2 x bfloat> +; CHECK-LLVM: %dataA = load <2 x bfloat>, ptr %addrA +; CHECK-LLVM: %dataB = load <2 x bfloat>, ptr %addrB +; CHECK-LLVM: %call = call spir_func bfloat @_Z3dotDv2_u6__bf16S_(<2 x bfloat> %dataA, <2 x bfloat> %dataB) + +declare spir_func bfloat @_Z3dotDv2_u6__bf16Dv2_S_(<2 x bfloat>, <2 x bfloat>) + +define spir_kernel void @test() { +entry: + %addrA = alloca <2 x bfloat> + %addrB = alloca <2 x bfloat> + %dataA = load <2 x bfloat>, ptr %addrA + %dataB = load <2 x bfloat>, ptr %addrB + %call = call spir_func bfloat @_Z3dotDv2_u6__bf16Dv2_S_(<2 x bfloat> %dataA, <2 x bfloat> %dataB) + ret void +} + +!opencl.ocl.version = !{!7} + +!7 = !{i32 2, i32 0} diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/extensions/KHR/SPV_KHR_bfloat16/cooperative_matrix_bfloat16.ll new/SPIRV-LLVM-Translator-20.1.3/test/extensions/KHR/SPV_KHR_bfloat16/cooperative_matrix_bfloat16.ll --- old/SPIRV-LLVM-Translator-20.1.2/test/extensions/KHR/SPV_KHR_bfloat16/cooperative_matrix_bfloat16.ll 1970-01-01 01:00:00.000000000 +0100 +++ new/SPIRV-LLVM-Translator-20.1.3/test/extensions/KHR/SPV_KHR_bfloat16/cooperative_matrix_bfloat16.ll 2025-05-15 10:45:53.000000000 +0200 @@ -0,0 +1,34 @@ +; RUN: llvm-as < %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_cooperative_matrix,+SPV_KHR_bfloat16 -o %t.spv +; RUN: llvm-spirv %t.spv -to-text -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV + +; RUN: llvm-spirv --spirv-ext=+SPV_KHR_cooperative_matrix,+SPV_KHR_bfloat16 -r %t.spv -o %t.rev.bc +; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM + +; CHECK-SPIRV-DAG: Capability CooperativeMatrixKHR +; CHECK-SPIRV-DAG: Capability BFloat16TypeKHR +; CHECK-SPIRV-DAG: Capability BFloat16CooperativeMatrixKHR +; CHECK-SPIRV-DAG: Extension "SPV_KHR_cooperative_matrix" +; CHECK-SPIRV-DAG: Extension "SPV_KHR_bfloat16" + +; CHECK-SPIRV-DAG: 4 TypeFloat [[#BFloatTy:]] 16 0 +; CHECK-SPIRV-DAG: TypeInt [[#Int32Ty:]] 32 0 +; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const12:]] 12 +; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const3:]] 3 +; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const2:]] 2 +; CHECK-SPIRV-DAG: TypeCooperativeMatrixKHR [[#MatTy:]] [[#BFloatTy]] [[#Const3]] [[#Const12]] [[#Const12]] [[#Const2]] +; CHECK-SPIRV-DAG: Constant [[#BFloatTy]] [[#]] 16256 +; CHECK-SPIRV: CompositeConstruct [[#MatTy]] + +; CHECK-LLVM: call spir_func target("spirv.CooperativeMatrixKHR", bfloat, 3, 12, 12, 2) @_Z26__spirv_CompositeConstructu6__bf16(bfloat 0xR3F80) + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" +target triple = "spir64-unknown-unknown" + +declare spir_func target("spirv.CooperativeMatrixKHR", bfloat, 3, 12, 12, 2) @_Z26__spirv_CompositeConstructu6__bf16(bfloat) + +define spir_kernel void @test() { + %mat = call spir_func target("spirv.CooperativeMatrixKHR", bfloat, 3, 12, 12, 2) @_Z26__spirv_CompositeConstructu6__bf16(bfloat 1.0) + ret void +} diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/negative/SPV_INTEL_bfloat16_conversion/f2bf16_inval_output_ty.spt new/SPIRV-LLVM-Translator-20.1.3/test/negative/SPV_INTEL_bfloat16_conversion/f2bf16_inval_output_ty.spt --- old/SPIRV-LLVM-Translator-20.1.2/test/negative/SPV_INTEL_bfloat16_conversion/f2bf16_inval_output_ty.spt 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/negative/SPV_INTEL_bfloat16_conversion/f2bf16_inval_output_ty.spt 2025-05-15 10:45:53.000000000 +0200 @@ -18,7 +18,7 @@ 6 Decorate 4 LinkageAttributes "_Z1f" Export 4 Decorate 8 Alignment 4 -4 TypeFloat 10 16 +3 TypeFloat 10 16 2 TypeVoid 2 3 TypeFunction 3 2 3 TypeFloat 6 32 diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/right_shift.spt new/SPIRV-LLVM-Translator-20.1.3/test/right_shift.spt --- old/SPIRV-LLVM-Translator-20.1.2/test/right_shift.spt 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/right_shift.spt 2025-05-15 10:45:53.000000000 +0200 @@ -31,7 +31,7 @@ 3 FunctionParameter 10 2 2 Label 17 -6 Load 5 18 3 2 0 +4 Load 5 18 3 5 CompositeExtract 4 19 18 0 5 ShiftRightArithmetic 4 20 19 12 5 ShiftLeftLogical 4 21 20 25 diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/selection_merge.spt new/SPIRV-LLVM-Translator-20.1.3/test/selection_merge.spt --- old/SPIRV-LLVM-Translator-20.1.2/test/selection_merge.spt 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/selection_merge.spt 2025-05-15 10:45:53.000000000 +0200 @@ -35,7 +35,7 @@ 2 Label 18 4 Variable 16 27 7 -6 Load 8 19 6 2 0 +4 Load 8 19 6 5 CompositeExtract 7 20 19 0 5 ShiftLeftLogical 7 21 20 13 5 ShiftRightArithmetic 7 22 21 13 diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/transcoding/OpenCL/convert_functions.ll new/SPIRV-LLVM-Translator-20.1.3/test/transcoding/OpenCL/convert_functions.ll --- old/SPIRV-LLVM-Translator-20.1.2/test/transcoding/OpenCL/convert_functions.ll 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/transcoding/OpenCL/convert_functions.ll 2025-05-15 10:45:53.000000000 +0200 @@ -14,9 +14,11 @@ ; CHECK-SPIRV: Name [[#Func1:]] "_Z20convert_uint_satfunc" ; CHECK-SPIRV: Name [[#Func2:]] "_Z21convert_float_rtzfunc" ; CHECK-SPIRV-DAG: TypeVoid [[#VoidTy:]] +; CHECK-SPIRV-DAG: TypeInt [[#CharTy:]] 8 ; CHECK-SPIRV-DAG: TypeFloat [[#FloatTy:]] 32 ; CHECK-SPIRV: Function [[#VoidTy]] [[#Func]] +; CHECK-SPIRV: SConvert [[#CharTy]] [[#ConvertId:]] [[#]] ; CHECK-SPIRV: ConvertSToF [[#FloatTy]] [[#ConvertId:]] [[#]] ; CHECK-SPIRV: FunctionCall [[#VoidTy]] [[#]] [[#Func]] [[#ConvertId]] ; CHECK-SPIRV: FunctionCall [[#VoidTy]] [[#]] [[#Func1]] [[#]] @@ -51,6 +53,11 @@ %x.addr = alloca i32, align 4 store i32 %x, ptr %x.addr, align 4 %0 = load i32, ptr %x.addr, align 4 +; We don't get the convert_char_rtei back, but that's fine because they are +; functionally equivalent anyway. The rounding information is lost when +; translating to SPIR-V. +; CHECK-LLVM: call spir_func i8 @_Z12convert_chari(i32 %[[#]]) + call spir_func signext i8 @_Z16convert_char_rtei(i32 noundef %0) #1 ; CHECK-LLVM: %[[Call:[a-z]+]] = sitofp i32 %[[#]] to float ; CHECK-LLVM: call spir_func void @_Z18convert_float_func(float %[[Call]]) ; CHECK-LLVM: call spir_func void @_Z20convert_uint_satfunc(i32 %[[#]]) @@ -63,6 +70,9 @@ } ; Function Attrs: convergent nounwind willreturn memory(none) +declare spir_func signext i8 @_Z16convert_char_rtei(i32 noundef) #1 + +; Function Attrs: convergent nounwind willreturn memory(none) declare spir_func float @_Z13convert_floati(i32 noundef) #1 attributes #0 = { convergent nounwind } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/SPIRV-LLVM-Translator-20.1.2/test/vector_times_scalar.spt new/SPIRV-LLVM-Translator-20.1.3/test/vector_times_scalar.spt --- old/SPIRV-LLVM-Translator-20.1.2/test/vector_times_scalar.spt 2025-04-10 14:21:56.000000000 +0200 +++ new/SPIRV-LLVM-Translator-20.1.3/test/vector_times_scalar.spt 2025-05-15 10:45:53.000000000 +0200 @@ -34,7 +34,7 @@ 3 FunctionParameter 13 4 2 Label 17 -6 Load 8 18 6 2 0 +4 Load 8 18 6 5 CompositeExtract 7 19 18 0 5 ShiftLeftLogical 7 20 19 10 5 ShiftRightArithmetic 7 21 20 10