Replace illegal address space casts asserts with errors

MrSidims · vmaksimo · commit 8dca6b92367a · 2020-12-07T18:04:00.000+03:00
Purpose of this patch is to make the translator a little bit more
user-friendly. For example following SYCL code:
  SYCL_EXTERNAL sycl::global_ptr&lt;int&gt; kernel_fun(int *ptr) {
    return sycl::global_ptr&lt;int&gt;((int*)&amp;ptr);
  }

is UB since it has an implicit cast from private address space to
global. It's hard to detect such UB in FE, hence this diagnostic is
moved to the translator (which already had it, but in form of an
assert).

Signed-off-by: Dmitry Sidorov &lt;dmitry.sidorov@intel.com&gt;
diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
@@ -888,18 +888,22 @@ SPIRV::SPIRVInstruction *LLVMToSPIRV::transUnaryInst(UnaryInstruction *U,
     const auto SrcAddrSpace = Cast->getSrcTy()->getPointerAddressSpace();
     const auto DestAddrSpace = Cast->getDestTy()->getPointerAddressSpace();
     if (DestAddrSpace == SPIRAS_Generic) {
-      assert(SrcAddrSpace != SPIRAS_Constant &&
-             "Casts from constant address space to generic are illegal");
+      getErrorLog().checkError(
+          SrcAddrSpace != SPIRAS_Constant, SPIRVEC_InvalidModule,
+          "Casts from constant address space to generic are illegal\n" +
+              toString(U));
       BOC = OpPtrCastToGeneric;
       // In SPIR-V only casts to/from generic are allowed. But with
       // SPV_INTEL_usm_storage_classes we can also have casts from global_device
       // and global_host to global addr space and vice versa.
     } else if (SrcAddrSpace == SPIRAS_GlobalDevice ||
                SrcAddrSpace == SPIRAS_GlobalHost) {
-      assert(
-          (DestAddrSpace == SPIRAS_Global || DestAddrSpace == SPIRAS_Generic) &&
-          "Casts from global_device/global_host only allowed to \
-             global/generic");
+      getErrorLog().checkError(DestAddrSpace == SPIRAS_Global ||
+                                   DestAddrSpace == SPIRAS_Generic,
+                               SPIRVEC_InvalidModule,
+                               "Casts from global_device/global_host only "
+                               "allowed to global/generic\n" +
+                                   toString(U));
       if (!BM->isAllowedToUseExtension(
               ExtensionID::SPV_INTEL_usm_storage_classes)) {
         if (DestAddrSpace == SPIRAS_Global)
@@ -910,10 +914,12 @@ SPIRV::SPIRVInstruction *LLVMToSPIRV::transUnaryInst(UnaryInstruction *U,
       }
     } else if (DestAddrSpace == SPIRAS_GlobalDevice ||
                DestAddrSpace == SPIRAS_GlobalHost) {
-      assert(
-          (SrcAddrSpace == SPIRAS_Global || SrcAddrSpace == SPIRAS_Generic) &&
-          "Casts to global_device/global_host only allowed from \
-             global/generic");
+      getErrorLog().checkError(SrcAddrSpace == SPIRAS_Global ||
+                                   SrcAddrSpace == SPIRAS_Generic,
+                               SPIRVEC_InvalidModule,
+                               "Casts to global_device/global_host only "
+                               "allowed from global/generic\n" +
+                                   toString(U));
       if (!BM->isAllowedToUseExtension(
               ExtensionID::SPV_INTEL_usm_storage_classes)) {
         if (SrcAddrSpace == SPIRAS_Global)
@@ -923,9 +929,15 @@ SPIRV::SPIRVInstruction *LLVMToSPIRV::transUnaryInst(UnaryInstruction *U,
         BOC = OpCrossWorkgroupCastToPtrINTEL;
       }
     } else {
-      assert(DestAddrSpace != SPIRAS_Constant &&
-             "Casts from generic address space to constant are illegal");
-      assert(SrcAddrSpace == SPIRAS_Generic);
+      getErrorLog().checkError(
+          SrcAddrSpace == SPIRAS_Generic, SPIRVEC_InvalidModule,
+          "Casts from private/local/global address space are allowed only to "
+          "generic\n" +
+              toString(U));
+      getErrorLog().checkError(
+          DestAddrSpace != SPIRAS_Constant, SPIRVEC_InvalidModule,
+          "Casts from generic address space to constant are illegal\n" +
+              toString(U));
       BOC = OpGenericCastToPtr;
     }
   } else {
diff --git a/llvm-spirv/test/negative/invalid-constant-generic-cast.ll b/llvm-spirv/test/negative/invalid-constant-generic-cast.ll
@@ -0,0 +1,30 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: not llvm-spirv %t.bc -o %t.spv 2>&1 | FileCheck %s
+
+; CHECK: InvalidModule: Invalid SPIR-V module: Casts from generic address space to constant are illegal
+
+source_filename = "sample.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-sycldevice"
+
+; Function Attrs: nofree norecurse nounwind writeonly
+define dso_local spir_func void @_Z10kernel_funPi(i32 addrspace(4)* %ptr) {
+entry:
+  %0 = addrspacecast i32 addrspace(4)* %ptr to i32 addrspace(2)*
+  ret void
+}
+
+!llvm.module.flags = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{!"clang version 12.0.0"}
+!4 = !{!5, !6, i64 0}
+!5 = !{!"_ZTSN2cl4sycl9multi_ptrIiLNS0_6access13address_spaceE1EEE", !6, i64 0}
+!6 = !{!"any pointer", !7, i64 0}
+!7 = !{!"omnipotent char", !8, i64 0}
+!8 = !{!"Simple C++ TBAA"}
diff --git a/llvm-spirv/test/negative/invalid-device-local-cast.ll b/llvm-spirv/test/negative/invalid-device-local-cast.ll
@@ -0,0 +1,30 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: not llvm-spirv %t.bc -o %t.spv 2>&1 | FileCheck %s
+
+; CHECK: InvalidModule: Invalid SPIR-V module: Casts from global_device/global_host only allowed to global/generic 
+
+source_filename = "sample.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-sycldevice"
+
+; Function Attrs: nofree norecurse nounwind writeonly
+define dso_local spir_func void @_Z10kernel_funPi(i32 addrspace(5)* %ptr) {
+entry:
+  %0 = addrspacecast i32 addrspace(5)* %ptr to i32 addrspace(3)*
+  ret void
+}
+
+!llvm.module.flags = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{!"clang version 12.0.0"}
+!4 = !{!5, !6, i64 0}
+!5 = !{!"_ZTSN2cl4sycl9multi_ptrIiLNS0_6access13address_spaceE1EEE", !6, i64 0}
+!6 = !{!"any pointer", !7, i64 0}
+!7 = !{!"omnipotent char", !8, i64 0}
+!8 = !{!"Simple C++ TBAA"}
diff --git a/llvm-spirv/test/negative/invalid-private-global-cast.ll b/llvm-spirv/test/negative/invalid-private-global-cast.ll
@@ -0,0 +1,30 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: not llvm-spirv %t.bc -o %t.spv 2>&1 | FileCheck %s
+
+; CHECK: InvalidModule: Invalid SPIR-V module: Casts from private/local/global address space are allowed only to generic
+
+source_filename = "sample.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-sycldevice"
+
+; Function Attrs: nofree norecurse nounwind writeonly
+define dso_local spir_func void @_Z10kernel_funPi(i32 addrspace(1)* %ptr) {
+entry:
+  %0 = addrspacecast i32 addrspace(1)* %ptr to i32*
+  ret void
+}
+
+!llvm.module.flags = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{!"clang version 12.0.0"}
+!4 = !{!5, !6, i64 0}
+!5 = !{!"_ZTSN2cl4sycl9multi_ptrIiLNS0_6access13address_spaceE1EEE", !6, i64 0}
+!6 = !{!"any pointer", !7, i64 0}
+!7 = !{!"omnipotent char", !8, i64 0}
+!8 = !{!"Simple C++ TBAA"}
