[SYCL-MLIR] Allow i32 ptr(memref) to be argument of sycl.constructor. (#7455)

During lowering of `atomic` and `multi-ptr`, constructors with pointer
type arguments are being called. Here is an example call generated by
clang:

` call spir_func void
@_ZN4sycl3_V19multi_ptrIjLNS0_6access13address_spaceE1ELNS2_9decoratedE1EEC2EPU3AS1j(%"class.sycl::_V1::multi_ptr"
addrspace(4)* noundef align 8 dereferenceable_or_null(8)
%agg.tmp2.ascast, i32 addrspace(1)* noundef %add.ptr) #8`

As can be seen, the second argument is a `i32` pointer. Currently our
`sycl.constructor` op (that will be replaced during lowering to this
function call) will not allow pointers to be arguments. This PR removes
that restriction.

Co-authored-by: arnamoy.bhattacharyya <[email protected]>

Author: arnamoy10

mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td

@@ -134,6 +134,7 @@ def SYCLMemref : AnyTypeOf<[
   VecMemRef,
 ]>;
 def IndexType : AnyTypeOf<[I32, I64, Index]>;
+def IntMemrefType : AnyTypeOf<[MemRefOf<[I32]>]>;
 def SYCLGetResult : AnyTypeOf<[I64, MemRefOf<[I64]>]>;
 def SYCLGetIDResult : AnyTypeOf<[I64, SYCL_IDType]>;
 def SYCLGetRangeResult : AnyTypeOf<[I64, SYCL_RangeType]>;
@@ -142,7 +143,7 @@ def SYCLGetRangeResult : AnyTypeOf<[I64, SYCL_RangeType]>;
 // CONSTRUCTOR OPERATION
 ////////////////////////////////////////////////////////////////////////////////
 
-def ConstructorArgs : AnyTypeOf<[SYCLMemref, IndexType, SYCL_IDType, SYCL_RangeType]>;
+def ConstructorArgs : AnyTypeOf<[SYCLMemref, IndexType, IntMemrefType, SYCL_IDType, SYCL_RangeType]>;
 def SYCLConstructorOp : SYCL_Op<"constructor", []> {
   let summary = "Generic constructor operation";
   let description = [{
@@ -156,6 +157,8 @@ def SYCLConstructorOp : SYCL_Op<"constructor", []> {
   );
   let results = (outs);
 
+  let hasVerifier = 1;
+
   let assemblyFormat = [{
     `(` $Args `)` attr-dict `:` functional-type($Args, results)
   }];

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

@@ -8,6 +8,7 @@
 
 #include "mlir/Dialect/SYCL/IR/SYCLOps.h"
 
+#include "mlir/Dialect/SYCL/IR/SYCLOpsTypes.h"
 #include "mlir/IR/OpImplementation.h"
 #include "llvm/ADT/TypeSwitch.h"
 
@@ -57,6 +58,15 @@ bool mlir::sycl::SYCLCastOp::areCastCompatible(::mlir::TypeRange Inputs,
   return false;
 }
 
+mlir::LogicalResult mlir::sycl::SYCLConstructorOp::verify() {
+  auto MT = getOperand(0).getType().dyn_cast<mlir::MemRefType>();
+  if (MT && isSYCLType(MT.getElementType()))
+    return success();
+
+  return emitOpError("The first argument of a sycl::constructor op has to be a "
+                     "MemRef to a SYCL type");
+}
+
 mlir::LogicalResult mlir::sycl::SYCLAccessorSubscriptOp::verify() {
   // Available only when: (Dimensions > 0)
   // reference operator[](id<Dimensions> index) const;

mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir

@@ -6,3 +6,10 @@ func.func @AccessorImplDevice(%arg0: memref<?x!sycl.accessor_impl_device<[1], (!
   sycl.constructor(%arg0, %arg1, %arg2, %arg2) {MangledFunctionName = @_ZN4sycl3_V16detail18AccessorImplDeviceILi1EEC1ENS0_2idILi1EEENS0_5rangeILi1EEES7_, TypeName = @AccessorImplDevice} : (memref<?x!sycl.accessor_impl_device<[1], (!sycl.id<1>, !sycl.range<1>, !sycl.range<1>)>>, !sycl.id<1>, !sycl.range<1>, !sycl.range<1>) -> ()
   return
 }
+
+// Ensure integer pointer can be arguments of sycl.constructor.
+// CHECK-LABEL: func.func @TestConstructorII32Ptr
+func.func @TestConstructorII32Ptr(%arg0: memref<?x!sycl.id<1>, 4>, %arg1: memref<?xi32, 1>) {
+  sycl.constructor(%arg0, %arg1) {MangledFunctionName = @_ZN4sycl3_V19multi_ptrIjLNS0_6access13address_spaceE1ELNS2_9decoratedE1EEC1EPU3AS1j, TypeName = @multi_ptr} : (memref<?x!sycl.id<1>, 4>, memref<?xi32, 1>) -> ()
+  return
+}