From 2324e55c004c319b9fba4be9e1cc64069b17f303 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Mon, 24 Oct 2022 22:35:59 +0200 Subject: [PATCH] [PATCH 48/79] [Backport to 15] [NFC] Remove JointMatrixINTEL W/S (#1658) It's not longer needed after https://github.com/intel/llvm/pull/6535 Signed-off-by: Sidorov, Dmitry Gbp-Pq: Name 0048-Backport-to-15-NFC-Remove-JointMatrixINTEL-W-S-1658.patch --- lib/SPIRV/SPIRVRegularizeLLVM.cpp | 115 ------------------ lib/SPIRV/SPIRVRegularizeLLVM.h | 1 - .../SPV_INTEL_joint_matrix/joint_matrix.ll | 40 +++--- .../joint_matrix_bfloat16.ll | 25 ++-- .../joint_matrix_element.ll | 24 ++-- ...rix_extract_insert_element_of_sycl_half.ll | 16 +-- .../joint_matrix_half.ll | 34 +++--- 7 files changed, 63 insertions(+), 192 deletions(-) diff --git a/lib/SPIRV/SPIRVRegularizeLLVM.cpp b/lib/SPIRV/SPIRVRegularizeLLVM.cpp index e1e2f70..443272a 100644 --- a/lib/SPIRV/SPIRVRegularizeLLVM.cpp +++ b/lib/SPIRV/SPIRVRegularizeLLVM.cpp @@ -347,118 +347,6 @@ Value *SPIRVRegularizeLLVMBase::extendBitInstBoolArg(Instruction *II) { } } -void SPIRVRegularizeLLVMBase::adaptStructTypes(StructType *ST) { - if (!ST->hasName()) - return; - StringRef STName = ST->getName(); - STName.consume_front("struct."); - STName.consume_front("__spv::"); - StringRef MangledName = STName.substr(0, STName.find('.')); - - // Representation in LLVM IR before the translator is a pointer array wrapped - // in a structure: - // %struct.__spirv_JointMatrixINTEL = type { [R x [C x [L x [S x type]]]]* } - // where R = Rows, C = Columnts, L = Layout + 1, S = Scope + 1 - // this '+1' for the Layout and Scope is required because both of them can - // be '0', but array size can not be '0'. - // The result should look like SPIR-V friendly LLVM IR: - // %spirv.JointMatrixINTEL._char_2_2_0_3 - // Here we check the structure name yet again. Another option would be to - // check SPIR-V friendly function calls (by their name) and obtain return - // or their parameter types, assuming, that the appropriate types are Matrix - // structure type. But in the near future, we will reuse Composite - // instructions to do, for example, matrix initialization directly on AMX - // register by OpCompositeConstruct. And we can't claim, that the Result type - // of OpCompositeConstruct instruction is always the joint matrix type, it's - // simply not true. - if (MangledName == "__spirv_JointMatrixINTEL" && !ST->isOpaquePointerTy()) { - auto *PtrTy = dyn_cast(ST->getElementType(0)); - assert(PtrTy && - "Expected a pointer to an array to represent joint matrix type"); - std::vector TypeLayout; - ArrayType *ArrayTy = - dyn_cast(PtrTy->getNonOpaquePointerElementType()); - assert(ArrayTy && "Expected a pointer element type of an array type to " - "represent joint matrix type"); - TypeLayout.push_back(ArrayTy->getNumElements()); - for (size_t I = 1; I != 4; ++I) { - ArrayTy = dyn_cast(ArrayTy->getElementType()); - assert(ArrayTy && - "Expected a element type to represent joint matrix type"); - TypeLayout.push_back(ArrayTy->getNumElements()); - } - // JointMatrixINTEL type can have optional 'Use' parameter, which is encoded - // as another array dimention. In case if it has default 'Unnecessary' (4) - // parameter - ignore it. - if (isa(ArrayTy->getElementType())) { - ArrayTy = cast(ArrayTy->getElementType()); - uint32_t UseInt = ArrayTy->getNumElements(); - assert(UseInt <= 4 && "Use parameter encoded in the array must be < 5 "); - if (UseInt != 4) - TypeLayout.push_back(UseInt); - } - - auto *ElemTy = ArrayTy->getElementType(); - std::string ElemTyStr; - if (ElemTy->isIntegerTy()) { - auto *IntElemTy = cast(ElemTy); - switch (IntElemTy->getBitWidth()) { - case 8: - ElemTyStr = "char"; - break; - case 16: - ElemTyStr = "short"; - break; - case 32: - ElemTyStr = "int"; - break; - case 64: - ElemTyStr = "long"; - break; - default: - ElemTyStr = "i" + std::to_string(IntElemTy->getBitWidth()); - } - } - // Check half type like this as well, but in DPC++ it most likelly will - // be a class - else if (ElemTy->isHalfTy()) - ElemTyStr = "half"; - else if (ElemTy->isFloatTy()) - ElemTyStr = "float"; - else if (ElemTy->isDoubleTy()) - ElemTyStr = "double"; - else { - // Half type is special: in DPC++ we use `class half` instead of `half` - // type natively supported by Clang. - auto *STElemTy = dyn_cast(ElemTy); - if (!STElemTy && !STElemTy->hasName()) - llvm_unreachable("Unexpected type for matrix!"); - if (isSYCLHalfType(ElemTy)) - ElemTyStr = "half"; - if (isSYCLBfloat16Type(ElemTy)) - ElemTyStr = "bfloat16"; - if (ElemTyStr.size() == 0) - llvm_unreachable("Unexpected type for matrix!"); - } - std::stringstream SPVName; - SPVName << kSPIRVTypeName::PrefixAndDelim - << kSPIRVTypeName::JointMatrixINTEL << kSPIRVTypeName::Delimiter - << kSPIRVTypeName::PostfixDelim << ElemTyStr - << kSPIRVTypeName::PostfixDelim << std::to_string(TypeLayout[0]) - << kSPIRVTypeName::PostfixDelim << std::to_string(TypeLayout[1]) - << kSPIRVTypeName::PostfixDelim << std::to_string(TypeLayout[2] - 1) - << kSPIRVTypeName::PostfixDelim - << std::to_string(TypeLayout[3] - 1); - if (TypeLayout.size() == 5) - SPVName << kSPIRVTypeName::PostfixDelim - << std::to_string(TypeLayout[4] - 1); - // Note, that this structure is not opaque and there is no way to make it - // opaque but to recreate it entirely and replace it everywhere. Lets - // keep the structure as is, dealing with it during SPIR-V generation. - ST->setName(SPVName.str()); - } -} - bool SPIRVRegularizeLLVMBase::runRegularizeLLVM(Module &Module) { M = &Module; Ctx = &M->getContext(); @@ -623,9 +511,6 @@ bool SPIRVRegularizeLLVMBase::regularize() { } } - for (StructType *ST : M->getIdentifiedStructTypes()) - adaptStructTypes(ST); - if (SPIRVDbgSaveRegularizedModule) saveLLVMModule(M, RegularizedModuleTmpFile); return true; diff --git a/lib/SPIRV/SPIRVRegularizeLLVM.h b/lib/SPIRV/SPIRVRegularizeLLVM.h index f8b3d08..1bcb5f0 100644 --- a/lib/SPIRV/SPIRVRegularizeLLVM.h +++ b/lib/SPIRV/SPIRVRegularizeLLVM.h @@ -105,7 +105,6 @@ public: Value *extendBitInstBoolArg(llvm::Instruction *OldInst); static std::string lowerLLVMIntrinsicName(llvm::IntrinsicInst *II); - void adaptStructTypes(llvm::StructType *ST); static char ID; private: diff --git a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix.ll b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix.ll index 959ea8e..89c3e21 100644 --- a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix.ll +++ b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix.ll @@ -1,16 +1,10 @@ ; RUN: llvm-as < %s -o %t.bc -; RUN: llvm-spirv %t.bc -s -o %t.pre.bc -; RUN: llvm-dis %t.pre.bc -o - | FileCheck %s --check-prefix=CHECK-PRE ; RUN: llvm-spirv %t.bc -spirv-ext=+SPV_INTEL_joint_matrix -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 -; CHECK-PRE: %spirv.JointMatrixINTEL._short_2_2_0_3 -; CHECK-PRE: %spirv.JointMatrixINTEL._char_2_16_0_3_0 -; CHECK-PRE: %spirv.JointMatrixINTEL._char_16_2_3_3 - ; CHECK-SPIRV: Capability JointMatrixINTEL ; CHECK-SPIRV: Extension "SPV_INTEL_joint_matrix" ; CHECK-SPIRV: Name [[#Kernel:]] "_ZTSZ4mainE11matrix_test" @@ -67,9 +61,9 @@ source_filename = "./joint_matrix_test.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 = "spir64-unknown-unknown" -%"struct.__spv::__spirv_JointMatrixINTEL" = type { [2 x [2 x [1 x [4 x [4 x i16]]]]]* } -%"struct.__spv::__spirv_JointMatrixINTEL.0" = type { [2 x [16 x [1 x [4 x [1 x i8]]]]]* } -%"struct.__spv::__spirv_JointMatrixINTEL.2" = type { [16 x [2 x [4 x [4 x i8]]]]* } +%spirv.JointMatrixINTEL._short_2_2_0_3 = type { [2 x [2 x [1 x [4 x [4 x i16]]]]]* } +%spirv.JointMatrixINTEL._char_2_16_0_3_0 = type { [2 x [16 x [1 x [4 x [1 x i8]]]]]* } +%spirv.JointMatrixINTEL._char_16_2_3_3 = type { [16 x [2 x [4 x [4 x i8]]]]* } $_ZTSZ4mainE11matrix_test = comdat any @@ -99,14 +93,14 @@ entry: %add.ptr.i51 = getelementptr inbounds i16, i16 addrspace(1)* %_arg_, i64 %mul6.i %add.ptr7.i52 = getelementptr inbounds i16, i16 addrspace(1)* %add.ptr.i51, i64 %sub5.i %add.ptr7.i = addrspacecast i16 addrspace(1)* %add.ptr7.i52 to i16 addrspace(4)* - %call8.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIsLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i16 addrspace(4)* %add.ptr7.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 + %call8.i = tail call spir_func %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIsLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i16 addrspace(4)* %add.ptr7.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 %add.ptr11.i53 = getelementptr inbounds i8, i8 addrspace(1)* %_arg_3, i64 %mul6.i %add.ptr16.i55 = getelementptr inbounds i8, i8 addrspace(1)* %_arg_5, i64 %sub5.i br label %for.cond.i for.cond.i: ; preds = %for.body.i, %entry %k.0.i = phi i32 [ 0, %entry ], [ %add.i, %for.body.i ] - %C.0.i = phi %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* [ %call8.i, %entry ], [ %call19.i, %for.body.i ] + %C.0.i = phi %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* [ %call8.i, %entry ], [ %call19.i, %for.body.i ] %cmp.i = icmp ult i32 %k.0.i, 32 br i1 %cmp.i, label %for.body.i, label %_ZZ4mainENKUlN2cl4sycl7nd_itemILi2EEEE_clES2_.exit @@ -114,45 +108,45 @@ for.body.i: ; preds = %for.cond.i %idx.ext.i = zext i32 %k.0.i to i64 %add.ptr12.i54 = getelementptr inbounds i8, i8 addrspace(1)* %add.ptr11.i53, i64 %idx.ext.i %add.ptr12.i = addrspacecast i8 addrspace(1)* %add.ptr12.i54 to i8 addrspace(4)* - %call13.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL.0" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIaLm2ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i8 addrspace(4)* %add.ptr12.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 + %call13.i = tail call spir_func %spirv.JointMatrixINTEL._char_2_16_0_3_0 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIaLm2ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i8 addrspace(4)* %add.ptr12.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 %mul14.i = shl nuw nsw i32 %k.0.i, 5 %idx.ext15.i = zext i32 %mul14.i to i64 %add.ptr17.i56 = getelementptr inbounds i8, i8 addrspace(1)* %add.ptr16.i55, i64 %idx.ext15.i %add.ptr17.i = addrspacecast i8 addrspace(1)* %add.ptr17.i56 to i8 addrspace(4)* - %call18.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL.2" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIaLm16ELm2ELN5__spv12MatrixLayoutE3ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i8 addrspace(4)* %add.ptr17.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 - %call19.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_JointMatrixMadINTELIasLm2ELm16ELm2ELN5__spv12MatrixLayoutE0ELS1_3ELS1_0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT6_EXT7_EEEPNS4_IT_XT1_EXT2_EXT4_EXT7_EEEPNS4_IS8_XT2_EXT3_EXT5_EXT7_EEES7_S3_(%"struct.__spv::__spirv_JointMatrixINTEL.0" addrspace(4)* %call13.i, %"struct.__spv::__spirv_JointMatrixINTEL.2" addrspace(4)* %call18.i, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %C.0.i, i32 3) #3 + %call18.i = tail call spir_func %spirv.JointMatrixINTEL._char_16_2_3_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIaLm16ELm2ELN5__spv12MatrixLayoutE3ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i8 addrspace(4)* %add.ptr17.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 + %call19.i = tail call spir_func %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* @_Z27__spirv_JointMatrixMadINTELIasLm2ELm16ELm2ELN5__spv12MatrixLayoutE0ELS1_3ELS1_0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT6_EXT7_EEEPNS4_IT_XT1_EXT2_EXT4_EXT7_EEEPNS4_IS8_XT2_EXT3_EXT5_EXT7_EEES7_S3_(%spirv.JointMatrixINTEL._char_2_16_0_3_0 addrspace(4)* %call13.i, %spirv.JointMatrixINTEL._char_16_2_3_3 addrspace(4)* %call18.i, %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* %C.0.i, i32 3) #3 %add.i = add nuw nsw i32 %k.0.i, 16 br label %for.cond.i, !llvm.loop !19 _ZZ4mainENKUlN2cl4sycl7nd_itemILi2EEEE_clES2_.exit: ; preds = %for.cond.i - tail call spir_func void @_Z29__spirv_JointMatrixStoreINTELIsLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(i16 addrspace(4)* %add.ptr7.i, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %C.0.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 - %C.0.i.new = call spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z26__spirv_CompositeConstructi(i32 42) #1 + tail call spir_func void @_Z29__spirv_JointMatrixStoreINTELIsLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(i16 addrspace(4)* %add.ptr7.i, %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* %C.0.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 + %C.0.i.new = call spir_func %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* @_Z26__spirv_CompositeConstructi(i32 42) #1 %ref.tmp = alloca i32, align 4 %ref.tmp.ascast = addrspacecast i32* %ref.tmp to i32 addrspace(4)* store i32 0, i32 addrspace(4)* %ref.tmp.ascast, align 4 %zero = load i32, i32 addrspace(4)* %ref.tmp.ascast, align 8 - %C.0.i.new.load = call spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z26__spirv_CompositeConstructi(i32 %zero) #1 + %C.0.i.new.load = call spir_func %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* @_Z26__spirv_CompositeConstructi(i32 %zero) #1 ret void } ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIsLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i16 addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIsLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i16 addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL.0" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIaLm2ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i8 addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._char_2_16_0_3_0 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIaLm2ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i8 addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL.2" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIaLm16ELm2ELN5__spv12MatrixLayoutE3ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i8 addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._char_16_2_3_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIaLm16ELm2ELN5__spv12MatrixLayoutE3ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(i8 addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_JointMatrixMadINTELIasLm2ELm16ELm2ELN5__spv12MatrixLayoutE0ELS1_3ELS1_0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT6_EXT7_EEEPNS4_IT_XT1_EXT2_EXT4_EXT7_EEEPNS4_IS8_XT2_EXT3_EXT5_EXT7_EEES7_S3_(%"struct.__spv::__spirv_JointMatrixINTEL.0" addrspace(4)*, %"struct.__spv::__spirv_JointMatrixINTEL.2" addrspace(4)*, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, i32) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* @_Z27__spirv_JointMatrixMadINTELIasLm2ELm16ELm2ELN5__spv12MatrixLayoutE0ELS1_3ELS1_0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT6_EXT7_EEEPNS4_IT_XT1_EXT2_EXT4_EXT7_EEEPNS4_IS8_XT2_EXT3_EXT5_EXT7_EEES7_S3_(%spirv.JointMatrixINTEL._char_2_16_0_3_0 addrspace(4)*, %spirv.JointMatrixINTEL._char_16_2_3_3 addrspace(4)*, %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)*, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func void @_Z29__spirv_JointMatrixStoreINTELIsLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(i16 addrspace(4)*, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func void @_Z29__spirv_JointMatrixStoreINTELIsLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(i16 addrspace(4)*, %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z26__spirv_CompositeConstructi(i32) #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._short_2_2_0_3 addrspace(4)* @_Z26__spirv_CompositeConstructi(i32) #1 ; Function Attrs: inaccessiblememonly nofree nosync nounwind willreturn declare void @llvm.assume(i1 noundef) #2 diff --git a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_bfloat16.ll b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_bfloat16.ll index 0cf2dd0..a7cd188 100644 --- a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_bfloat16.ll +++ b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_bfloat16.ll @@ -11,7 +11,6 @@ ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc ; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM -; CHECK-REGULARIZED: %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 ; CHECK-REGULARIZED: %[[Alloca:.*]] = alloca %"class.cl::sycl::ext::intel::experimental::bfloat16", align 2 ; CHECK-REGULARIZED: %[[ASCast:.*]] = addrspacecast %"class.cl::sycl::ext::intel::experimental::bfloat16"* %[[Alloca]] to %"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* ; CHECK-REGULARIZED: %[[GEP1:.*]] = getelementptr inbounds %"class.cl::sycl::ext::intel::experimental::bfloat16", %"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* %[[ASCast]], i64 0, i32 0 @@ -65,7 +64,7 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::item.0" = type { %"struct.cl::sycl::detail::ItemBase.1" } %"struct.cl::sycl::detail::ItemBase.1" = type { %"class.cl::sycl::range", %"class.cl::sycl::id" } %"class.cl::sycl::group" = type { %"class.cl::sycl::range", %"class.cl::sycl::range", %"class.cl::sycl::range", %"class.cl::sycl::id" } -%"struct.__spv::__spirv_JointMatrixINTEL" = type { [8 x [16 x [1 x [4 x %"class.cl::sycl::ext::intel::experimental::bfloat16"]]]] addrspace(4)* } +%spirv.JointMatrixINTEL._bfloat16_8_16_0_3 = type opaque $_ZZZ17matrix_verify_addIN2cl4sycl3ext5intel12experimental8bfloat16ELm16ELm16EEvNS1_5queueER10big_matrixIT_XT0_EXT1_EERNS1_8nd_rangeILi2EEEfENKUlRNS1_7handlerEE_clESF_ENKUlNS1_7nd_itemILi2EEEE_clESI_ = comdat any @@ -101,7 +100,7 @@ entry: %call.i.i.i = tail call spir_func noundef zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float noundef 5.000000e+00) #6 %value.i.i = getelementptr inbounds %"class.cl::sycl::ext::intel::experimental::bfloat16", %"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* %agg.tmp.ascast.i, i64 0, i32 0 store i16 %call.i.i.i, i16 addrspace(4)* %value.i.i, align 2, !tbaa !9 - %call.i = tail call spir_func noundef %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z26__spirv_CompositeConstructIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESB_(%"class.cl::sycl::ext::intel::experimental::bfloat16"* noundef nonnull byval(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2 %agg.tmp.i) #7 + %call.i = tail call spir_func noundef %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* @_Z26__spirv_CompositeConstructIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESB_(%"class.cl::sycl::ext::intel::experimental::bfloat16"* noundef nonnull byval(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2 %agg.tmp.i) #7 call void @llvm.lifetime.end.p0i8(i64 2, i8* nonnull %4) %ref.tmp.ascast.i = addrspacecast %"class.cl::sycl::ext::intel::experimental::bfloat16"* %ref.tmp.i to %"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* %5 = bitcast %"class.cl::sycl::ext::intel::experimental::bfloat16"* %ref.tmp.i to i8* @@ -112,10 +111,10 @@ entry: br label %for.cond for.cond: ; preds = %for.body, %entry - %sub_a.sroa.0.0 = phi %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* [ %call.i, %entry ], [ %call.i58, %for.body ] + %sub_a.sroa.0.0 = phi %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* [ %call.i, %entry ], [ %call.i58, %for.body ] %i.0 = phi i32 [ 0, %entry ], [ %inc, %for.body ] %conv = zext i32 %i.0 to i64 - %call.i41 = call spir_func noundef i64 @_Z38__spirv_JointMatrixWorkItemLengthINTELIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEmPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEE(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* noundef %sub_a.sroa.0.0) #7 + %call.i41 = call spir_func noundef i64 @_Z38__spirv_JointMatrixWorkItemLengthINTELIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEmPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEE(%spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* noundef %sub_a.sroa.0.0) #7 %cmp = icmp ugt i64 %call.i41, %conv br i1 %cmp, label %for.body, label %for.cond.cleanup @@ -129,13 +128,13 @@ for.cond.cleanup: ; preds = %for.cond %div = and i64 %sub5, -8 %add.ptr.i45 = getelementptr inbounds %"class.cl::sycl::ext::intel::experimental::bfloat16", %"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(1)* %add.ptr.i, i64 %div %call.ascast.i = addrspacecast %"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(1)* %add.ptr.i45 to %"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* - call spir_func void @_Z29__spirv_JointMatrixStoreINTELIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEvPT_PNS6_24__spirv_JointMatrixINTELISA_XT0_EXT1_EXT2_EXT3_EEEmS7_S9_i(%"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* noundef %call.ascast.i, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* noundef %sub_a.sroa.0.0, i64 noundef 16, i32 noundef 0, i32 noundef 3, i32 noundef 0) #7 + call spir_func void @_Z29__spirv_JointMatrixStoreINTELIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEvPT_PNS6_24__spirv_JointMatrixINTELISA_XT0_EXT1_EXT2_EXT3_EEEmS7_S9_i(%"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* noundef %call.ascast.i, %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* noundef %sub_a.sroa.0.0, i64 noundef 16, i32 noundef 0, i32 noundef 3, i32 noundef 0) #7 ret void for.body: ; preds = %for.cond %call.i.i = call spir_func noundef zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float noundef 2.000000e+00) #6 call void @llvm.lifetime.start.p0i8(i64 2, i8* nonnull %5) #8, !noalias !16 - call spir_func void @_Z28__spirv_VectorExtractDynamicIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EET_PNS6_24__spirv_JointMatrixINTELISA_XT0_EXT1_EXT2_EXT3_EEEm(%"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* sret(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2 %ref.tmp.ascast.i, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* noundef %sub_a.sroa.0.0, i64 noundef %conv) #7, !noalias !16 + call spir_func void @_Z28__spirv_VectorExtractDynamicIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EET_PNS6_24__spirv_JointMatrixINTELISA_XT0_EXT1_EXT2_EXT3_EEEm(%"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* sret(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2 %ref.tmp.ascast.i, %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* noundef %sub_a.sroa.0.0, i64 noundef %conv) #7, !noalias !16 %10 = load i16, i16 addrspace(4)* %value.i.i.i, align 2, !tbaa !19, !noalias !20 %call.i.i.i.i = call spir_func noundef float @_Z27__spirv_ConvertBF16ToFINTELt(i16 noundef zeroext %10) #6, !noalias !20 %call.i.i3.i.i = call spir_func noundef float @_Z27__spirv_ConvertBF16ToFINTELt(i16 noundef zeroext %call.i.i) #6, !noalias !20 @@ -144,7 +143,7 @@ for.body: ; preds = %for.cond call void @llvm.lifetime.end.p0i8(i64 2, i8* nonnull %5) #8, !noalias !16 call void @llvm.lifetime.start.p0i8(i64 2, i8* nonnull %6) store i16 %call.i.i4.i.i, i16 addrspace(4)* %8, align 2, !tbaa !19 - %call.i58 = call spir_func noundef %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_VectorInsertDynamicIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESD_SB_m(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* noundef %sub_a.sroa.0.0, %"class.cl::sycl::ext::intel::experimental::bfloat16"* noundef nonnull byval(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2 %agg.tmp.i54, i64 noundef %conv) #7 + %call.i58 = call spir_func noundef %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* @_Z27__spirv_VectorInsertDynamicIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESD_SB_m(%spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* noundef %sub_a.sroa.0.0, %"class.cl::sycl::ext::intel::experimental::bfloat16"* noundef nonnull byval(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2 %agg.tmp.i54, i64 noundef %conv) #7 call void @llvm.lifetime.end.p0i8(i64 2, i8* nonnull %6) %inc = add nuw nsw i32 %i.0, 1 br label %for.cond, !llvm.loop !23 @@ -157,13 +156,13 @@ declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #0 declare void @llvm.assume(i1 noundef) #3 ; Function Attrs: convergent -declare dso_local spir_func noundef %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z26__spirv_CompositeConstructIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESB_(%"class.cl::sycl::ext::intel::experimental::bfloat16"* noundef byval(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2) local_unnamed_addr #4 +declare dso_local spir_func noundef %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* @_Z26__spirv_CompositeConstructIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESB_(%"class.cl::sycl::ext::intel::experimental::bfloat16"* noundef byval(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2) local_unnamed_addr #4 ; Function Attrs: convergent -declare dso_local spir_func noundef i64 @_Z38__spirv_JointMatrixWorkItemLengthINTELIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEmPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEE(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* noundef) local_unnamed_addr #4 +declare dso_local spir_func noundef i64 @_Z38__spirv_JointMatrixWorkItemLengthINTELIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEmPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEE(%spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* noundef) local_unnamed_addr #4 ; Function Attrs: convergent -declare dso_local spir_func void @_Z28__spirv_VectorExtractDynamicIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EET_PNS6_24__spirv_JointMatrixINTELISA_XT0_EXT1_EXT2_EXT3_EEEm(%"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* sret(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* noundef, i64 noundef) local_unnamed_addr #4 +declare dso_local spir_func void @_Z28__spirv_VectorExtractDynamicIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EET_PNS6_24__spirv_JointMatrixINTELISA_XT0_EXT1_EXT2_EXT3_EEEm(%"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* sret(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2, %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* noundef, i64 noundef) local_unnamed_addr #4 ; Function Attrs: convergent nounwind declare dso_local spir_func noundef float @_Z27__spirv_ConvertBF16ToFINTELt(i16 noundef zeroext) local_unnamed_addr #5 @@ -172,10 +171,10 @@ declare dso_local spir_func noundef float @_Z27__spirv_ConvertBF16ToFINTELt(i16 declare dso_local spir_func noundef zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float noundef) local_unnamed_addr #5 ; Function Attrs: convergent -declare dso_local spir_func noundef %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_VectorInsertDynamicIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESD_SB_m(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* noundef, %"class.cl::sycl::ext::intel::experimental::bfloat16"* noundef byval(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2, i64 noundef) local_unnamed_addr #4 +declare dso_local spir_func noundef %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* @_Z27__spirv_VectorInsertDynamicIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESD_SB_m(%spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* noundef, %"class.cl::sycl::ext::intel::experimental::bfloat16"* noundef byval(%"class.cl::sycl::ext::intel::experimental::bfloat16") align 2, i64 noundef) local_unnamed_addr #4 ; Function Attrs: convergent -declare dso_local spir_func void @_Z29__spirv_JointMatrixStoreINTELIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEvPT_PNS6_24__spirv_JointMatrixINTELISA_XT0_EXT1_EXT2_EXT3_EEEmS7_S9_i(%"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* noundef, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* noundef, i64 noundef, i32 noundef, i32 noundef, i32 noundef) local_unnamed_addr #4 +declare dso_local spir_func void @_Z29__spirv_JointMatrixStoreINTELIN2cl4sycl3ext5intel12experimental8bfloat16ELm8ELm16ELN5__spv12MatrixLayoutE0ELNS6_5Scope4FlagE3EEvPT_PNS6_24__spirv_JointMatrixINTELISA_XT0_EXT1_EXT2_EXT3_EEEmS7_S9_i(%"class.cl::sycl::ext::intel::experimental::bfloat16" addrspace(4)* noundef, %spirv.JointMatrixINTEL._bfloat16_8_16_0_3 addrspace(4)* noundef, i64 noundef, i32 noundef, i32 noundef, i32 noundef) local_unnamed_addr #4 attributes #0 = { argmemonly nofree nosync nounwind willreturn } attributes #1 = { convergent inlinehint norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } diff --git a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_element.ll b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_element.ll index 11db9a0..f6519cf 100644 --- a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_element.ll +++ b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_element.ll @@ -29,7 +29,7 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } %"class.cl::sycl::detail::array" = type { [1 x i64] } %"class.cl::sycl::id" = type { %"class.cl::sycl::detail::array" } -%"struct.__spv::__spirv_JointMatrixINTEL" = type { [16 x [16 x [1 x [4 x float]]]] addrspace(4)* } +%spirv.JointMatrixINTEL._float_16_16_0_3 = type opaque $_ZTSN2cl4sycl6detail23__sycl_service_kernel__16AssertInfoCopierE = comdat any @@ -54,43 +54,43 @@ declare extern_weak dso_local spir_func void @__devicelib_assert_read(i8 addrspa ; Function Attrs: convergent norecurse define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E6matrix() local_unnamed_addr #0 comdat !kernel_arg_buffer_location !6 { entry: - %call9.i.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(float addrspace(4)* addrspacecast (float addrspace(1)* null to float addrspace(4)*), i64 1, i32 0, i32 3, i32 0) #2 + %call9.i.i = tail call spir_func %spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(float addrspace(4)* addrspacecast (float addrspace(1)* null to float addrspace(4)*), i64 1, i32 0, i32 3, i32 0) #2 br label %for.cond.i for.cond.i: ; preds = %for.body.i, %entry - %A.sroa.0.0.i = phi %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* [ %call9.i.i, %entry ], [ %call5.i.i, %for.body.i ] + %A.sroa.0.0.i = phi %spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)* [ %call9.i.i, %entry ], [ %call5.i.i, %for.body.i ] %i.0.i = phi i32 [ 0, %entry ], [ %inc.i, %for.body.i ] %conv.i = zext i32 %i.0.i to i64 - %call.i12.i = tail call spir_func i64 @_Z38__spirv_JointMatrixWorkItemLengthINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEmPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEE(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %A.sroa.0.0.i) #2 + %call.i12.i = tail call spir_func i64 @_Z38__spirv_JointMatrixWorkItemLengthINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEmPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEE(%spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)* %A.sroa.0.0.i) #2 %cmp.i = icmp ugt i64 %call.i12.i, %conv.i br i1 %cmp.i, label %for.body.i, label %_ZZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_ENKUlNS0_7nd_itemILi2EEEE_clES5_.exit for.body.i: ; preds = %for.cond.i - %call.i.i = tail call spir_func float @_Z28__spirv_VectorExtractDynamicIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EmET_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEET4_(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %A.sroa.0.0.i, i64 %conv.i) #2 + %call.i.i = tail call spir_func float @_Z28__spirv_VectorExtractDynamicIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EmET_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEET4_(%spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)* %A.sroa.0.0.i, i64 %conv.i) #2 %mul.i.i = fmul float %call.i.i, 5.000000e+00 - %call5.i.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_VectorInsertDynamicIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EmEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEES7_T4_S5_(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %A.sroa.0.0.i, float %mul.i.i, i64 %conv.i) #2 + %call5.i.i = tail call spir_func %spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)* @_Z27__spirv_VectorInsertDynamicIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EmEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEES7_T4_S5_(%spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)* %A.sroa.0.0.i, float %mul.i.i, i64 %conv.i) #2 %inc.i = add nuw nsw i32 %i.0.i, 1 br label %for.cond.i, !llvm.loop !7 _ZZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_ENKUlNS0_7nd_itemILi2EEEE_clES5_.exit: ; preds = %for.cond.i - tail call spir_func void @_Z29__spirv_JointMatrixStoreINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(float addrspace(4)* addrspacecast (float addrspace(1)* null to float addrspace(4)*), %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %A.sroa.0.0.i, i64 1, i32 0, i32 3, i32 0) #2 + tail call spir_func void @_Z29__spirv_JointMatrixStoreINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(float addrspace(4)* addrspacecast (float addrspace(1)* null to float addrspace(4)*), %spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)* %A.sroa.0.0.i, i64 1, i32 0, i32 3, i32 0) #2 ret void } ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(float addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(float addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func i64 @_Z38__spirv_JointMatrixWorkItemLengthINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEmPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEE(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*) local_unnamed_addr #1 +declare dso_local spir_func i64 @_Z38__spirv_JointMatrixWorkItemLengthINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEmPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEE(%spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)*) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func float @_Z28__spirv_VectorExtractDynamicIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EmET_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEET4_(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, i64) local_unnamed_addr #1 +declare dso_local spir_func float @_Z28__spirv_VectorExtractDynamicIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EmET_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEET4_(%spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)*, i64) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_VectorInsertDynamicIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EmEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEES7_T4_S5_(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, float, i64) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)* @_Z27__spirv_VectorInsertDynamicIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EmEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEES7_T4_S5_(%spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)*, float, i64) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func void @_Z29__spirv_JointMatrixStoreINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(float addrspace(4)*, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func void @_Z29__spirv_JointMatrixStoreINTELIfLm16ELm16ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(float addrspace(4)*, %spirv.JointMatrixINTEL._float_16_16_0_3 addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="/work/tmp/matrix-slice.cpp" "uniform-work-group-size"="true" } attributes #1 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } diff --git a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_extract_insert_element_of_sycl_half.ll b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_extract_insert_element_of_sycl_half.ll index a105dff..1192be6 100644 --- a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_extract_insert_element_of_sycl_half.ll +++ b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_extract_insert_element_of_sycl_half.ll @@ -44,8 +44,8 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256: target triple = "spir64-unknown-unknown" %"class.cl::sycl::detail::half_impl::half" = type { half } -%"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix" = type { %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* } -%"struct.__spv::__spirv_JointMatrixINTEL" = type { [8 x [16 x [1 x [4 x %"class.cl::sycl::detail::half_impl::half"]]]] addrspace(4)* } +%"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix" = type { %spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)* } +%spirv.JointMatrixINTEL._half_8_16_0_3 = type opaque %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element" = type { %"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix" addrspace(4)*, i64 } $_ZN2cl4sycl3ext6oneapi12experimental6matrixplERKNS4_10wi_elementINS0_6detail9half_impl4halfELm8ELm16ELNS4_13matrix_layoutE0ENS2_9sub_groupEEERKS8_ = comdat any @@ -63,11 +63,11 @@ entry: %M = getelementptr inbounds %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element", %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element" addrspace(4)* %0, i32 0, i32 0 %1 = load %"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix" addrspace(4)*, %"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix" addrspace(4)* addrspace(4)* %M, align 8, !tbaa !15 %spvm = getelementptr inbounds %"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix", %"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix" addrspace(4)* %1, i32 0, i32 0 - %2 = load %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* addrspace(4)* %spvm, align 8, !tbaa !13 + %2 = load %spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)*, %spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)* addrspace(4)* %spvm, align 8, !tbaa !13 %3 = load %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element" addrspace(4)*, %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element" addrspace(4)* addrspace(4)* %lhs.addr.ascast, align 8, !tbaa !8 %idx = getelementptr inbounds %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element", %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element" addrspace(4)* %3, i32 0, i32 1 %4 = load i64, i64 addrspace(4)* %idx, align 8, !tbaa !17 - call spir_func void @_Z28__spirv_VectorExtractDynamicIN2cl4sycl6detail9half_impl4halfELm8ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EET_PNS5_24__spirv_JointMatrixINTELIS9_XT0_EXT1_EXT2_EXT3_EEEm(%"class.cl::sycl::detail::half_impl::half" addrspace(4)* sret(%"class.cl::sycl::detail::half_impl::half") align 2 %ref.tmp1.ascast, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %2, i64 %4) #2 + call spir_func void @_Z28__spirv_VectorExtractDynamicIN2cl4sycl6detail9half_impl4halfELm8ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EET_PNS5_24__spirv_JointMatrixINTELIS9_XT0_EXT1_EXT2_EXT3_EEEm(%"class.cl::sycl::detail::half_impl::half" addrspace(4)* sret(%"class.cl::sycl::detail::half_impl::half") align 2 %ref.tmp1.ascast, %spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)* %2, i64 %4) #2 ret void } @@ -82,19 +82,19 @@ entry: %M = getelementptr inbounds %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element", %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element" addrspace(4)* %this1, i32 0, i32 0 %0 = load %"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix" addrspace(4)*, %"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix" addrspace(4)* addrspace(4)* %M, align 8, !tbaa !15 %spvm = getelementptr inbounds %"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix", %"struct.cl::sycl::ext::oneapi::experimental::matrix::joint_matrix" addrspace(4)* %0, i32 0, i32 0 - %1 = load %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* addrspace(4)* %spvm, align 8, !tbaa !13 + %1 = load %spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)*, %spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)* addrspace(4)* %spvm, align 8, !tbaa !13 %idx = getelementptr inbounds %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element", %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element" addrspace(4)* %this1, i32 0, i32 1 %2 = load i64, i64 addrspace(4)* %idx, align 8, !tbaa !17 %agg.tmp.ascast.ascast = addrspacecast %"class.cl::sycl::detail::half_impl::half" addrspace(4)* %agg.tmp.ascast to %"class.cl::sycl::detail::half_impl::half"* - %call = call spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_VectorInsertDynamicIN2cl4sycl6detail9half_impl4halfELm8ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESC_SA_m(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %1, %"class.cl::sycl::detail::half_impl::half"* byval(%"class.cl::sycl::detail::half_impl::half") align 2 %agg.tmp.ascast.ascast, i64 %2) #2 + %call = call spir_func %spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)* @_Z27__spirv_VectorInsertDynamicIN2cl4sycl6detail9half_impl4halfELm8ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESC_SA_m(%spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)* %1, %"class.cl::sycl::detail::half_impl::half"* byval(%"class.cl::sycl::detail::half_impl::half") align 2 %agg.tmp.ascast.ascast, i64 %2) #2 ret %"class.cl::sycl::ext::oneapi::experimental::matrix::wi_element" addrspace(4)* %this1 } ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_VectorInsertDynamicIN2cl4sycl6detail9half_impl4halfELm8ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESC_SA_m(%"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, %"class.cl::sycl::detail::half_impl::half"* byval(%"class.cl::sycl::detail::half_impl::half") align 2, i64) #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)* @_Z27__spirv_VectorInsertDynamicIN2cl4sycl6detail9half_impl4halfELm8ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEESC_SA_m(%spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)*, %"class.cl::sycl::detail::half_impl::half"* byval(%"class.cl::sycl::detail::half_impl::half") align 2, i64) #1 ; Function Attrs: convergent -declare dso_local spir_func void @_Z28__spirv_VectorExtractDynamicIN2cl4sycl6detail9half_impl4halfELm8ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EET_PNS5_24__spirv_JointMatrixINTELIS9_XT0_EXT1_EXT2_EXT3_EEEm(%"class.cl::sycl::detail::half_impl::half" addrspace(4)* sret(%"class.cl::sycl::detail::half_impl::half") align 2, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, i64) #1 +declare dso_local spir_func void @_Z28__spirv_VectorExtractDynamicIN2cl4sycl6detail9half_impl4halfELm8ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EET_PNS5_24__spirv_JointMatrixINTELIS9_XT0_EXT1_EXT2_EXT3_EEEm(%"class.cl::sycl::detail::half_impl::half" addrspace(4)* sret(%"class.cl::sycl::detail::half_impl::half") align 2, %spirv.JointMatrixINTEL._half_8_16_0_3 addrspace(4)*, i64) #1 attributes #0 = { convergent mustprogress norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } attributes #1 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } diff --git a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_half.ll b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_half.ll index 94aa124..0b2b5dd 100644 --- a/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_half.ll +++ b/test/transcoding/SPV_INTEL_joint_matrix/joint_matrix_half.ll @@ -1,16 +1,10 @@ ; RUN: llvm-as < %s -o %t.bc -; RUN: llvm-spirv %t.bc -s -o %t.pre.bc -; RUN: llvm-dis %t.pre.bc -o - | FileCheck %s --check-prefix=CHECK-PRE ; RUN: llvm-spirv %t.bc -spirv-ext=+SPV_INTEL_joint_matrix -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 -; CHECK-PRE: %spirv.JointMatrixINTEL._float_2_2_0_3 -; CHECK-PRE: %spirv.JointMatrixINTEL._half_2_16_0_3 -; CHECK-PRE: %spirv.JointMatrixINTEL._half_16_2_3_3 - ; CHECK-SPIRV-DAG: TypeFloat [[#FloatTy:]] 32 ; CHECK-SPIRV-DAG: TypeFloat [[#HalfTy:]] 16 ; CHECK-SPIRV-DAG: TypeInt [[#IntTy:]] 32 0 @@ -36,9 +30,9 @@ target triple = "spir64-unknown-unknown" %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } %"class.cl::sycl::detail::half_impl::half" = type { half } -%"struct.__spv::__spirv_JointMatrixINTEL" = type { [2 x [2 x [1 x [4 x float]]]]* } -%"struct.__spv::__spirv_JointMatrixINTEL.0" = type { [2 x [16 x [1 x [4 x %"class.cl::sycl::detail::half_impl::half"]]]]* } -%"struct.__spv::__spirv_JointMatrixINTEL.1" = type { [16 x [2 x [4 x [4 x %"class.cl::sycl::detail::half_impl::half"]]]]* } +%spirv.JointMatrixINTEL._float_2_2_0_3 = type opaque +%spirv.JointMatrixINTEL._half_2_16_0_3 = type opaque +%spirv.JointMatrixINTEL._half_16_2_3_3 = type opaque $_ZTSN2cl4sycl6detail16AssertInfoCopierE = comdat any @@ -86,14 +80,14 @@ entry: %add.ptr.i51 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %mul6.i %add.ptr7.i52 = getelementptr inbounds float, float addrspace(1)* %add.ptr.i51, i64 %sub5.i %add.ptr7.i = addrspacecast float addrspace(1)* %add.ptr7.i52 to float addrspace(4)* - %call8.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIfLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(float addrspace(4)* %add.ptr7.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 + %call8.i = tail call spir_func %spirv.JointMatrixINTEL._float_2_2_0_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIfLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(float addrspace(4)* %add.ptr7.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 %add.ptr11.i53 = getelementptr inbounds %"class.cl::sycl::detail::half_impl::half", %"class.cl::sycl::detail::half_impl::half" addrspace(1)* %_arg_3, i64 %mul6.i %add.ptr16.i55 = getelementptr inbounds %"class.cl::sycl::detail::half_impl::half", %"class.cl::sycl::detail::half_impl::half" addrspace(1)* %_arg_5, i64 %sub5.i br label %for.cond.i for.cond.i: ; preds = %for.body.i, %entry %k.0.i = phi i32 [ 0, %entry ], [ %add.i, %for.body.i ] - %C.0.i = phi %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* [ %call8.i, %entry ], [ %call19.i, %for.body.i ] + %C.0.i = phi %spirv.JointMatrixINTEL._float_2_2_0_3 addrspace(4)* [ %call8.i, %entry ], [ %call19.i, %for.body.i ] %cmp.i = icmp ult i32 %k.0.i, 32 br i1 %cmp.i, label %for.body.i, label %_ZZ4mainENKUlN2cl4sycl7nd_itemILi2EEEE_clES2_.exit @@ -101,35 +95,35 @@ for.body.i: ; preds = %for.cond.i %idx.ext46.i = zext i32 %k.0.i to i64 %add.ptr12.i54 = getelementptr inbounds %"class.cl::sycl::detail::half_impl::half", %"class.cl::sycl::detail::half_impl::half" addrspace(1)* %add.ptr11.i53, i64 %idx.ext46.i %add.ptr12.i = addrspacecast %"class.cl::sycl::detail::half_impl::half" addrspace(1)* %add.ptr12.i54 to %"class.cl::sycl::detail::half_impl::half" addrspace(4)* - %call13.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL.0" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIN2cl4sycl6detail9half_impl4halfELm2ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPSA_mS6_S8_i(%"class.cl::sycl::detail::half_impl::half" addrspace(4)* %add.ptr12.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 + %call13.i = tail call spir_func %spirv.JointMatrixINTEL._half_2_16_0_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIN2cl4sycl6detail9half_impl4halfELm2ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPSA_mS6_S8_i(%"class.cl::sycl::detail::half_impl::half" addrspace(4)* %add.ptr12.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 %mul14.i = shl nuw nsw i32 %k.0.i, 5 %idx.ext1547.i = zext i32 %mul14.i to i64 %add.ptr17.i56 = getelementptr inbounds %"class.cl::sycl::detail::half_impl::half", %"class.cl::sycl::detail::half_impl::half" addrspace(1)* %add.ptr16.i55, i64 %idx.ext1547.i %add.ptr17.i = addrspacecast %"class.cl::sycl::detail::half_impl::half" addrspace(1)* %add.ptr17.i56 to %"class.cl::sycl::detail::half_impl::half" addrspace(4)* - %call18.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL.1" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIN2cl4sycl6detail9half_impl4halfELm16ELm2ELN5__spv12MatrixLayoutE3ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPSA_mS6_S8_i(%"class.cl::sycl::detail::half_impl::half" addrspace(4)* %add.ptr17.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 - %call19.i = tail call spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_JointMatrixMadINTELIN2cl4sycl6detail9half_impl4halfEfLm2ELm16ELm2ELN5__spv12MatrixLayoutE0ELS6_3ELS6_0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT6_EXT7_EEEPNS9_IT_XT1_EXT2_EXT4_EXT7_EEEPNS9_ISD_XT2_EXT3_EXT5_EXT7_EEESC_S8_(%"struct.__spv::__spirv_JointMatrixINTEL.0" addrspace(4)* %call13.i, %"struct.__spv::__spirv_JointMatrixINTEL.1" addrspace(4)* %call18.i, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %C.0.i, i32 3) #3 + %call18.i = tail call spir_func %spirv.JointMatrixINTEL._half_16_2_3_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIN2cl4sycl6detail9half_impl4halfELm16ELm2ELN5__spv12MatrixLayoutE3ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPSA_mS6_S8_i(%"class.cl::sycl::detail::half_impl::half" addrspace(4)* %add.ptr17.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 + %call19.i = tail call spir_func %spirv.JointMatrixINTEL._float_2_2_0_3 addrspace(4)* @_Z27__spirv_JointMatrixMadINTELIN2cl4sycl6detail9half_impl4halfEfLm2ELm16ELm2ELN5__spv12MatrixLayoutE0ELS6_3ELS6_0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT6_EXT7_EEEPNS9_IT_XT1_EXT2_EXT4_EXT7_EEEPNS9_ISD_XT2_EXT3_EXT5_EXT7_EEESC_S8_(%spirv.JointMatrixINTEL._half_2_16_0_3 addrspace(4)* %call13.i, %spirv.JointMatrixINTEL._half_16_2_3_3 addrspace(4)* %call18.i, %spirv.JointMatrixINTEL._float_2_2_0_3 addrspace(4)* %C.0.i, i32 3) #3 %add.i = add nuw nsw i32 %k.0.i, 16 br label %for.cond.i, !llvm.loop !20 _ZZ4mainENKUlN2cl4sycl7nd_itemILi2EEEE_clES2_.exit: ; preds = %for.cond.i - tail call spir_func void @_Z29__spirv_JointMatrixStoreINTELIfLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(float addrspace(4)* %add.ptr7.i, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* %C.0.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 + tail call spir_func void @_Z29__spirv_JointMatrixStoreINTELIfLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(float addrspace(4)* %add.ptr7.i, %spirv.JointMatrixINTEL._float_2_2_0_3 addrspace(4)* %C.0.i, i64 %_arg_1, i32 0, i32 3, i32 0) #3 ret void } ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIfLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(float addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._float_2_2_0_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIfLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPS5_mS1_S3_i(float addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL.0" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIN2cl4sycl6detail9half_impl4halfELm2ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPSA_mS6_S8_i(%"class.cl::sycl::detail::half_impl::half" addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._half_2_16_0_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIN2cl4sycl6detail9half_impl4halfELm2ELm16ELN5__spv12MatrixLayoutE0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPSA_mS6_S8_i(%"class.cl::sycl::detail::half_impl::half" addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL.1" addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIN2cl4sycl6detail9half_impl4halfELm16ELm2ELN5__spv12MatrixLayoutE3ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPSA_mS6_S8_i(%"class.cl::sycl::detail::half_impl::half" addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._half_16_2_3_3 addrspace(4)* @_Z28__spirv_JointMatrixLoadINTELIN2cl4sycl6detail9half_impl4halfELm16ELm2ELN5__spv12MatrixLayoutE3ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT_XT0_EXT1_EXT2_EXT3_EEEPSA_mS6_S8_i(%"class.cl::sycl::detail::half_impl::half" addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)* @_Z27__spirv_JointMatrixMadINTELIN2cl4sycl6detail9half_impl4halfEfLm2ELm16ELm2ELN5__spv12MatrixLayoutE0ELS6_3ELS6_0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT6_EXT7_EEEPNS9_IT_XT1_EXT2_EXT4_EXT7_EEEPNS9_ISD_XT2_EXT3_EXT5_EXT7_EEESC_S8_(%"struct.__spv::__spirv_JointMatrixINTEL.0" addrspace(4)*, %"struct.__spv::__spirv_JointMatrixINTEL.1" addrspace(4)*, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, i32) local_unnamed_addr #1 +declare dso_local spir_func %spirv.JointMatrixINTEL._float_2_2_0_3 addrspace(4)* @_Z27__spirv_JointMatrixMadINTELIN2cl4sycl6detail9half_impl4halfEfLm2ELm16ELm2ELN5__spv12MatrixLayoutE0ELS6_3ELS6_0ELNS5_5Scope4FlagE3EEPNS5_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT6_EXT7_EEEPNS9_IT_XT1_EXT2_EXT4_EXT7_EEEPNS9_ISD_XT2_EXT3_EXT5_EXT7_EEESC_S8_(%spirv.JointMatrixINTEL._half_2_16_0_3 addrspace(4)*, %spirv.JointMatrixINTEL._half_16_2_3_3 addrspace(4)*, %spirv.JointMatrixINTEL._float_2_2_0_3 addrspace(4)*, i32) local_unnamed_addr #1 ; Function Attrs: convergent -declare dso_local spir_func void @_Z29__spirv_JointMatrixStoreINTELIfLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(float addrspace(4)*, %"struct.__spv::__spirv_JointMatrixINTEL" addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 +declare dso_local spir_func void @_Z29__spirv_JointMatrixStoreINTELIfLm2ELm2ELN5__spv12MatrixLayoutE0ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS4_XT0_EXT1_EXT2_EXT3_EEEmS1_S3_i(float addrspace(4)*, %spirv.JointMatrixINTEL._float_2_2_0_3 addrspace(4)*, i64, i32, i32, i32) local_unnamed_addr #1 ; Function Attrs: inaccessiblememonly nofree nosync nounwind willreturn declare void @llvm.assume(i1 noundef) #2 -- 2.30.2