Support SYCL 2020 namespace (::sycl) for SYCL Half/BFloat16 types

In SYCL 2020 the "::cl::sycl" has been changed to "::sycl". Update SYCL
half type support to account for that.

Author: aelovikov-intel

lib/SPIRV/SPIRVUtil.cpp

@@ -288,7 +288,7 @@ bool isSYCLHalfType(llvm::Type *Ty) {
       return false;
     StringRef Name = ST->getName();
     Name.consume_front("class.");
-    if ((Name.startswith("cl::sycl::") ||
+    if ((Name.startswith("sycl::") || Name.startswith("cl::sycl::") ||
          Name.startswith("__sycl_internal::")) &&
         Name.endswith("::half")) {
       return true;
@@ -303,7 +303,7 @@ bool isSYCLBfloat16Type(llvm::Type *Ty) {
       return false;
     StringRef Name = ST->getName();
     Name.consume_front("class.");
-    if ((Name.startswith("cl::sycl::") ||
+    if ((Name.startswith("sycl::") || Name.startswith("cl::sycl::") ||
          Name.startswith("__sycl_internal::")) &&
         Name.endswith("::bfloat16")) {
       return true;

test/transcoding/SPV_INTEL_joint_matrix/sycl_2020_namespace.ll

@@ -0,0 +1,24 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-ext=+SPV_INTEL_joint_matrix -o %t.spv
+; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s
+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"
+
+; Ensure that ::sycl::_V1.*{half|bfloat16} are parsed as SYCL types.
+
+; CHECK-DAG: TypeFloat [[#HalfTy:]] 16
+; CHECK-DAG: TypeInt [[#BFloat16Ty:]] 16
+
+%"class.sycl::_V1::anything::half" = type { half }
+%"class.sycl::_V1::anything::bfloat16" = type { i16 }
+
+%"struct.__spv::__spirv_JointMatrixINTEL.half" = type { [2 x [2 x [1 x [4 x %"class.sycl::_V1::anything::half"]]]]* }
+%"struct.__spv::__spirv_JointMatrixINTEL.bfloat16" = type { [2 x [2 x [1 x [4 x %"class.sycl::_V1::anything::bfloat16"]]]]* }
+
+define spir_func void @foo(%"struct.__spv::__spirv_JointMatrixINTEL.half" *) {
+  ret void
+}
+
+define spir_func void @bar(%"struct.__spv::__spirv_JointMatrixINTEL.bfloat16" *) {
+  ret void
+}