-
Notifications
You must be signed in to change notification settings - Fork 15.3k
[MLIR][NVVM][NFC] Re-order mem_scope and shared_space attrs #168348
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
[MLIR][NVVM][NFC] Re-order mem_scope and shared_space attrs #168348
Conversation
The mbarrier Ops also require access to the mem_scope and shared_space attributes. To maintain consistency, this patch moves their definitions to the beginning of the file alongside the other attribute definitions. Signed-off-by: Durgadoss R <durgadossr@nvidia.com>
|
@llvm/pr-subscribers-mlir-llvm Author: Durgadoss R (durga4github) ChangesThe mbarrier Ops also require access to the mem_scope and shared_space attributes. Full diff: https://github.com/llvm/llvm-project/pull/168348.diff 1 Files Affected:
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index d4ef5104d3c1f..456d816205b58 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -228,6 +228,33 @@ def NVVMMemorySpaceAttr :
let assemblyFormat = "`<` $value `>`";
}
+// Attrs describing the scope of the Memory Operation
+def MemScopeKindCTA : I32EnumAttrCase<"CTA", 0, "cta">;
+def MemScopeKindCluster : I32EnumAttrCase<"CLUSTER", 1, "cluster">;
+def MemScopeKindGPU : I32EnumAttrCase<"GPU", 2, "gpu">;
+def MemScopeKindSYS : I32EnumAttrCase<"SYS", 3, "sys">;
+
+def MemScopeKind : I32EnumAttr<"MemScopeKind", "NVVM Memory Scope kind",
+ [MemScopeKindCTA, MemScopeKindCluster, MemScopeKindGPU, MemScopeKindSYS]> {
+ let genSpecializedAttr = 0;
+ let cppNamespace = "::mlir::NVVM";
+}
+def MemScopeKindAttr : EnumAttr<NVVM_Dialect, MemScopeKind, "mem_scope"> {
+ let assemblyFormat = "`<` $value `>`";
+}
+
+// Attrs to disambiguate the cta or cluster space within shared memory
+def SharedSpaceCTA : I32EnumAttrCase<"shared_cta", 0, "cta">;
+def SharedSpaceCluster : I32EnumAttrCase<"shared_cluster", 1, "cluster">;
+def SharedSpace : I32EnumAttr<"SharedSpace", "Shared memory space",
+ [SharedSpaceCTA, SharedSpaceCluster]> {
+ let genSpecializedAttr = 0;
+ let cppNamespace = "::mlir::NVVM";
+}
+def SharedSpaceAttr : EnumAttr<NVVM_Dialect, SharedSpace, "shared_space"> {
+ let assemblyFormat = "`<` $value `>`";
+}
+
//===----------------------------------------------------------------------===//
// NVVM intrinsic operations
//===----------------------------------------------------------------------===//
@@ -1107,17 +1134,6 @@ def NVVM_FenceScClusterOp : NVVM_Op<"fence.sc.cluster"> {
let assemblyFormat = "attr-dict";
}
-def SharedSpaceCTA : I32EnumAttrCase<"shared_cta", 0, "cta">;
-def SharedSpaceCluster : I32EnumAttrCase<"shared_cluster", 1, "cluster">;
-def SharedSpace : I32EnumAttr<"SharedSpace", "Shared memory space",
- [SharedSpaceCTA, SharedSpaceCluster]> {
- let genSpecializedAttr = 0;
- let cppNamespace = "::mlir::NVVM";
-}
-def SharedSpaceAttr : EnumAttr<NVVM_Dialect, SharedSpace, "shared_space"> {
- let assemblyFormat = "`<` $value `>`";
-}
-
def ProxyAlias : I32EnumAttrCase<"alias", 0, "alias">;
def ProxyAsync : I32EnumAttrCase<"async", 1, "async">;
def ProxyAsyncGlobal : I32EnumAttrCase<"async_global", 2, "async.global">;
@@ -1158,21 +1174,6 @@ def NVVM_FenceProxyOp : NVVM_PTXBuilder_Op<"fence.proxy">,
let hasVerifier = 1;
}
-// Attrs describing the scope of the Memory Operation
-def MemScopeKindCTA : I32EnumAttrCase<"CTA", 0, "cta">;
-def MemScopeKindCluster : I32EnumAttrCase<"CLUSTER", 1, "cluster">;
-def MemScopeKindGPU : I32EnumAttrCase<"GPU", 2, "gpu">;
-def MemScopeKindSYS : I32EnumAttrCase<"SYS", 3, "sys">;
-
-def MemScopeKind : I32EnumAttr<"MemScopeKind", "NVVM Memory Scope kind",
- [MemScopeKindCTA, MemScopeKindCluster, MemScopeKindGPU, MemScopeKindSYS]> {
- let genSpecializedAttr = 0;
- let cppNamespace = "::mlir::NVVM";
-}
-def MemScopeKindAttr : EnumAttr<NVVM_Dialect, MemScopeKind, "mem_scope"> {
- let assemblyFormat = "`<` $value `>`";
-}
-
def NVVM_FenceProxyAcquireOp : NVVM_Op<"fence.proxy.acquire">,
Arguments<(ins MemScopeKindAttr:$scope, LLVM_PointerGeneric:$addr, I32:$size,
DefaultValuedAttr<ProxyKindAttr,
|
|
@llvm/pr-subscribers-mlir Author: Durgadoss R (durga4github) ChangesThe mbarrier Ops also require access to the mem_scope and shared_space attributes. Full diff: https://github.com/llvm/llvm-project/pull/168348.diff 1 Files Affected:
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index d4ef5104d3c1f..456d816205b58 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -228,6 +228,33 @@ def NVVMMemorySpaceAttr :
let assemblyFormat = "`<` $value `>`";
}
+// Attrs describing the scope of the Memory Operation
+def MemScopeKindCTA : I32EnumAttrCase<"CTA", 0, "cta">;
+def MemScopeKindCluster : I32EnumAttrCase<"CLUSTER", 1, "cluster">;
+def MemScopeKindGPU : I32EnumAttrCase<"GPU", 2, "gpu">;
+def MemScopeKindSYS : I32EnumAttrCase<"SYS", 3, "sys">;
+
+def MemScopeKind : I32EnumAttr<"MemScopeKind", "NVVM Memory Scope kind",
+ [MemScopeKindCTA, MemScopeKindCluster, MemScopeKindGPU, MemScopeKindSYS]> {
+ let genSpecializedAttr = 0;
+ let cppNamespace = "::mlir::NVVM";
+}
+def MemScopeKindAttr : EnumAttr<NVVM_Dialect, MemScopeKind, "mem_scope"> {
+ let assemblyFormat = "`<` $value `>`";
+}
+
+// Attrs to disambiguate the cta or cluster space within shared memory
+def SharedSpaceCTA : I32EnumAttrCase<"shared_cta", 0, "cta">;
+def SharedSpaceCluster : I32EnumAttrCase<"shared_cluster", 1, "cluster">;
+def SharedSpace : I32EnumAttr<"SharedSpace", "Shared memory space",
+ [SharedSpaceCTA, SharedSpaceCluster]> {
+ let genSpecializedAttr = 0;
+ let cppNamespace = "::mlir::NVVM";
+}
+def SharedSpaceAttr : EnumAttr<NVVM_Dialect, SharedSpace, "shared_space"> {
+ let assemblyFormat = "`<` $value `>`";
+}
+
//===----------------------------------------------------------------------===//
// NVVM intrinsic operations
//===----------------------------------------------------------------------===//
@@ -1107,17 +1134,6 @@ def NVVM_FenceScClusterOp : NVVM_Op<"fence.sc.cluster"> {
let assemblyFormat = "attr-dict";
}
-def SharedSpaceCTA : I32EnumAttrCase<"shared_cta", 0, "cta">;
-def SharedSpaceCluster : I32EnumAttrCase<"shared_cluster", 1, "cluster">;
-def SharedSpace : I32EnumAttr<"SharedSpace", "Shared memory space",
- [SharedSpaceCTA, SharedSpaceCluster]> {
- let genSpecializedAttr = 0;
- let cppNamespace = "::mlir::NVVM";
-}
-def SharedSpaceAttr : EnumAttr<NVVM_Dialect, SharedSpace, "shared_space"> {
- let assemblyFormat = "`<` $value `>`";
-}
-
def ProxyAlias : I32EnumAttrCase<"alias", 0, "alias">;
def ProxyAsync : I32EnumAttrCase<"async", 1, "async">;
def ProxyAsyncGlobal : I32EnumAttrCase<"async_global", 2, "async.global">;
@@ -1158,21 +1174,6 @@ def NVVM_FenceProxyOp : NVVM_PTXBuilder_Op<"fence.proxy">,
let hasVerifier = 1;
}
-// Attrs describing the scope of the Memory Operation
-def MemScopeKindCTA : I32EnumAttrCase<"CTA", 0, "cta">;
-def MemScopeKindCluster : I32EnumAttrCase<"CLUSTER", 1, "cluster">;
-def MemScopeKindGPU : I32EnumAttrCase<"GPU", 2, "gpu">;
-def MemScopeKindSYS : I32EnumAttrCase<"SYS", 3, "sys">;
-
-def MemScopeKind : I32EnumAttr<"MemScopeKind", "NVVM Memory Scope kind",
- [MemScopeKindCTA, MemScopeKindCluster, MemScopeKindGPU, MemScopeKindSYS]> {
- let genSpecializedAttr = 0;
- let cppNamespace = "::mlir::NVVM";
-}
-def MemScopeKindAttr : EnumAttr<NVVM_Dialect, MemScopeKind, "mem_scope"> {
- let assemblyFormat = "`<` $value `>`";
-}
-
def NVVM_FenceProxyAcquireOp : NVVM_Op<"fence.proxy.acquire">,
Arguments<(ins MemScopeKindAttr:$scope, LLVM_PointerGeneric:$addr, I32:$size,
DefaultValuedAttr<ProxyKindAttr,
|
The mbarrier Ops also require access to the
mem_scopeandshared_spaceattributes.Hence, this patch moves their definitions to the beginning of the file alongside
the other attribute definitions.