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 <[email protected]>
+
+- 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
