Resolve type mismatch when calling ushort OCL built-ins for bfloat16 types

Resolve "Calling a function with a bad signature!" assertion in debug
builds caused by passing %"class.sycl::_V1::ext::oneapi::bfloat16" type parameter
to built-ins, which accept i16 in SubGroupFuncsResolution pass.

Author: YuriPlyakhin

IGC/Compiler/Optimizer/OpenCLPasses/SubGroupFuncs/SubGroupFuncsResolution.cpp

@@ -333,6 +333,22 @@ void SubGroupFuncsResolution::mediaBlockWrite(llvm::CallInst& CI)
     CI.eraseFromParent();
 }
 
+// If CI parameter is %"class.sycl::_V1::ext::oneapi::bfloat16" type, which is { i16 },
+// then we need to cast it to i16 type before calling simdBlockRead intrinsic.
+static inline Value* castSYCLBFloat16toi16(PointerType* PtrTy, Value* Ptr, CallInst& CI, LLVMContext& C)
+{
+    if (StructType* ST = dyn_cast<StructType>(IGCLLVM::getNonOpaquePtrEltTy(PtrTy)))
+    {
+        // check if ST has only field and this field is i16 type
+        if (ST->getNumElements() == 1 && ST->getElementType(0)->isIntegerTy(16))
+        {
+            return CastInst::CreatePointerCast(Ptr, PointerType::get(Type::getInt16Ty(C), PtrTy->getAddressSpace()), "", &CI);
+        }
+    }
+
+    return Ptr;
+}
+
 void SubGroupFuncsResolution::simdBlockRead(llvm::CallInst& CI, bool hasCacheControls)
 {
     // Creates intrinsics that will be lowered in the CodeGen and will handle the simd_block_read
@@ -341,7 +357,7 @@ void SubGroupFuncsResolution::simdBlockRead(llvm::CallInst& CI, bool hasCacheCon
     PointerType* PtrTy = dyn_cast<PointerType>(Ptr->getType());
     IGC_ASSERT_MESSAGE(PtrTy, "simdBlockRead has non-pointer type!");
     SmallVector<Value*, 1> args;
-    args.push_back(Ptr);
+    args.push_back(castSYCLBFloat16toi16(PtrTy, Ptr, CI, C));
     SmallVector<Type*, 3>  types;
     types.push_back(nullptr); types.push_back(nullptr);
     GenISAIntrinsic::ID  genIntrinID = GenISAIntrinsic::GenISA_simdBlockRead;
@@ -447,7 +463,7 @@ void SubGroupFuncsResolution::simdBlockWrite(llvm::CallInst& CI, bool hasCacheCo
     SmallVector<Type*, 2>  types;
     Value* dataArg = CI.getArgOperand(1);
 
-    args.push_back(CI.getArgOperand(0));
+    args.push_back(castSYCLBFloat16toi16(PtrTy, Ptr, CI, C));
     args.push_back(dataArg);
 
     switch (dataArg->getType()->getScalarType()->getScalarSizeInBits())

IGC/Compiler/tests/SubGroupFuncsResolution/simd_blockreadwrite_bfloat16.ll

@@ -0,0 +1,52 @@
+;=========================== begin_copyright_notice ============================
+;
+; Copyright (C) 2024 Intel Corporation
+;
+; SPDX-License-Identifier: MIT
+;
+;============================ end_copyright_notice =============================
+;
+; RUN: igc_opt --platformdg2 --igc-sub-group-func-resolution -S %s 2>&1 | FileCheck %s
+; ------------------------------------------------
+; SubGroupFuncsResolution
+; ------------------------------------------------
+; This test checks that SubGroupFuncsResolution pass resolves mismatch
+; between bfloat16 type passed from SYCL and built-ins accepting i16 type
+; ------------------------------------------------
+
+%"class.sycl::_V1::ext::oneapi::bfloat16" = type { i16 }
+
+define spir_kernel void @test_bfloat16(%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(3)* %dst, %"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(1)* %src) #0 {
+; CHECK-LABEL: @test_bfloat16(
+; CHECK-NEXT:  entry:
+; CHECK-NEXT:    [[TMP0:%.*]] = bitcast %"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(1)* [[SRC:%.*]] to i16 addrspace(1)*
+; CHECK-NEXT:    [[TMP1:%.*]] = call <2 x i16> @llvm.genx.GenISA.simdBlockRead.v2i16.p1i16(i16 addrspace(1)* [[TMP0]])
+; CHECK-NEXT:    [[TMP2:%.*]] = bitcast %"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(3)* [[DST:%.*]] to i16 addrspace(3)*
+; CHECK-NEXT:    call void @llvm.genx.GenISA.simdBlockWrite.p3i16.v2i16(i16 addrspace(3)* [[TMP2]], <2 x i16> [[TMP1]])
+; CHECK-NEXT:    [[TMP3:%.*]] = bitcast %"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(1)* [[SRC]] to i16 addrspace(1)*
+; CHECK-NEXT:    [[TMP4:%.*]] = call <16 x i16> @llvm.genx.GenISA.simdBlockRead.v16i16.p1i16(i16 addrspace(1)* [[TMP3]])
+; CHECK-NEXT:    [[TMP5:%.*]] = bitcast %"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(3)* [[DST]] to i16 addrspace(3)*
+; CHECK-NEXT:    call void @llvm.genx.GenISA.simdBlockWrite.p3i16.v16i16(i16 addrspace(3)* [[TMP5]], <16 x i16> [[TMP4]])
+; CHECK-NEXT:    ret void
+;
+entry:
+  %0 = call spir_func <2 x i16> @__builtin_IB_simd_block_read_2_global_h(%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(1)* %src) #0
+  call spir_func void @__builtin_IB_simd_block_write_2_local_h(%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(3)* %dst, <2 x i16> %0) #0
+  %1 = call spir_func <16 x i16> @__builtin_IB_simd_block_read_16_global_h(%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(1)* %src) #0
+  call spir_func void @__builtin_IB_simd_block_write_16_local_h(%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(3)* %dst, <16 x i16> %1) #0
+  ret void
+}
+
+declare spir_func <2 x i16> @__builtin_IB_simd_block_read_2_global_h(%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(1)*) #0
+declare spir_func void @__builtin_IB_simd_block_write_2_local_h(%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(3)*, <2 x i16>) #0
+declare spir_func <16 x i16> @__builtin_IB_simd_block_read_16_global_h(%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(1)*) #0
+declare spir_func void @__builtin_IB_simd_block_write_16_local_h(%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(3)*, <16 x i16>) #0
+
+attributes #0 = { convergent noinline nounwind optnone }
+
+!igc.functions = !{!3}
+
+!3 = !{void (%"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(3)*, %"class.sycl::_V1::ext::oneapi::bfloat16" addrspace(1)*)* @test_bfloat16, !4}
+!4 = !{!5, !6}
+!5 = !{!"function_type", i32 0}
+!6 = !{!"sub_group_size", i32 8}