[SYCL-MLIR] Opaque pointer support in SYCL-to-LLVM conversion

mlir-sycl/include/mlir/Conversion/SYCLPasses.td

@@ -30,8 +30,11 @@ def ConvertSYCLToLLVM : Pass<"convert-sycl-to-llvm", "ModuleOp"> {
   ];
   let options = [
     Option<"indexBitwidth", "index-bitwidth", "unsigned",
-           /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
-           "Bitwidth of the index type, 0 to use size of machine word">
+            /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
+            "Bitwidth of the index type, 0 to use size of machine word">,
+    Option<"useOpaquePointers", "use-opaque-pointers", "bool",
+            /*default=*/"false", "Generate LLVM IR using opaque pointers "
+            "instead of typed pointers">,
   ];
 }
 

mlir-sycl/include/mlir/Conversion/SYCLToLLVM/DialectBuilder.h

@@ -72,7 +72,8 @@ class LLVMBuilder : public DialectBuilder {
 public:
   LLVMBuilder(OpBuilder &b, Location loc) : DialectBuilder(b, loc) {}
 
-  LLVM::AllocaOp genAlloca(Type type, Value size, int64_t align) const;
+  LLVM::AllocaOp genAlloca(Type type, Type elemType, Value size,
+                           int64_t align) const;
   LLVM::BitcastOp genBitcast(Type type, Value val) const;
   LLVM::ExtractValueOp genExtractValue(Type type, Value container,
                                        ArrayRef<int64_t> pos) const;

mlir-sycl/lib/Conversion/SYCLToLLVM/DialectBuilder.cpp

@@ -103,9 +103,9 @@ func::CallOp FuncBuilder::genCall(StringRef funcName, TypeRange resTypes,
 // LLVMBuilder
 //===----------------------------------------------------------------------===//
 
-LLVM::AllocaOp LLVMBuilder::genAlloca(Type type, Value size,
+LLVM::AllocaOp LLVMBuilder::genAlloca(Type type, Type elemType, Value size,
                                       int64_t align) const {
-  return create<LLVM::AllocaOp>(type, size, align);
+  return create<LLVM::AllocaOp>(type, elemType, size, align);
 }
 
 LLVM::BitcastOp LLVMBuilder::genBitcast(Type type, Value val) const {

mlir-sycl/lib/Dialect/IR/SYCLOps.cpp

@@ -146,6 +146,10 @@ LogicalResult SYCLAccessorSubscriptOp::verify() {
         .Case<MemRefType>(
             [&](auto Ty) { return VerifyElemType(Ty.getElementType()); })
         .Case<LLVM::LLVMPointerType>([&](auto Ty) {
+          if (!Ty.getElementType()) {
+            // With opaque pointers, there is no element type to inspect.
+            return success();
+          }
           const Type ElemType = Ty.getElementType();
           return (!isa<LLVM::LLVMStructType>(ElemType))
                      ? emitOpError(

mlir-sycl/test/Conversion/SYCLToLLVM/sycl-addrspacecast-to-llvm-typed-pointer.mlir

@@ -0,0 +1,23 @@
+// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm='use-opaque-pointers=0' -verify-diagnostics %s | FileCheck %s
+
+// CHECK-LABEL: llvm.func @PtrCastToGeneric(%arg0: !llvm.ptr<i32>) -> !llvm.ptr<i32, 4> {
+// CHECK-NEXT:    %0 = llvm.addrspacecast %arg0 : !llvm.ptr<i32> to !llvm.ptr<i32, 4>
+// CHECK-NEXT:    llvm.return %0 : !llvm.ptr<i32, 4>
+// CHECK-NEXT:  }
+
+func.func @PtrCastToGeneric(%arg0: memref<?xi32>) -> memref<?xi32, 4> {
+  %0 = sycl.addrspacecast %arg0 : memref<?xi32> to memref<?xi32, 4>
+  return %0 : memref<?xi32, 4>
+}
+
+// -----
+
+// CHECK-LABEL: llvm.func @GenericCastToPtr(%arg0: !llvm.ptr<i32, 4>) -> !llvm.ptr<i32> {
+// CHECK-NEXT:    %0 = llvm.addrspacecast %arg0 : !llvm.ptr<i32, 4> to !llvm.ptr<i32>
+// CHECK-NEXT:    llvm.return %0 : !llvm.ptr<i32>
+// CHECK-NEXT:  }
+
+func.func @GenericCastToPtr(%arg0: memref<?xi32, 4>) -> memref<?xi32> {
+  %0 = sycl.addrspacecast %arg0 : memref<?xi32, 4> to memref<?xi32>
+  return %0 : memref<?xi32>
+}

mlir-sycl/test/Conversion/SYCLToLLVM/sycl-addrspacecast-to-llvm.mlir

@@ -1,8 +1,8 @@
-// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm -verify-diagnostics %s | FileCheck %s
+// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm='use-opaque-pointers=1' -verify-diagnostics %s | FileCheck %s
 
-// CHECK-LABEL: llvm.func @PtrCastToGeneric(%arg0: !llvm.ptr<i32>) -> !llvm.ptr<i32, 4> {
-// CHECK-NEXT:    %0 = llvm.addrspacecast %arg0 : !llvm.ptr<i32> to !llvm.ptr<i32, 4>
-// CHECK-NEXT:    llvm.return %0 : !llvm.ptr<i32, 4>
+// CHECK-LABEL: llvm.func @PtrCastToGeneric(%arg0: !llvm.ptr) -> !llvm.ptr<4> {
+// CHECK-NEXT:    %0 = llvm.addrspacecast %arg0 : !llvm.ptr to !llvm.ptr<4>
+// CHECK-NEXT:    llvm.return %0 : !llvm.ptr<4>
 // CHECK-NEXT:  }
 
 func.func @PtrCastToGeneric(%arg0: memref<?xi32>) -> memref<?xi32, 4> {
@@ -12,9 +12,9 @@ func.func @PtrCastToGeneric(%arg0: memref<?xi32>) -> memref<?xi32, 4> {
 
 // -----
 
-// CHECK-LABEL: llvm.func @GenericCastToPtr(%arg0: !llvm.ptr<i32, 4>) -> !llvm.ptr<i32> {
-// CHECK-NEXT:    %0 = llvm.addrspacecast %arg0 : !llvm.ptr<i32, 4> to !llvm.ptr<i32>
-// CHECK-NEXT:    llvm.return %0 : !llvm.ptr<i32>
+// CHECK-LABEL: llvm.func @GenericCastToPtr(%arg0: !llvm.ptr<4>) -> !llvm.ptr {
+// CHECK-NEXT:    %0 = llvm.addrspacecast %arg0 : !llvm.ptr<4> to !llvm.ptr
+// CHECK-NEXT:    llvm.return %0 : !llvm.ptr
 // CHECK-NEXT:  }
 
 func.func @GenericCastToPtr(%arg0: memref<?xi32, 4>) -> memref<?xi32> {

mlir-sycl/test/Conversion/SYCLToLLVM/sycl-call-to-llvm-typed-pointer.mlir

@@ -0,0 +1,38 @@
+// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm='use-opaque-pointers=0' -verify-diagnostics %s | FileCheck %s
+
+//===-------------------------------------------------------------------------------------------------===//
+// sycl.call with non void return type
+//===-------------------------------------------------------------------------------------------------===//
+
+// CHECK: llvm.func @foo() -> [[RET_TYPE:i32]]
+// CHECK: llvm.func @test() -> [[RET_TYPE]] {
+// CHECK-NEXT:  %0 = llvm.call @foo() : () -> [[RET_TYPE]]
+// CHECK-NEXT:  llvm.return %0 : [[RET_TYPE]]
+
+func.func private @foo() -> (i32)
+
+func.func @test() -> (i32) {
+  %0 = sycl.call @foo() {MangledFunctionName = @foo, TypeName = @accessor} : () -> i32
+  return %0 : i32
+}
+
+// -----
+
+//===-------------------------------------------------------------------------------------------------===//
+// Member functions for sycl::accessor
+//===-------------------------------------------------------------------------------------------------===//
+
+!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
+!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
+!sycl_accessor_1_i32_rw_gb = !sycl.accessor<[1, i32, read_write, global_buffer], (!sycl.accessor_impl_device<[1], (!sycl_id_1_, !sycl_range_1_, !sycl_range_1_)>, !llvm.struct<(ptr<i32, 1>)>)>
+
+// CHECK: llvm.func @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE([[ARG_TYPES:!llvm.ptr<struct<"class.sycl::_V1::accessor.1",.*]])
+func.func private @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE(memref<?x!sycl_accessor_1_i32_rw_gb>, memref<?xi32>, !sycl_range_1_, !sycl_range_1_, !sycl_id_1_)
+
+func.func @accessorInit1(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>, %arg1: memref<?xi32>, %arg2: !sycl_range_1_, %arg3: !sycl_range_1_, %arg4: !sycl_id_1_) {
+  // CHECK: llvm.call @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE({{.*}}) : ([[ARG_TYPES]]) -> ()
+  sycl.call @__init(%arg0, %arg1, %arg2, %arg3, %arg4) {MangledFunctionName = @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE, TypeName = @accessor} : (memref<?x!sycl_accessor_1_i32_rw_gb>, memref<?xi32>, !sycl_range_1_, !sycl_range_1_, !sycl_id_1_) -> ()
+  return
+}
+
+// -----

mlir-sycl/test/Conversion/SYCLToLLVM/sycl-call-to-llvm.mlir

@@ -1,4 +1,4 @@
-// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm -verify-diagnostics %s | FileCheck %s
+// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm='use-opaque-pointers=1' -verify-diagnostics %s | FileCheck %s
 
 //===-------------------------------------------------------------------------------------------------===//
 // sycl.call with non void return type
@@ -24,9 +24,9 @@ func.func @test() -> (i32) {
 
 !sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
 !sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
-!sycl_accessor_1_i32_rw_gb = !sycl.accessor<[1, i32, read_write, global_buffer], (!sycl.accessor_impl_device<[1], (!sycl_id_1_, !sycl_range_1_, !sycl_range_1_)>, !llvm.struct<(ptr<i32, 1>)>)>
+!sycl_accessor_1_i32_rw_gb = !sycl.accessor<[1, i32, read_write, global_buffer], (!sycl.accessor_impl_device<[1], (!sycl_id_1_, !sycl_range_1_, !sycl_range_1_)>, !llvm.struct<(ptr<1>)>)>
 
-// CHECK: llvm.func @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE([[ARG_TYPES:!llvm.ptr<struct<"class.sycl::_V1::accessor.1",.*]])
+// CHECK: llvm.func @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE([[ARG_TYPES:!llvm.ptr,.*]])
 func.func private @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE(memref<?x!sycl_accessor_1_i32_rw_gb>, memref<?xi32>, !sycl_range_1_, !sycl_range_1_, !sycl_id_1_)
 
 func.func @accessorInit1(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>, %arg1: memref<?xi32>, %arg2: !sycl_range_1_, %arg3: !sycl_range_1_, %arg4: !sycl_id_1_) {

mlir-sycl/test/Conversion/SYCLToLLVM/sycl-cast-to-llvm-typed-pointer.mlir

@@ -0,0 +1,92 @@
+// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm='use-opaque-pointers=0' -verify-diagnostics %s | FileCheck %s
+
+!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
+!sycl_range_1_ = !sycl.range<[1], (!sycl_array_1_)>
+func.func @cast_sycl_range_to_array(%arg0: memref<?x!sycl_range_1_>) -> memref<?x!sycl_array_1_> {
+  // CHECK-LABEL: llvm.func @cast_sycl_range_to_array(
+  // CHECK-SAME:                                      [[SRC:%.*]]: !llvm.ptr<[[RANGE1:.*]]>) -> !llvm.ptr<[[ARRAY1:.*]]>
+  // CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[RANGE1]]> to !llvm.ptr<[[ARRAY1]]>
+  // CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[ARRAY1]]>
+
+  %0 = "sycl.cast"(%arg0) : (memref<?x!sycl_range_1_>) -> memref<?x!sycl_array_1_>
+  func.return %0 : memref<?x!sycl_array_1_>
+}
+
+// -----
+
+!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
+!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
+func.func @cast_sycl_id_to_array(%arg0: memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_> {
+  // CHECK-LABEL: llvm.func @cast_sycl_id_to_array(
+  // CHECK-SAME:                                   [[SRC:%.*]]: !llvm.ptr<[[ID1:.*]]>) -> !llvm.ptr<[[ARRAY1]]>
+  // CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[ID1]]> to !llvm.ptr<[[ARRAY1]]>
+  // CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[ARRAY1]]>
+
+  %0 = "sycl.cast"(%arg0) : (memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_>
+  func.return %0: memref<?x!sycl_array_1_>
+}
+
+// -----
+
+!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
+!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
+!sycl_accessor_1_i32_rw_gb = !sycl.accessor<[1, i32, read_write, global_buffer], (!sycl.accessor_impl_device<[1], (!sycl_id_1_, !sycl_range_1_, !sycl_range_1_)>, !llvm.struct<(ptr<i32, 1>)>)>
+func.func @cast_sycl_accessor_to_accessor_common(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl.accessor_common> {
+  // CHECK-LABEL: llvm.func @cast_sycl_accessor_to_accessor_common(
+  // CHECK-SAME:                                                   [[SRC:%.*]]: !llvm.ptr<[[ACC1:.*]]>) -> !llvm.ptr<[[COMMON:.*]]>
+  // CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[ACC1]]> to !llvm.ptr<[[COMMON]]>
+  // CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[COMMON]]>
+
+  %0 = "sycl.cast"(%arg0) : (memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl.accessor_common>
+  func.return %0: memref<?x!sycl.accessor_common>
+}
+
+!sycl_LocalAccessorBaseDevice_1_ = !sycl.LocalAccessorBaseDevice<[1], (!sycl_range_1_, !sycl_range_1_, !sycl_id_1_)>
+!sycl_local_accessor_base_1_i32_rw = !sycl.local_accessor_base<[1, i32, read_write], (!sycl_LocalAccessorBaseDevice_1_, memref<?xi32, 3>)>
+func.func @cast_sycl_accessor_to_local_accessor_base(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl_local_accessor_base_1_i32_rw> {
+  // CHECK-LABEL: llvm.func @cast_sycl_accessor_to_local_accessor_base(
+  // CHECK-SAME:                                                       [[SRC:%.*]]: !llvm.ptr<[[ACC1]]>) -> !llvm.ptr<[[LOCALBASE:.*]]>
+  // CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[ACC1]]> to !llvm.ptr<[[LOCALBASE]]>
+  // CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[LOCALBASE]]>
+
+  %0 = "sycl.cast"(%arg0) : (memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl_local_accessor_base_1_i32_rw>
+  func.return %0: memref<?x!sycl_local_accessor_base_1_i32_rw>
+}
+
+func.func @cast_sycl_accessor_to_owner_less_base(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl.owner_less_base> {
+  // CHECK-LABEL: llvm.func @cast_sycl_accessor_to_owner_less_base(
+  // CHECK-SAME:                                                   [[SRC:%.*]]: !llvm.ptr<[[ACC1]]>) -> !llvm.ptr<[[OWNERLESSBASE:.*]]>
+  // CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[ACC1]]> to !llvm.ptr<[[OWNERLESSBASE]]>
+  // CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[OWNERLESSBASE]]>
+
+  %0 = "sycl.cast"(%arg0) : (memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl.owner_less_base>
+  func.return %0: memref<?x!sycl.owner_less_base>
+}
+
+// -----
+
+!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
+!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
+!sycl_LocalAccessorBaseDevice_1_ = !sycl.LocalAccessorBaseDevice<[1], (!sycl_range_1_, !sycl_range_1_, !sycl_id_1_)>
+!sycl_local_accessor_base_1_i32_rw = !sycl.local_accessor_base<[1, i32, read_write], (!sycl_LocalAccessorBaseDevice_1_, memref<?xi32, 3>)>
+func.func @cast_sycl_local_accessor_base_to_accessor_common(%arg0: memref<?x!sycl_local_accessor_base_1_i32_rw>) -> memref<?x!sycl.accessor_common> {
+  // CHECK-LABEL: llvm.func @cast_sycl_local_accessor_base_to_accessor_common(
+  // CHECK-SAME:                                                              [[SRC:%.*]]: !llvm.ptr<[[LAB1:.*]]>) -> !llvm.ptr<[[COMMON]]>
+  // CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[LAB1]]> to !llvm.ptr<[[COMMON]]>
+  // CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[COMMON]]
+  %0 = "sycl.cast"(%arg0) : (memref<?x!sycl_local_accessor_base_1_i32_rw>) -> memref<?x!sycl.accessor_common>
+  func.return %0: memref<?x!sycl.accessor_common>
+}
+
+!sycl_local_accessor_1_i32_rw = !sycl.local_accessor<[1, i32], (!sycl_local_accessor_base_1_i32_rw)>
+func.func @cast_sycl_local_accessor_to_local_accessor_base(%arg0: memref<?x!sycl_local_accessor_1_i32_rw>) -> memref<?x!sycl_local_accessor_base_1_i32_rw> {
+  // CHECK-LABEL: llvm.func @cast_sycl_local_accessor_to_local_accessor_base(
+  // CHECK-SAME:                                                             [[SRC:%.*]]: !llvm.ptr<[[LA1:.*]]>) -> !llvm.ptr<[[LAB1]]>
+  // CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[LA1]]> to !llvm.ptr<[[LAB1]]>
+  // CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[LAB1]]
+
+  %0 = "sycl.cast"(%arg0) : (memref<?x!sycl_local_accessor_1_i32_rw>) -> memref<?x!sycl_local_accessor_base_1_i32_rw>
+  func.return %0: memref<?x!sycl_local_accessor_base_1_i32_rw>  
+}
+
+