From f510aa8ae7be679f3593f4d32622b971e536bf49 Mon Sep 17 00:00:00 2001 From: "arnamoy.bhattacharyya" Date: Fri, 18 Nov 2022 15:22:12 -0500 Subject: [PATCH 1/9] [SYCL-MLIR] Allow i32 ptr(memref) to be argument of sycl.constructor. --- mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td | 4 +++- mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir | 7 +++++++ 2 files changed, 10 insertions(+), 1 deletion(-) diff --git a/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td b/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td index c3eea43508735..e243ae6b0ee83 100644 --- a/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td +++ b/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td @@ -118,6 +118,7 @@ def NDItemMemRef : MemRefOf<[SYCL_NdItemType]>; def GroupMemRef : MemRefOf<[SYCL_GroupType]>; def VecMemRef : MemRefOf<[SYCL_VecType]>; + def SYCLMemref : AnyTypeOf<[ IDMemRef, AccessorCommonMemRef, @@ -134,6 +135,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 +144,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 = [{ diff --git a/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir b/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir index 071a858cd2bd9..e4dec7a0a00f9 100644 --- a/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir +++ b/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir @@ -6,3 +6,10 @@ func.func @AccessorImplDevice(%arg0: memref, !sycl.range<1>, !sycl.range<1>)>>, !sycl.id<1>, !sycl.range<1>, !sycl.range<1>) -> () return } + +// Ensure sycl.id and sycl.range types can be arguments of sycl.constructor. +// CHECK-LABEL: func.func @TestConstructorII32Ptr +func.func @TestConstructorII32Ptr(%arg0: memref, 4>, %arg1: memref) { + sycl.constructor(%arg0, %arg1) {MangledFunctionName = @_ZN4sycl3_V19multi_ptrIjLNS0_6access13address_spaceE1ELNS2_9decoratedE1EEC1EPU3AS1j, Type = @multi_ptr} : (memref, 4>, memref) -> () + return +} From 7384caf3a9edead105200c9ea02a727560f7a384 Mon Sep 17 00:00:00 2001 From: "arnamoy.bhattacharyya" Date: Fri, 18 Nov 2022 15:28:08 -0500 Subject: [PATCH 2/9] minor changes --- mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td | 1 - mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir | 2 +- 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td b/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td index e243ae6b0ee83..cb70d64cdb3c8 100644 --- a/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td +++ b/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td @@ -118,7 +118,6 @@ def NDItemMemRef : MemRefOf<[SYCL_NdItemType]>; def GroupMemRef : MemRefOf<[SYCL_GroupType]>; def VecMemRef : MemRefOf<[SYCL_VecType]>; - def SYCLMemref : AnyTypeOf<[ IDMemRef, AccessorCommonMemRef, diff --git a/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir b/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir index e4dec7a0a00f9..5a63a9448f5a9 100644 --- a/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir +++ b/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir @@ -7,7 +7,7 @@ func.func @AccessorImplDevice(%arg0: memref, 4>, %arg1: memref) { sycl.constructor(%arg0, %arg1) {MangledFunctionName = @_ZN4sycl3_V19multi_ptrIjLNS0_6access13address_spaceE1ELNS2_9decoratedE1EEC1EPU3AS1j, Type = @multi_ptr} : (memref, 4>, memref) -> () From 0eb1acdc86ad31d76635629f695dfd821d1b16c0 Mon Sep 17 00:00:00 2001 From: "arnamoy.bhattacharyya" Date: Sat, 19 Nov 2022 08:53:02 -0500 Subject: [PATCH 3/9] Add verifier --- mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td | 2 ++ mlir-sycl/lib/Dialect/IR/SYCLOps.cpp | 6 ++++++ 2 files changed, 8 insertions(+) diff --git a/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td b/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td index cb70d64cdb3c8..0b1cf83c4ba1a 100644 --- a/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td +++ b/mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td @@ -157,6 +157,8 @@ def SYCLConstructorOp : SYCL_Op<"constructor", []> { ); let results = (outs); + let hasVerifier = 1; + let assemblyFormat = [{ `(` $Args `)` attr-dict `:` functional-type($Args, results) }]; diff --git a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp index fc541429d4ddd..228245887f3dd 100644 --- a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp +++ b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp @@ -57,6 +57,12 @@ bool mlir::sycl::SYCLCastOp::areCastCompatible(::mlir::TypeRange Inputs, return false; } +mlir::LogicalResult mlir::sycl::SYCLConstructorOp::verify() { + if (!getOperand(0).getType().dyn_cast()) + return emitOpError("The first argument of a sycl::constructor op has to be a MemRef"); + return success(); +} + mlir::LogicalResult mlir::sycl::SYCLAccessorSubscriptOp::verify() { // Available only when: (Dimensions > 0) // reference operator[](id index) const; From f190b2daca4a3f614711a4e46728f613da61cf05 Mon Sep 17 00:00:00 2001 From: "arnamoy.bhattacharyya" Date: Sat, 19 Nov 2022 08:53:46 -0500 Subject: [PATCH 4/9] clang-format --- mlir-sycl/lib/Dialect/IR/SYCLOps.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp index 228245887f3dd..8969ef1150792 100644 --- a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp +++ b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp @@ -58,8 +58,9 @@ bool mlir::sycl::SYCLCastOp::areCastCompatible(::mlir::TypeRange Inputs, } mlir::LogicalResult mlir::sycl::SYCLConstructorOp::verify() { - if (!getOperand(0).getType().dyn_cast()) - return emitOpError("The first argument of a sycl::constructor op has to be a MemRef"); + if (!getOperand(0).getType().dyn_cast()) + return emitOpError( + "The first argument of a sycl::constructor op has to be a MemRef"); return success(); } From 5151fa7a02a34941a31e66139ff7b0098c8f887b Mon Sep 17 00:00:00 2001 From: "arnamoy.bhattacharyya" Date: Mon, 21 Nov 2022 07:15:30 -0500 Subject: [PATCH 5/9] Make sure that the first arg is a SYCL memref --- mlir-sycl/lib/Dialect/IR/SYCLOps.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp index 8969ef1150792..9a541fd2b75a6 100644 --- a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp +++ b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp @@ -10,6 +10,7 @@ #include "mlir/IR/OpImplementation.h" #include "llvm/ADT/TypeSwitch.h" +#include "mlir/Dialect/SYCL/IR/SYCLOpsTypes.h" bool mlir::sycl::SYCLCastOp::areCastCompatible(::mlir::TypeRange Inputs, ::mlir::TypeRange Outputs) { @@ -58,10 +59,11 @@ bool mlir::sycl::SYCLCastOp::areCastCompatible(::mlir::TypeRange Inputs, } mlir::LogicalResult mlir::sycl::SYCLConstructorOp::verify() { - if (!getOperand(0).getType().dyn_cast()) - return emitOpError( - "The first argument of a sycl::constructor op has to be a MemRef"); - return success(); + if (getOperand(0).getType().dyn_cast() && isSYCLType(getOperand(0).getType().cast().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() { From fb5e720ca8ca9e679536ba65ea997472b3c77489 Mon Sep 17 00:00:00 2001 From: "arnamoy.bhattacharyya" Date: Mon, 21 Nov 2022 07:16:08 -0500 Subject: [PATCH 6/9] clang-format --- mlir-sycl/lib/Dialect/IR/SYCLOps.cpp | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp index 9a541fd2b75a6..2e79efcfa13bc 100644 --- a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp +++ b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp @@ -8,9 +8,9 @@ #include "mlir/Dialect/SYCL/IR/SYCLOps.h" +#include "mlir/Dialect/SYCL/IR/SYCLOpsTypes.h" #include "mlir/IR/OpImplementation.h" #include "llvm/ADT/TypeSwitch.h" -#include "mlir/Dialect/SYCL/IR/SYCLOpsTypes.h" bool mlir::sycl::SYCLCastOp::areCastCompatible(::mlir::TypeRange Inputs, ::mlir::TypeRange Outputs) { @@ -59,11 +59,13 @@ bool mlir::sycl::SYCLCastOp::areCastCompatible(::mlir::TypeRange Inputs, } mlir::LogicalResult mlir::sycl::SYCLConstructorOp::verify() { - if (getOperand(0).getType().dyn_cast() && isSYCLType(getOperand(0).getType().cast().getElementType())) + if (getOperand(0).getType().dyn_cast() && + isSYCLType( + getOperand(0).getType().cast().getElementType())) return success(); - - return emitOpError( - "The first argument of a sycl::constructor op has to be a MemRef to a SYCL type"); + + 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() { From 7d5ee16932f807c22e9a9b09a67fd9cbcf6a174a Mon Sep 17 00:00:00 2001 From: "arnamoy.bhattacharyya" Date: Mon, 21 Nov 2022 09:01:06 -0500 Subject: [PATCH 7/9] Fix test failure --- mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir b/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir index 5a63a9448f5a9..10d1e893e02cd 100644 --- a/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir +++ b/mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir @@ -10,6 +10,6 @@ func.func @AccessorImplDevice(%arg0: memref, 4>, %arg1: memref) { - sycl.constructor(%arg0, %arg1) {MangledFunctionName = @_ZN4sycl3_V19multi_ptrIjLNS0_6access13address_spaceE1ELNS2_9decoratedE1EEC1EPU3AS1j, Type = @multi_ptr} : (memref, 4>, memref) -> () + sycl.constructor(%arg0, %arg1) {MangledFunctionName = @_ZN4sycl3_V19multi_ptrIjLNS0_6access13address_spaceE1ELNS2_9decoratedE1EEC1EPU3AS1j, TypeName = @multi_ptr} : (memref, 4>, memref) -> () return } From fdf120cfd57ec445d01570c6d879596eb401ba67 Mon Sep 17 00:00:00 2001 From: "arnamoy.bhattacharyya" Date: Mon, 21 Nov 2022 09:59:21 -0500 Subject: [PATCH 8/9] Reveiewers comments --- mlir-sycl/lib/Dialect/IR/SYCLOps.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp index 2e79efcfa13bc..2a71403772de3 100644 --- a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp +++ b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp @@ -59,9 +59,10 @@ bool mlir::sycl::SYCLCastOp::areCastCompatible(::mlir::TypeRange Inputs, } mlir::LogicalResult mlir::sycl::SYCLConstructorOp::verify() { - if (getOperand(0).getType().dyn_cast() && + auto MT = getOperand(0).getType().dyn_cast(); + if (MT && isSYCLType( - getOperand(0).getType().cast().getElementType())) + MT.getElementType())) return success(); return emitOpError("The first argument of a sycl::constructor op has to be a " From 7a0ef49cf42011e10ff5dfa3cc20bcf506755f57 Mon Sep 17 00:00:00 2001 From: "arnamoy.bhattacharyya" Date: Mon, 21 Nov 2022 09:59:53 -0500 Subject: [PATCH 9/9] clang-format --- mlir-sycl/lib/Dialect/IR/SYCLOps.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp index 2a71403772de3..e8b26c70bc3df 100644 --- a/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp +++ b/mlir-sycl/lib/Dialect/IR/SYCLOps.cpp @@ -60,9 +60,7 @@ bool mlir::sycl::SYCLCastOp::areCastCompatible(::mlir::TypeRange Inputs, mlir::LogicalResult mlir::sycl::SYCLConstructorOp::verify() { auto MT = getOperand(0).getType().dyn_cast(); - if (MT && - isSYCLType( - MT.getElementType())) + if (MT && isSYCLType(MT.getElementType())) return success(); return emitOpError("The first argument of a sycl::constructor op has to be a "