diff --git a/nixos/modules/profiles/installation-device.nix b/nixos/modules/profiles/installation-device.nix
index 980720691a4..32884f4b875 100644
--- a/nixos/modules/profiles/installation-device.nix
+++ b/nixos/modules/profiles/installation-device.nix
@@ -52,9 +52,9 @@ with lib;
     services.getty.helpLine = ''
       The "nixos" and "root" accounts have empty passwords.
 
-      An ssh daemon is running. You then must set a password
-      for either "root" or "nixos" with `passwd` or add an ssh key
-      to /home/nixos/.ssh/authorized_keys be able to login.
+      To log in over ssh you must set a password for either "nixos" or "root"
+      with `passwd` (prefix with `sudo` for "root"), or add your public key to
+      /home/nixos/.ssh/authorized_keys or /root/.ssh/authorized_keys.
 
       If you need a wireless connection, type
       `sudo systemctl start wpa_supplicant` and configure a
@@ -65,8 +65,8 @@ with lib;
       start the graphical user interface.
     '';
 
-    # We run sshd by default. Login via root is only possible after adding a
-    # password via "passwd" or by adding a ssh key to /home/nixos/.ssh/authorized_keys.
+    # We run sshd by default. Login is only possible after adding a
+    # password via "passwd" or by adding a ssh key to ~/.ssh/authorized_keys.
     # The latter one is particular useful if keys are manually added to
     # installation device for head-less systems i.e. arm boards by manually
     # mounting the storage in a different system.
diff --git a/nixos/modules/system/boot/resolved.nix b/nixos/modules/system/boot/resolved.nix
index 0ab2a875975..4e7201833db 100644
--- a/nixos/modules/system/boot/resolved.nix
+++ b/nixos/modules/system/boot/resolved.nix
@@ -16,7 +16,9 @@ in
       default = false;
       type = types.bool;
       description = lib.mdDoc ''
-        Whether to enable the systemd DNS resolver daemon.
+        Whether to enable the systemd DNS resolver daemon, `systemd-resolved`.
+
+        Search for `services.resolved` to see all options.
       '';
     };
 
diff --git a/pkgs/applications/audio/tauon/default.nix b/pkgs/applications/audio/tauon/default.nix
index 6441182e72a..427c1e9e5dc 100644
--- a/pkgs/applications/audio/tauon/default.nix
+++ b/pkgs/applications/audio/tauon/default.nix
@@ -25,13 +25,13 @@
 
 stdenv.mkDerivation rec {
   pname = "tauon";
-  version = "7.6.2";
+  version = "7.6.3";
 
   src = fetchFromGitHub {
     owner = "Taiko2k";
     repo = "TauonMusicBox";
     rev = "v${version}";
-    hash = "sha256-x/tHCDplC45XEaBaf0aQ0w8AS1SorXtYilJoiOcBDtM=";
+    hash = "sha256-cNR4Ffn9HvgL5KV4FUSnbzEh6VfoKaIbfpb18/qKEns=";
   };
 
   postUnpack = ''
diff --git a/pkgs/applications/editors/vim/plugins/overrides.nix b/pkgs/applications/editors/vim/plugins/overrides.nix
index e15d497a80a..c7a4b2fc96e 100644
--- a/pkgs/applications/editors/vim/plugins/overrides.nix
+++ b/pkgs/applications/editors/vim/plugins/overrides.nix
@@ -706,7 +706,7 @@ self: super: {
   });
 
   noice-nvim = super.noice-nvim.overrideAttrs(old: {
-    dependencies = with self; [ nui-nvim nvim-notify ];
+    dependencies = with self; [ nui-nvim ];
   });
 
   null-ls-nvim = super.null-ls-nvim.overrideAttrs (old: {
diff --git a/pkgs/applications/misc/binocle/default.nix b/pkgs/applications/misc/binocle/default.nix
index 6a999790635..8ec1490bdcc 100644
--- a/pkgs/applications/misc/binocle/default.nix
+++ b/pkgs/applications/misc/binocle/default.nix
@@ -16,16 +16,16 @@
 
 rustPlatform.buildRustPackage rec {
   pname = "binocle";
-  version = "0.3.0";
+  version = "0.3.1";
 
   src = fetchFromGitHub {
     owner = "sharkdp";
     repo = pname;
     rev = "v${version}";
-    sha256 = "0b0hf2aq34kxxj0la0yar5sp44k6mqcbyailp6j6q0mksf1l74bc";
+    sha256 = "sha256-L4l8Gl7Ok/TWqHjNujPx8Qk3UWebs0SgOQNyBNtpnZo=";
   };
 
-  cargoSha256 = "sha256-CZWAHWZYaL54Rl6Jrp8B6w6HK+2fIKQle2x4mGHv2/o=";
+  cargoHash = "sha256-9d0MNQ7jEJKpGbjVtl1XBoOBEVNKDgFouSMrcZ7tXNU=";
 
   nativeBuildInputs = [
     makeWrapper
diff --git a/pkgs/applications/misc/lyx/default.nix b/pkgs/applications/misc/lyx/default.nix
index 921ae570bca..10dbb2c6a74 100644
--- a/pkgs/applications/misc/lyx/default.nix
+++ b/pkgs/applications/misc/lyx/default.nix
@@ -18,9 +18,9 @@ mkDerivation rec {
   '';
 
   # LaTeX is used from $PATH, as people often want to have it with extra pkgs
-  nativeBuildInputs = [ pkg-config makeWrapper ];
+  nativeBuildInputs = [ pkg-config makeWrapper python3 ];
   buildInputs = [
-    qtbase qtsvg python3 file/*for libmagic*/ bc
+    qtbase qtsvg file/*for libmagic*/ bc
     hunspell # enchant
   ];
 
diff --git a/pkgs/applications/networking/browsers/chromium/upstream-info.json b/pkgs/applications/networking/browsers/chromium/upstream-info.json
index 19b777b5e36..e34595db624 100644
--- a/pkgs/applications/networking/browsers/chromium/upstream-info.json
+++ b/pkgs/applications/networking/browsers/chromium/upstream-info.json
@@ -19,22 +19,22 @@
     }
   },
   "beta": {
-    "version": "112.0.5615.49",
-    "sha256": "0hgzbbmz40235binfn3vkkxzvwxilaxg04dclqrz980z7hvkgzfx",
-    "sha256bin64": "146gd9csj08d1ygwwh6gyqqbi7d34mhv3vv7wv4a8z9hn7jhdifg",
+    "version": "113.0.5672.24",
+    "sha256": "1z7yv5lqi1n4ycymkf0kz1v8ig06n430a37ik8hri3a6db8r9lv8",
+    "sha256bin64": "1y9jaw47dgphqr2l8yc36s7k9lf4qrbmfll1d2d1zdjd5ma3slab",
     "deps": {
       "gn": {
-        "version": "2023-02-17",
+        "version": "2023-03-18",
         "url": "https://gn.googlesource.com/gn",
-        "rev": "b25a2f8c2d33f02082f0f258350f5e22c0973108",
-        "sha256": "075p4jwk1apvwmqmvhwfw5f669ci7nxwjq9mz5aa2g5lz4fkdm4c"
+        "rev": "41fef642de70ecdcaaa26be96d56a0398f95abd4",
+        "sha256": "12w4g2dl58283allclpi1c4i6ih9v2xvdb9hpbmfda12v8lizmlq"
       }
     }
   },
   "dev": {
-    "version": "113.0.5672.24",
-    "sha256": "1z7yv5lqi1n4ycymkf0kz1v8ig06n430a37ik8hri3a6db8r9lv8",
-    "sha256bin64": "0byksvk781gmh5fmjmc77jh19gvkzadf78yr9b4c42las44g4pn4",
+    "version": "114.0.5696.0",
+    "sha256": "0hcf971azy3jjsl211qk53b6nk0f4pzf2dwfv3rjh4f6fa3nbwai",
+    "sha256bin64": "1ssj633vdg3ghd87az28q262ly1fk7rc3k70l1531xvj2030ijrl",
     "deps": {
       "gn": {
         "version": "2023-03-18",
diff --git a/pkgs/applications/networking/remote/rustdesk/Cargo.lock b/pkgs/applications/networking/remote/rustdesk/Cargo.lock
new file mode 100644
index 00000000000..656b0fd9b06
--- /dev/null
+++ b/pkgs/applications/networking/remote/rustdesk/Cargo.lock
@@ -0,0 +1,5193 @@
+# This file is automatically @generated by Cargo.
+# It is not intended for manual editing.
+version = 3
+
+[[package]]
+name = "addr2line"
+version = "0.19.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a76fd60b23679b7d19bd066031410fb7e458ccc5e958eb5c325888ce4baedc97"
+dependencies = [
+ "gimli",
+]
+
+[[package]]
+name = "adler"
+version = "1.0.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f26201604c87b1e01bd3d98f8d5d9a8fcbb815e8cedb41ffccbeb4bf593a35fe"
+
+[[package]]
+name = "adler32"
+version = "1.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "aae1277d39aeec15cb388266ecc24b11c80469deae6067e17a1a7aa9e5c1f234"
+
+[[package]]
+name = "aho-corasick"
+version = "0.7.20"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "cc936419f96fa211c1b9166887b38e5e40b19958e5b895be7c1f93adec7071ac"
+dependencies = [
+ "memchr",
+]
+
+[[package]]
+name = "alsa"
+version = "0.6.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5915f52fe2cf65e83924d037b6c5290b7cee097c6b5c8700746e6168a343fd6b"
+dependencies = [
+ "alsa-sys",
+ "bitflags",
+ "libc",
+ "nix 0.23.2",
+]
+
+[[package]]
+name = "alsa-sys"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "db8fee663d06c4e303404ef5f40488a53e062f89ba8bfed81f42325aafad1527"
+dependencies = [
+ "libc",
+ "pkg-config",
+]
+
+[[package]]
+name = "android_log-sys"
+version = "0.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "85965b6739a430150bdd138e2374a98af0c3ee0d030b3bb7fc3bddff58d0102e"
+
+[[package]]
+name = "android_logger"
+version = "0.11.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8619b80c242aa7bd638b5c7ddd952addeecb71f69c75e33f1d47b2804f8f883a"
+dependencies = [
+ "android_log-sys",
+ "env_logger 0.10.0",
+ "log",
+ "once_cell",
+]
+
+[[package]]
+name = "android_system_properties"
+version = "0.1.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "819e7219dbd41043ac279b19830f2efc897156490d7fd6ea916720117ee66311"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "ansi_term"
+version = "0.12.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d52a9bb7ec0cf484c551830a7ce27bd20d67eac647e1befb56b0be4ee39a55d2"
+dependencies = [
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "anyhow"
+version = "1.0.70"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7de8ce5e0f9f8d88245311066a578d72b7af3e7088f32783804676302df237e4"
+
+[[package]]
+name = "arboard"
+version = "2.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "dc120354d1b5ec6d7aaf4876b602def75595937b5e15d356eb554ab5177e08bb"
+dependencies = [
+ "clipboard-win",
+ "core-graphics",
+ "image",
+ "log",
+ "objc",
+ "objc-foundation",
+ "objc_id",
+ "parking_lot 0.12.1",
+ "thiserror",
+ "winapi 0.3.9",
+ "x11rb",
+]
+
+[[package]]
+name = "async-trait"
+version = "0.1.68"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b9ccdd8f2a161be9bd5c023df56f1b2a0bd1d83872ae53b71a84a12c9bf6e842"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 2.0.13",
+]
+
+[[package]]
+name = "atk"
+version = "0.9.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "812b4911e210bd51b24596244523c856ca749e6223c50a7fbbba3f89ee37c426"
+dependencies = [
+ "atk-sys 0.10.0",
+ "bitflags",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+]
+
+[[package]]
+name = "atk"
+version = "0.15.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2c3d816ce6f0e2909a96830d6911c2aff044370b1ef92d7f267b43bae5addedd"
+dependencies = [
+ "atk-sys 0.15.1",
+ "bitflags",
+ "glib 0.15.12",
+ "libc",
+]
+
+[[package]]
+name = "atk-sys"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f530e4af131d94cc4fa15c5c9d0348f0ef28bac64ba660b6b2a1cf2605dedfce"
+dependencies = [
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "atk-sys"
+version = "0.15.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "58aeb089fb698e06db8089971c7ee317ab9644bade33383f63631437b03aafb6"
+dependencies = [
+ "glib-sys 0.15.10",
+ "gobject-sys 0.15.10",
+ "libc",
+ "system-deps 6.0.4",
+]
+
+[[package]]
+name = "atty"
+version = "0.2.14"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d9b39be18770d11421cdb1b9947a45dd3f37e93092cbf377614828a319d5fee8"
+dependencies = [
+ "hermit-abi 0.1.19",
+ "libc",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "autocfg"
+version = "0.1.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0dde43e75fd43e8a1bf86103336bc699aa8d17ad1be60c76c0bdfd4828e19b78"
+dependencies = [
+ "autocfg 1.1.0",
+]
+
+[[package]]
+name = "autocfg"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d468802bab17cbc0cc575e9b053f41e72aa36bfa6b7f55e3529ffa43161b97fa"
+
+[[package]]
+name = "backtrace"
+version = "0.3.67"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "233d376d6d185f2a3093e58f283f60f880315b6c60075b01f36b3b85154564ca"
+dependencies = [
+ "addr2line",
+ "cc",
+ "cfg-if 1.0.0",
+ "libc",
+ "miniz_oxide 0.6.2",
+ "object",
+ "rustc-demangle",
+]
+
+[[package]]
+name = "base64"
+version = "0.13.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9e1b586273c5702936fe7b7d6896644d8be71e6314cfe09d3167c95f712589e8"
+
+[[package]]
+name = "base64"
+version = "0.21.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a4a4ddaa51a5bc52a6948f74c06d20aaaddb71924eab79b8c97a8c556e942d6a"
+
+[[package]]
+name = "bindgen"
+version = "0.59.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2bd2a9a458e8f4304c52c43ebb0cfbd520289f8379a52e329a38afda99bf8eb8"
+dependencies = [
+ "bitflags",
+ "cexpr",
+ "clang-sys",
+ "clap 2.34.0",
+ "env_logger 0.9.3",
+ "lazy_static",
+ "lazycell",
+ "log",
+ "peeking_take_while",
+ "proc-macro2",
+ "quote",
+ "regex",
+ "rustc-hash",
+ "shlex",
+ "which 4.4.0",
+]
+
+[[package]]
+name = "bindgen"
+version = "0.64.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c4243e6031260db77ede97ad86c27e501d646a27ab57b59a574f725d98ab1fb4"
+dependencies = [
+ "bitflags",
+ "cexpr",
+ "clang-sys",
+ "lazy_static",
+ "lazycell",
+ "peeking_take_while",
+ "proc-macro2",
+ "quote",
+ "regex",
+ "rustc-hash",
+ "shlex",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "bitflags"
+version = "1.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bef38d45163c2f1dde094a7dfd33ccf595c92905c8f8f4fdc18d06fb1037718a"
+
+[[package]]
+name = "block"
+version = "0.1.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a"
+
+[[package]]
+name = "block-buffer"
+version = "0.10.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3078c7629b62d3f0439517fa394996acacc5cbc91c5a20d8c658e77abd503a71"
+dependencies = [
+ "generic-array",
+]
+
+[[package]]
+name = "bumpalo"
+version = "3.12.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0d261e256854913907f67ed06efbc3338dfe6179796deefc1ff763fc1aee5535"
+
+[[package]]
+name = "bytemuck"
+version = "1.13.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "17febce684fd15d89027105661fec94afb475cb995fbc59d2865198446ba2eea"
+
+[[package]]
+name = "byteorder"
+version = "1.4.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "14c189c53d098945499cdfa7ecc63567cf3886b3332b312a5b4585d8d3a6a610"
+
+[[package]]
+name = "bytes"
+version = "1.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "89b2fd2a0dcf38d7971e2194b6b6eebab45ae01067456a7fd93d5547a61b70be"
+
+[[package]]
+name = "cairo-rs"
+version = "0.9.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c5c0f2e047e8ca53d0ff249c54ae047931d7a6ebe05d00af73e0ffeb6e34bdb8"
+dependencies = [
+ "bitflags",
+ "cairo-sys-rs 0.10.0",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "thiserror",
+]
+
+[[package]]
+name = "cairo-rs"
+version = "0.15.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c76ee391b03d35510d9fa917357c7f1855bd9a6659c95a1b392e33f49b3369bc"
+dependencies = [
+ "bitflags",
+ "cairo-sys-rs 0.15.1",
+ "glib 0.15.12",
+ "libc",
+ "thiserror",
+]
+
+[[package]]
+name = "cairo-sys-rs"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2ed2639b9ad5f1d6efa76de95558e11339e7318426d84ac4890b86c03e828ca7"
+dependencies = [
+ "glib-sys 0.10.1",
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "cairo-sys-rs"
+version = "0.15.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3c55d429bef56ac9172d25fecb85dc8068307d17acd74b377866b7a1ef25d3c8"
+dependencies = [
+ "glib-sys 0.15.10",
+ "libc",
+ "system-deps 6.0.4",
+]
+
+[[package]]
+name = "cc"
+version = "1.0.79"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "50d30906286121d95be3d479533b458f87493b30a4b5f79a607db8f5d11aa91f"
+dependencies = [
+ "jobserver",
+]
+
+[[package]]
+name = "cesu8"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6d43a04d8753f35258c91f8ec639f792891f748a1edbd759cf1dcea3382ad83c"
+
+[[package]]
+name = "cexpr"
+version = "0.6.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6fac387a98bb7c37292057cffc56d62ecb629900026402633ae9160df93a8766"
+dependencies = [
+ "nom",
+]
+
+[[package]]
+name = "cfg-expr"
+version = "0.14.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a35b255461940a32985c627ce82900867c61db1659764d3675ea81963f72a4c6"
+dependencies = [
+ "smallvec",
+]
+
+[[package]]
+name = "cfg-if"
+version = "0.1.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4785bdd1c96b2a846b2bd7cc02e86b6b3dbf14e7e53446c4f54c92a361040822"
+
+[[package]]
+name = "cfg-if"
+version = "1.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "baf1de4339761588bc0619e3cbc0120ee582ebb74b53b4efbf79117bd2da40fd"
+
+[[package]]
+name = "chrono"
+version = "0.4.24"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4e3c5919066adf22df73762e50cffcde3a758f2a848b113b586d1f86728b673b"
+dependencies = [
+ "iana-time-zone",
+ "num-integer",
+ "num-traits 0.2.15",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "clang-sys"
+version = "1.6.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c688fc74432808e3eb684cae8830a86be1d66a2bd58e1f248ed0960a590baf6f"
+dependencies = [
+ "glob",
+ "libc",
+ "libloading",
+]
+
+[[package]]
+name = "clap"
+version = "2.34.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a0610544180c38b88101fecf2dd634b174a62eef6946f84dfc6a7127512b381c"
+dependencies = [
+ "ansi_term",
+ "atty",
+ "bitflags",
+ "strsim 0.8.0",
+ "textwrap 0.11.0",
+ "unicode-width",
+ "vec_map",
+]
+
+[[package]]
+name = "clap"
+version = "3.2.23"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "71655c45cb9845d3270c9d6df84ebe72b4dad3c2ba3f7023ad47c144e4e473a5"
+dependencies = [
+ "atty",
+ "bitflags",
+ "clap_lex",
+ "indexmap",
+ "strsim 0.10.0",
+ "termcolor",
+ "textwrap 0.16.0",
+]
+
+[[package]]
+name = "clap_lex"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2850f2f5a82cbf437dd5af4d49848fbdfc27c157c3d010345776f952765261c5"
+dependencies = [
+ "os_str_bytes",
+]
+
+[[package]]
+name = "clipboard"
+version = "0.1.0"
+dependencies = [
+ "cc",
+ "hbb_common",
+ "lazy_static",
+ "serde 1.0.159",
+ "serde_derive",
+ "thiserror",
+]
+
+[[package]]
+name = "clipboard-master"
+version = "3.1.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "459887701008d8ee21f8de7f45f0f0707417c7eea3311973f6e67222bd686b7a"
+dependencies = [
+ "objc",
+ "objc-foundation",
+ "objc_id",
+ "winapi 0.3.9",
+ "windows-win",
+ "x11-clipboard",
+]
+
+[[package]]
+name = "clipboard-win"
+version = "4.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7191c27c2357d9b7ef96baac1773290d4ca63b24205b82a3fd8a0637afcf0362"
+dependencies = [
+ "error-code",
+ "str-buf",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "cloudabi"
+version = "0.0.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ddfc5b9aa5d4507acaf872de71051dfd0e309860e88966e1051e462a077aac4f"
+dependencies = [
+ "bitflags",
+]
+
+[[package]]
+name = "cmake"
+version = "0.1.50"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a31c789563b815f77f4250caee12365734369f942439b7defd71e18a48197130"
+dependencies = [
+ "cc",
+]
+
+[[package]]
+name = "cocoa"
+version = "0.24.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f425db7937052c684daec3bd6375c8abe2d146dca4b8b143d6db777c39138f3a"
+dependencies = [
+ "bitflags",
+ "block",
+ "cocoa-foundation",
+ "core-foundation",
+ "core-graphics",
+ "foreign-types",
+ "libc",
+ "objc",
+]
+
+[[package]]
+name = "cocoa-foundation"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "931d3837c286f56e3c58423ce4eba12d08db2374461a785c86f672b08b5650d6"
+dependencies = [
+ "bitflags",
+ "block",
+ "core-foundation",
+ "core-graphics-types",
+ "foreign-types",
+ "libc",
+ "objc",
+]
+
+[[package]]
+name = "codespan-reporting"
+version = "0.11.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3538270d33cc669650c4b093848450d380def10c331d38c768e34cac80576e6e"
+dependencies = [
+ "termcolor",
+ "unicode-width",
+]
+
+[[package]]
+name = "color_quant"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3d7b894f5411737b7867f4827955924d7c254fc9f4d91a6aad6b097804b1018b"
+
+[[package]]
+name = "combine"
+version = "4.6.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "35ed6e9d84f0b51a7f52daf1c7d71dd136fd7a3f41a8462b8cdb8c78d920fad4"
+dependencies = [
+ "bytes",
+ "memchr",
+]
+
+[[package]]
+name = "confy"
+version = "0.4.0"
+source = "git+https://github.com/open-trade/confy#630cc28a396cb7d01eefdd9f3824486fe4d8554b"
+dependencies = [
+ "directories-next",
+ "serde 1.0.159",
+ "thiserror",
+ "toml 0.5.11",
+]
+
+[[package]]
+name = "core-foundation"
+version = "0.9.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "194a7a9e6de53fa55116934067c844d9d749312f75c6f6d0980e8c252f8c2146"
+dependencies = [
+ "core-foundation-sys",
+ "libc",
+]
+
+[[package]]
+name = "core-foundation-sys"
+version = "0.8.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e496a50fda8aacccc86d7529e2c1e0892dbd0f898a6b5645b5561b89c3210efa"
+
+[[package]]
+name = "core-graphics"
+version = "0.22.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2581bbab3b8ffc6fcbd550bf46c355135d16e9ff2a6ea032ad6b9bf1d7efe4fb"
+dependencies = [
+ "bitflags",
+ "core-foundation",
+ "core-graphics-types",
+ "foreign-types",
+ "libc",
+]
+
+[[package]]
+name = "core-graphics-types"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3a68b68b3446082644c91ac778bf50cd4104bfb002b5a6a7c44cca5a2c70788b"
+dependencies = [
+ "bitflags",
+ "core-foundation",
+ "foreign-types",
+ "libc",
+]
+
+[[package]]
+name = "coreaudio-rs"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "11894b20ebfe1ff903cbdc52259693389eea03b94918a2def2c30c3bf227ad88"
+dependencies = [
+ "bitflags",
+ "coreaudio-sys",
+]
+
+[[package]]
+name = "coreaudio-sys"
+version = "0.2.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f034b2258e6c4ade2f73bf87b21047567fb913ee9550837c2316d139b0262b24"
+dependencies = [
+ "bindgen 0.64.0",
+]
+
+[[package]]
+name = "cpal"
+version = "0.13.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "74117836a5124f3629e4b474eed03e479abaf98988b4bb317e29f08cfe0e4116"
+dependencies = [
+ "alsa",
+ "core-foundation-sys",
+ "coreaudio-rs",
+ "jni",
+ "js-sys",
+ "lazy_static",
+ "libc",
+ "mach",
+ "ndk",
+ "ndk-glue",
+ "nix 0.23.2",
+ "oboe",
+ "parking_lot 0.11.2",
+ "stdweb",
+ "thiserror",
+ "web-sys",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "cpufeatures"
+version = "0.2.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "280a9f2d8b3a38871a3c8a46fb80db65e5e5ed97da80c4d08bf27fb63e35e181"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "crc32fast"
+version = "1.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b540bd8bc810d3885c6ea91e2018302f68baba2129ab3e88f32389ee9370880d"
+dependencies = [
+ "cfg-if 1.0.0",
+]
+
+[[package]]
+name = "crossbeam-channel"
+version = "0.5.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "cf2b3e8478797446514c91ef04bafcb59faba183e621ad488df88983cc14128c"
+dependencies = [
+ "cfg-if 1.0.0",
+ "crossbeam-utils",
+]
+
+[[package]]
+name = "crossbeam-deque"
+version = "0.8.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ce6fd6f855243022dcecf8702fef0c297d4338e226845fe067f6341ad9fa0cef"
+dependencies = [
+ "cfg-if 1.0.0",
+ "crossbeam-epoch",
+ "crossbeam-utils",
+]
+
+[[package]]
+name = "crossbeam-epoch"
+version = "0.9.14"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "46bd5f3f85273295a9d14aedfb86f6aadbff6d8f5295c4a9edb08e819dcf5695"
+dependencies = [
+ "autocfg 1.1.0",
+ "cfg-if 1.0.0",
+ "crossbeam-utils",
+ "memoffset 0.8.0",
+ "scopeguard",
+]
+
+[[package]]
+name = "crossbeam-queue"
+version = "0.3.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d1cfb3ea8a53f37c40dea2c7bedcbd88bdfae54f5e2175d6ecaff1c988353add"
+dependencies = [
+ "cfg-if 1.0.0",
+ "crossbeam-utils",
+]
+
+[[package]]
+name = "crossbeam-utils"
+version = "0.8.15"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3c063cd8cc95f5c377ed0d4b49a4b21f632396ff690e8470c29b3359b346984b"
+dependencies = [
+ "cfg-if 1.0.0",
+]
+
+[[package]]
+name = "crypto-common"
+version = "0.1.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1bfb12502f3fc46cca1bb51ac28df9d618d813cdc3d2f25b9fe775a34af26bb3"
+dependencies = [
+ "generic-array",
+ "typenum",
+]
+
+[[package]]
+name = "ctrlc"
+version = "3.2.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bbcf33c2a618cbe41ee43ae6e9f2e48368cd9f9db2896f10167d8d762679f639"
+dependencies = [
+ "nix 0.26.2",
+ "windows-sys 0.45.0",
+]
+
+[[package]]
+name = "cxx"
+version = "1.0.94"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f61f1b6389c3fe1c316bf8a4dccc90a38208354b330925bce1f74a6c4756eb93"
+dependencies = [
+ "cc",
+ "cxxbridge-flags",
+ "cxxbridge-macro",
+ "link-cplusplus",
+]
+
+[[package]]
+name = "cxx-build"
+version = "1.0.94"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "12cee708e8962df2aeb38f594aae5d827c022b6460ac71a7a3e2c3c2aae5a07b"
+dependencies = [
+ "cc",
+ "codespan-reporting",
+ "once_cell",
+ "proc-macro2",
+ "quote",
+ "scratch",
+ "syn 2.0.13",
+]
+
+[[package]]
+name = "cxxbridge-flags"
+version = "1.0.94"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7944172ae7e4068c533afbb984114a56c46e9ccddda550499caa222902c7f7bb"
+
+[[package]]
+name = "cxxbridge-macro"
+version = "1.0.94"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2345488264226bf682893e25de0769f3360aac9957980ec49361b083ddaa5bc5"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 2.0.13",
+]
+
+[[package]]
+name = "darling"
+version = "0.13.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a01d95850c592940db9b8194bc39f4bc0e89dee5c4265e4b1807c34a9aba453c"
+dependencies = [
+ "darling_core",
+ "darling_macro",
+]
+
+[[package]]
+name = "darling_core"
+version = "0.13.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "859d65a907b6852c9361e3185c862aae7fafd2887876799fa55f5f99dc40d610"
+dependencies = [
+ "fnv",
+ "ident_case",
+ "proc-macro2",
+ "quote",
+ "strsim 0.10.0",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "darling_macro"
+version = "0.13.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9c972679f83bdf9c42bd905396b6c3588a843a17f0f16dfcfa3e2c5d57441835"
+dependencies = [
+ "darling_core",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "dasp"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7381b67da416b639690ac77c73b86a7b5e64a29e31d1f75fb3b1102301ef355a"
+dependencies = [
+ "dasp_envelope",
+ "dasp_frame",
+ "dasp_interpolate",
+ "dasp_peak",
+ "dasp_ring_buffer",
+ "dasp_rms",
+ "dasp_sample",
+ "dasp_signal",
+ "dasp_slice",
+ "dasp_window",
+]
+
+[[package]]
+name = "dasp_envelope"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8ec617ce7016f101a87fe85ed44180839744265fae73bb4aa43e7ece1b7668b6"
+dependencies = [
+ "dasp_frame",
+ "dasp_peak",
+ "dasp_ring_buffer",
+ "dasp_rms",
+ "dasp_sample",
+]
+
+[[package]]
+name = "dasp_frame"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b2a3937f5fe2135702897535c8d4a5553f8b116f76c1529088797f2eee7c5cd6"
+dependencies = [
+ "dasp_sample",
+]
+
+[[package]]
+name = "dasp_interpolate"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7fc975a6563bb7ca7ec0a6c784ead49983a21c24835b0bc96eea11ee407c7486"
+dependencies = [
+ "dasp_frame",
+ "dasp_ring_buffer",
+ "dasp_sample",
+]
+
+[[package]]
+name = "dasp_peak"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5cf88559d79c21f3d8523d91250c397f9a15b5fc72fbb3f87fdb0a37b79915bf"
+dependencies = [
+ "dasp_frame",
+ "dasp_sample",
+]
+
+[[package]]
+name = "dasp_ring_buffer"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "07d79e19b89618a543c4adec9c5a347fe378a19041699b3278e616e387511ea1"
+
+[[package]]
+name = "dasp_rms"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a6c5dcb30b7e5014486e2822537ea2beae50b19722ffe2ed7549ab03774575aa"
+dependencies = [
+ "dasp_frame",
+ "dasp_ring_buffer",
+ "dasp_sample",
+]
+
+[[package]]
+name = "dasp_sample"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0c87e182de0887fd5361989c677c4e8f5000cd9491d6d563161a8f3a5519fc7f"
+
+[[package]]
+name = "dasp_signal"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "aa1ab7d01689c6ed4eae3d38fe1cea08cba761573fbd2d592528d55b421077e7"
+dependencies = [
+ "dasp_envelope",
+ "dasp_frame",
+ "dasp_interpolate",
+ "dasp_peak",
+ "dasp_ring_buffer",
+ "dasp_rms",
+ "dasp_sample",
+ "dasp_window",
+]
+
+[[package]]
+name = "dasp_slice"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4e1c7335d58e7baedafa516cb361360ff38d6f4d3f9d9d5ee2a2fc8e27178fa1"
+dependencies = [
+ "dasp_frame",
+ "dasp_sample",
+]
+
+[[package]]
+name = "dasp_window"
+version = "0.11.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "99ded7b88821d2ce4e8b842c9f1c86ac911891ab89443cc1de750cae764c5076"
+dependencies = [
+ "dasp_sample",
+]
+
+[[package]]
+name = "dbus"
+version = "0.9.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1bb21987b9fb1613058ba3843121dd18b163b254d8a6e797e144cbac14d96d1b"
+dependencies = [
+ "libc",
+ "libdbus-sys",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "deflate"
+version = "0.8.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "73770f8e1fe7d64df17ca66ad28994a0a623ea497fa69486e14984e715c5d174"
+dependencies = [
+ "adler32",
+ "byteorder",
+]
+
+[[package]]
+name = "digest"
+version = "0.10.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8168378f4e5023e7218c89c891c0fd8ecdb5e5e4f18cb78f38cf245dd021e76f"
+dependencies = [
+ "block-buffer",
+ "crypto-common",
+]
+
+[[package]]
+name = "directories-next"
+version = "2.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "339ee130d97a610ea5a5872d2bbb130fdf68884ff09d3028b81bec8a1ac23bbc"
+dependencies = [
+ "cfg-if 1.0.0",
+ "dirs-sys-next",
+]
+
+[[package]]
+name = "dirs-next"
+version = "2.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b98cf8ebf19c3d1b223e151f99a4f9f0690dca41414773390fc824184ac833e1"
+dependencies = [
+ "cfg-if 1.0.0",
+ "dirs-sys-next",
+]
+
+[[package]]
+name = "dirs-sys-next"
+version = "0.1.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4ebda144c4fe02d1f7ea1a7d9641b6fc6b580adcfa024ae48797ecdeb6825b4d"
+dependencies = [
+ "libc",
+ "redox_users",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "dispatch"
+version = "0.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bd0c93bb4b0c6d9b77f4435b0ae98c24d17f1c45b2ff844c6151a07256ca923b"
+
+[[package]]
+name = "docopt"
+version = "1.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7f3f119846c823f9eafcf953a8f6ffb6ed69bf6240883261a7f13b634579a51f"
+dependencies = [
+ "lazy_static",
+ "regex",
+ "serde 1.0.159",
+ "strsim 0.10.0",
+]
+
+[[package]]
+name = "dtoa"
+version = "0.4.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "56899898ce76aaf4a0f24d914c97ea6ed976d42fec6ad33fcbb0a1103e07b2b0"
+
+[[package]]
+name = "ed25519"
+version = "1.5.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "91cff35c70bba8a626e3185d8cd48cc11b5437e1a5bcd15b9b5fa3c64b6dfee7"
+dependencies = [
+ "signature",
+]
+
+[[package]]
+name = "either"
+version = "1.8.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7fcaabb2fef8c910e7f4c7ce9f67a1283a1715879a7c230ca9d6d1ae31f16d91"
+
+[[package]]
+name = "enigo"
+version = "0.0.14"
+dependencies = [
+ "core-graphics",
+ "libc",
+ "log",
+ "objc",
+ "pkg-config",
+ "serde 1.0.159",
+ "serde_derive",
+ "unicode-segmentation",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "enum-map"
+version = "2.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "988f0d17a0fa38291e5f41f71ea8d46a5d5497b9054d5a759fae2cbb819f2356"
+dependencies = [
+ "enum-map-derive",
+]
+
+[[package]]
+name = "enum-map-derive"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2a4da76b3b6116d758c7ba93f7ec6a35d2e2cf24feda76c6e38a375f4d5c59f2"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "env_logger"
+version = "0.9.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a12e6657c4c97ebab115a42dcee77225f7f482cdd841cf7088c657a42e9e00e7"
+dependencies = [
+ "atty",
+ "humantime",
+ "log",
+ "regex",
+ "termcolor",
+]
+
+[[package]]
+name = "env_logger"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "85cdab6a89accf66733ad5a1693a4dcced6aeff64602b634530dd73c1f3ee9f0"
+dependencies = [
+ "log",
+ "regex",
+]
+
+[[package]]
+name = "err-derive"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c34a887c8df3ed90498c1c437ce21f211c8e27672921a8ffa293cb8d6d4caa9e"
+dependencies = [
+ "proc-macro-error",
+ "proc-macro2",
+ "quote",
+ "rustversion",
+ "syn 1.0.109",
+ "synstructure",
+]
+
+[[package]]
+name = "errno"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4bcfec3a70f97c962c307b2d2c56e358cf1d00b558d74262b5f929ee8cc7e73a"
+dependencies = [
+ "errno-dragonfly",
+ "libc",
+ "windows-sys 0.48.0",
+]
+
+[[package]]
+name = "errno-dragonfly"
+version = "0.1.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "aa68f1b12764fab894d2755d2518754e71b4fd80ecfb822714a1206c2aab39bf"
+dependencies = [
+ "cc",
+ "libc",
+]
+
+[[package]]
+name = "error-code"
+version = "2.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "64f18991e7bf11e7ffee451b5318b5c1a73c52d0d0ada6e5a3017c8c1ced6a21"
+dependencies = [
+ "libc",
+ "str-buf",
+]
+
+[[package]]
+name = "failure"
+version = "0.1.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d32e9bd16cc02eae7db7ef620b392808b89f6a5e16bb3497d159c6b92a0f4f86"
+dependencies = [
+ "backtrace",
+]
+
+[[package]]
+name = "fastrand"
+version = "1.9.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e51093e27b0797c359783294ca4f0a911c270184cb10f85783b118614a1501be"
+dependencies = [
+ "instant",
+]
+
+[[package]]
+name = "field-offset"
+version = "0.3.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a3cf3a800ff6e860c863ca6d4b16fd999db8b752819c1606884047b73e468535"
+dependencies = [
+ "memoffset 0.8.0",
+ "rustc_version",
+]
+
+[[package]]
+name = "filetime"
+version = "0.2.21"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5cbc844cecaee9d4443931972e1289c8ff485cb4cc2767cb03ca139ed6885153"
+dependencies = [
+ "cfg-if 1.0.0",
+ "libc",
+ "redox_syscall 0.2.16",
+ "windows-sys 0.48.0",
+]
+
+[[package]]
+name = "flate2"
+version = "1.0.25"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a8a2db397cb1c8772f31494cb8917e48cd1e64f0fa7efac59fbd741a0a8ce841"
+dependencies = [
+ "crc32fast",
+ "miniz_oxide 0.6.2",
+]
+
+[[package]]
+name = "flexi_logger"
+version = "0.22.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0c76a80dd14a27fc3d8bc696502132cb52b3f227256fd8601166c3a35e45f409"
+dependencies = [
+ "ansi_term",
+ "atty",
+ "chrono",
+ "crossbeam-channel",
+ "crossbeam-queue",
+ "glob",
+ "lazy_static",
+ "log",
+ "regex",
+ "rustversion",
+ "thiserror",
+ "time",
+]
+
+[[package]]
+name = "fnv"
+version = "1.0.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3f9eec918d3f24069decb9af1554cad7c880e2da24a9afd88aca000531ab82c1"
+
+[[package]]
+name = "foreign-types"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f6f339eb8adc052cd2ca78910fda869aefa38d22d5cb648e6485e4d3fc06f3b1"
+dependencies = [
+ "foreign-types-shared",
+]
+
+[[package]]
+name = "foreign-types-shared"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "00b0228411908ca8685dba7fc2cdd70ec9990a6e753e89b6ac91a84c40fbaf4b"
+
+[[package]]
+name = "fuchsia-cprng"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a06f77d526c1a601b7c4cdd98f54b5eaabffc14d5f2f0296febdc7f357c6d3ba"
+
+[[package]]
+name = "fuchsia-zircon"
+version = "0.3.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2e9763c69ebaae630ba35f74888db465e49e259ba1bc0eda7d06f4a067615d82"
+dependencies = [
+ "bitflags",
+ "fuchsia-zircon-sys",
+]
+
+[[package]]
+name = "fuchsia-zircon-sys"
+version = "0.3.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3dcaa9ae7725d12cdb85b3ad99a434db70b468c09ded17e012d86b5c1010f7a7"
+
+[[package]]
+name = "futures"
+version = "0.3.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "23342abe12aba583913b2e62f22225ff9c950774065e4bfb61a19cd9770fec40"
+dependencies = [
+ "futures-channel",
+ "futures-core",
+ "futures-executor",
+ "futures-io",
+ "futures-sink",
+ "futures-task",
+ "futures-util",
+]
+
+[[package]]
+name = "futures-channel"
+version = "0.3.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "955518d47e09b25bbebc7a18df10b81f0c766eaf4c4f1cccef2fca5f2a4fb5f2"
+dependencies = [
+ "futures-core",
+ "futures-sink",
+]
+
+[[package]]
+name = "futures-core"
+version = "0.3.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4bca583b7e26f571124fe5b7561d49cb2868d79116cfa0eefce955557c6fee8c"
+
+[[package]]
+name = "futures-executor"
+version = "0.3.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ccecee823288125bd88b4d7f565c9e58e41858e47ab72e8ea2d64e93624386e0"
+dependencies = [
+ "futures-core",
+ "futures-task",
+ "futures-util",
+]
+
+[[package]]
+name = "futures-io"
+version = "0.3.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4fff74096e71ed47f8e023204cfd0aa1289cd54ae5430a9523be060cdb849964"
+
+[[package]]
+name = "futures-macro"
+version = "0.3.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "89ca545a94061b6365f2c7355b4b32bd20df3ff95f02da9329b34ccc3bd6ee72"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 2.0.13",
+]
+
+[[package]]
+name = "futures-sink"
+version = "0.3.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f43be4fe21a13b9781a69afa4985b0f6ee0e1afab2c6f454a8cf30e2b2237b6e"
+
+[[package]]
+name = "futures-task"
+version = "0.3.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "76d3d132be6c0e6aa1534069c705a74a5997a356c0dc2f86a47765e5617c5b65"
+
+[[package]]
+name = "futures-util"
+version = "0.3.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "26b01e40b772d54cf6c6d721c1d1abd0647a0106a12ecaa1c186273392a69533"
+dependencies = [
+ "futures-channel",
+ "futures-core",
+ "futures-io",
+ "futures-macro",
+ "futures-sink",
+ "futures-task",
+ "memchr",
+ "pin-project-lite",
+ "pin-utils",
+ "slab",
+]
+
+[[package]]
+name = "fxhash"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c31b6d751ae2c7f11320402d34e41349dd1016f8d5d45e48c4312bc8625af50c"
+dependencies = [
+ "byteorder",
+]
+
+[[package]]
+name = "gdk"
+version = "0.13.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "db00839b2a68a7a10af3fa28dfb3febaba3a20c3a9ac2425a33b7df1f84a6b7d"
+dependencies = [
+ "bitflags",
+ "cairo-rs 0.9.1",
+ "cairo-sys-rs 0.10.0",
+ "gdk-pixbuf 0.9.0",
+ "gdk-sys 0.10.0",
+ "gio 0.9.1",
+ "gio-sys 0.10.1",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "pango 0.9.1",
+]
+
+[[package]]
+name = "gdk"
+version = "0.15.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a6e05c1f572ab0e1f15be94217f0dc29088c248b14f792a5ff0af0d84bcda9e8"
+dependencies = [
+ "bitflags",
+ "cairo-rs 0.15.12",
+ "gdk-pixbuf 0.15.11",
+ "gdk-sys 0.15.1",
+ "gio 0.15.12",
+ "glib 0.15.12",
+ "libc",
+ "pango 0.15.10",
+]
+
+[[package]]
+name = "gdk-pixbuf"
+version = "0.9.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8f6dae3cb99dd49b758b88f0132f8d401108e63ae8edd45f432d42cdff99998a"
+dependencies = [
+ "gdk-pixbuf-sys 0.10.0",
+ "gio 0.9.1",
+ "gio-sys 0.10.1",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+]
+
+[[package]]
+name = "gdk-pixbuf"
+version = "0.15.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ad38dd9cc8b099cceecdf41375bb6d481b1b5a7cd5cd603e10a69a9383f8619a"
+dependencies = [
+ "bitflags",
+ "gdk-pixbuf-sys 0.15.10",
+ "gio 0.15.12",
+ "glib 0.15.12",
+ "libc",
+]
+
+[[package]]
+name = "gdk-pixbuf-sys"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3bfe468a7f43e97b8d193a762b6c5cf67a7d36cacbc0b9291dbcae24bfea1e8f"
+dependencies = [
+ "gio-sys 0.10.1",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "gdk-pixbuf-sys"
+version = "0.15.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "140b2f5378256527150350a8346dbdb08fadc13453a7a2d73aecd5fab3c402a7"
+dependencies = [
+ "gio-sys 0.15.10",
+ "glib-sys 0.15.10",
+ "gobject-sys 0.15.10",
+ "libc",
+ "system-deps 6.0.4",
+]
+
+[[package]]
+name = "gdk-sys"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0a9653cfc500fd268015b1ac055ddbc3df7a5c9ea3f4ccef147b3957bd140d69"
+dependencies = [
+ "cairo-sys-rs 0.10.0",
+ "gdk-pixbuf-sys 0.10.0",
+ "gio-sys 0.10.1",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "pango-sys 0.10.0",
+ "pkg-config",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "gdk-sys"
+version = "0.15.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "32e7a08c1e8f06f4177fb7e51a777b8c1689f743a7bc11ea91d44d2226073a88"
+dependencies = [
+ "cairo-sys-rs 0.15.1",
+ "gdk-pixbuf-sys 0.15.10",
+ "gio-sys 0.15.10",
+ "glib-sys 0.15.10",
+ "gobject-sys 0.15.10",
+ "libc",
+ "pango-sys 0.15.10",
+ "pkg-config",
+ "system-deps 6.0.4",
+]
+
+[[package]]
+name = "generic-array"
+version = "0.14.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "85649ca51fd72272d7821adaf274ad91c288277713d9c18820d8499a7ff69e9a"
+dependencies = [
+ "typenum",
+ "version_check",
+]
+
+[[package]]
+name = "gethostname"
+version = "0.2.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c1ebd34e35c46e00bb73e81363248d627782724609fe1b6396f553f68fe3862e"
+dependencies = [
+ "libc",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "getrandom"
+version = "0.2.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c85e1d9ab2eadba7e5040d4e09cbd6d072b76a557ad64e797c2cb9d4da21d7e4"
+dependencies = [
+ "cfg-if 1.0.0",
+ "libc",
+ "wasi",
+]
+
+[[package]]
+name = "gimli"
+version = "0.27.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ad0a93d233ebf96623465aad4046a8d3aa4da22d4f4beba5388838c8a434bbb4"
+
+[[package]]
+name = "gio"
+version = "0.9.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1fb60242bfff700772dae5d9e3a1f7aa2e4ebccf18b89662a16acb2822568561"
+dependencies = [
+ "bitflags",
+ "futures",
+ "futures-channel",
+ "futures-core",
+ "futures-io",
+ "futures-util",
+ "gio-sys 0.10.1",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "once_cell",
+ "thiserror",
+]
+
+[[package]]
+name = "gio"
+version = "0.15.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "68fdbc90312d462781a395f7a16d96a2b379bb6ef8cd6310a2df272771c4283b"
+dependencies = [
+ "bitflags",
+ "futures-channel",
+ "futures-core",
+ "futures-io",
+ "gio-sys 0.15.10",
+ "glib 0.15.12",
+ "libc",
+ "once_cell",
+ "thiserror",
+]
+
+[[package]]
+name = "gio-sys"
+version = "0.10.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5e24fb752f8f5d2cf6bbc2c606fd2bc989c81c5e2fe321ab974d54f8b6344eac"
+dependencies = [
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "system-deps 1.3.2",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "gio-sys"
+version = "0.15.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "32157a475271e2c4a023382e9cab31c4584ee30a97da41d3c4e9fdd605abcf8d"
+dependencies = [
+ "glib-sys 0.15.10",
+ "gobject-sys 0.15.10",
+ "libc",
+ "system-deps 6.0.4",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "glib"
+version = "0.10.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0c685013b7515e668f1b57a165b009d4d28cb139a8a989bbd699c10dad29d0c5"
+dependencies = [
+ "bitflags",
+ "futures-channel",
+ "futures-core",
+ "futures-executor",
+ "futures-task",
+ "futures-util",
+ "glib-macros 0.10.1",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "once_cell",
+]
+
+[[package]]
+name = "glib"
+version = "0.15.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "edb0306fbad0ab5428b0ca674a23893db909a98582969c9b537be4ced78c505d"
+dependencies = [
+ "bitflags",
+ "futures-channel",
+ "futures-core",
+ "futures-executor",
+ "futures-task",
+ "glib-macros 0.15.13",
+ "glib-sys 0.15.10",
+ "gobject-sys 0.15.10",
+ "libc",
+ "once_cell",
+ "smallvec",
+ "thiserror",
+]
+
+[[package]]
+name = "glib-macros"
+version = "0.10.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "41486a26d1366a8032b160b59065a59fb528530a46a49f627e7048fb8c064039"
+dependencies = [
+ "anyhow",
+ "heck 0.3.3",
+ "itertools",
+ "proc-macro-crate 0.1.5",
+ "proc-macro-error",
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "glib-macros"
+version = "0.15.13"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "10c6ae9f6fa26f4fb2ac16b528d138d971ead56141de489f8111e259b9df3c4a"
+dependencies = [
+ "anyhow",
+ "heck 0.4.1",
+ "proc-macro-crate 1.3.1",
+ "proc-macro-error",
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "glib-sys"
+version = "0.10.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c7e9b997a66e9a23d073f2b1abb4dbfc3925e0b8952f67efd8d9b6e168e4cdc1"
+dependencies = [
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "glib-sys"
+version = "0.15.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ef4b192f8e65e9cf76cbf4ea71fa8e3be4a0e18ffe3d68b8da6836974cc5bad4"
+dependencies = [
+ "libc",
+ "system-deps 6.0.4",
+]
+
+[[package]]
+name = "glob"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d2fabcfbdc87f4758337ca535fb41a6d701b65693ce38287d856d1674551ec9b"
+
+[[package]]
+name = "gobject-sys"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "952133b60c318a62bf82ee75b93acc7e84028a093e06b9e27981c2b6fe68218c"
+dependencies = [
+ "glib-sys 0.10.1",
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "gobject-sys"
+version = "0.15.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0d57ce44246becd17153bd035ab4d32cfee096a657fc01f2231c9278378d1e0a"
+dependencies = [
+ "glib-sys 0.15.10",
+ "libc",
+ "system-deps 6.0.4",
+]
+
+[[package]]
+name = "gstreamer"
+version = "0.16.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9ff5d0f7ff308ae37e6eb47b6ded17785bdea06e438a708cd09e0288c1862f33"
+dependencies = [
+ "bitflags",
+ "cfg-if 1.0.0",
+ "futures-channel",
+ "futures-core",
+ "futures-util",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "gstreamer-sys",
+ "libc",
+ "muldiv",
+ "num-rational",
+ "once_cell",
+ "paste",
+ "pretty-hex",
+ "thiserror",
+]
+
+[[package]]
+name = "gstreamer-app"
+version = "0.16.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "cc80888271338c3ede875d8cafc452eb207476ff5539dcbe0018a8f5b827af0e"
+dependencies = [
+ "bitflags",
+ "futures-core",
+ "futures-sink",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "gstreamer",
+ "gstreamer-app-sys",
+ "gstreamer-base",
+ "gstreamer-sys",
+ "libc",
+ "once_cell",
+]
+
+[[package]]
+name = "gstreamer-app-sys"
+version = "0.9.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "813f64275c9e7b33b828b9efcf9dfa64b95996766d4de996e84363ac65b87e3d"
+dependencies = [
+ "glib-sys 0.10.1",
+ "gstreamer-base-sys",
+ "gstreamer-sys",
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "gstreamer-base"
+version = "0.16.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bafd01c56f59cb10f4b5a10f97bb4bdf8c2b2784ae5b04da7e2d400cf6e6afcf"
+dependencies = [
+ "bitflags",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "gstreamer",
+ "gstreamer-base-sys",
+ "gstreamer-sys",
+ "libc",
+]
+
+[[package]]
+name = "gstreamer-base-sys"
+version = "0.9.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a4b7b6dc2d6e160a1ae28612f602bd500b3fa474ce90bf6bb2f08072682beef5"
+dependencies = [
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "gstreamer-sys",
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "gstreamer-sys"
+version = "0.9.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "fc1f154082d01af5718c5f8a8eb4f565a4ea5586ad8833a8fc2c2aa6844b601d"
+dependencies = [
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "gstreamer-video"
+version = "0.16.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f7bbb1485d87469849ec45c08e03c2f280d3ea20ff3c439d03185be54e3ce98e"
+dependencies = [
+ "bitflags",
+ "futures-channel",
+ "futures-util",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "gstreamer",
+ "gstreamer-base",
+ "gstreamer-base-sys",
+ "gstreamer-sys",
+ "gstreamer-video-sys",
+ "libc",
+ "once_cell",
+]
+
+[[package]]
+name = "gstreamer-video-sys"
+version = "0.9.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "92347e46438007d6a2386302125f62cb9df6769cdacb931af5c0f12c1ee21de4"
+dependencies = [
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "gstreamer-base-sys",
+ "gstreamer-sys",
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "gtk"
+version = "0.9.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2f022f2054072b3af07666341984562c8e626a79daa8be27b955d12d06a5ad6a"
+dependencies = [
+ "atk 0.9.0",
+ "bitflags",
+ "cairo-rs 0.9.1",
+ "cairo-sys-rs 0.10.0",
+ "cc",
+ "gdk 0.13.2",
+ "gdk-pixbuf 0.9.0",
+ "gdk-pixbuf-sys 0.10.0",
+ "gdk-sys 0.10.0",
+ "gio 0.9.1",
+ "gio-sys 0.10.1",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "gtk-sys 0.10.0",
+ "libc",
+ "once_cell",
+ "pango 0.9.1",
+ "pango-sys 0.10.0",
+ "pkg-config",
+]
+
+[[package]]
+name = "gtk"
+version = "0.15.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "92e3004a2d5d6d8b5057d2b57b3712c9529b62e82c77f25c1fecde1fd5c23bd0"
+dependencies = [
+ "atk 0.15.1",
+ "bitflags",
+ "cairo-rs 0.15.12",
+ "field-offset",
+ "futures-channel",
+ "gdk 0.15.4",
+ "gdk-pixbuf 0.15.11",
+ "gio 0.15.12",
+ "glib 0.15.12",
+ "gtk-sys 0.15.3",
+ "gtk3-macros",
+ "libc",
+ "once_cell",
+ "pango 0.15.10",
+ "pkg-config",
+]
+
+[[package]]
+name = "gtk-sys"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "89acda6f084863307d948ba64a4b1ef674e8527dddab147ee4cdcc194c880457"
+dependencies = [
+ "atk-sys 0.10.0",
+ "cairo-sys-rs 0.10.0",
+ "gdk-pixbuf-sys 0.10.0",
+ "gdk-sys 0.10.0",
+ "gio-sys 0.10.1",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "pango-sys 0.10.0",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "gtk-sys"
+version = "0.15.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d5bc2f0587cba247f60246a0ca11fe25fb733eabc3de12d1965fc07efab87c84"
+dependencies = [
+ "atk-sys 0.15.1",
+ "cairo-sys-rs 0.15.1",
+ "gdk-pixbuf-sys 0.15.10",
+ "gdk-sys 0.15.1",
+ "gio-sys 0.15.10",
+ "glib-sys 0.15.10",
+ "gobject-sys 0.15.10",
+ "libc",
+ "pango-sys 0.15.10",
+ "system-deps 6.0.4",
+]
+
+[[package]]
+name = "gtk3-macros"
+version = "0.15.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "684c0456c086e8e7e9af73ec5b84e35938df394712054550e81558d21c44ab0d"
+dependencies = [
+ "anyhow",
+ "proc-macro-crate 1.3.1",
+ "proc-macro-error",
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "hashbrown"
+version = "0.12.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8a9ee70c43aaf417c914396645a0fa852624801b24ebb7ae78fe8272889ac888"
+
+[[package]]
+name = "hbb_common"
+version = "0.1.0"
+dependencies = [
+ "anyhow",
+ "bytes",
+ "confy",
+ "directories-next",
+ "dirs-next",
+ "env_logger 0.9.3",
+ "filetime",
+ "futures",
+ "futures-util",
+ "lazy_static",
+ "log",
+ "mac_address",
+ "protobuf",
+ "protobuf-codegen-pure",
+ "quinn",
+ "rand 0.8.5",
+ "regex",
+ "serde 1.0.159",
+ "serde_derive",
+ "serde_json 1.0.95",
+ "socket2 0.3.19",
+ "sodiumoxide",
+ "tokio",
+ "tokio-socks",
+ "tokio-util 0.6.10",
+ "toml 0.5.11",
+ "winapi 0.3.9",
+ "zstd",
+]
+
+[[package]]
+name = "heck"
+version = "0.3.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6d621efb26863f0e9924c6ac577e8275e5e6b77455db64ffa6c65c904e9e132c"
+dependencies = [
+ "unicode-segmentation",
+]
+
+[[package]]
+name = "heck"
+version = "0.4.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "95505c38b4572b2d910cecb0281560f54b440a19336cbbcb27bf6ce6adc6f5a8"
+
+[[package]]
+name = "hermit-abi"
+version = "0.1.19"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "62b467343b94ba476dcb2500d242dadbb39557df889310ac77c5d99100aaac33"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "hermit-abi"
+version = "0.2.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ee512640fe35acbfb4bb779db6f0d80704c2cacfa2e39b601ef3e3f47d1ae4c7"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "hermit-abi"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "fed44880c466736ef9a5c5b5facefb5ed0785676d0c02d612db14e54f0d84286"
+
+[[package]]
+name = "hound"
+version = "3.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4d13cdbd5dbb29f9c88095bbdc2590c9cba0d0a1269b983fef6b2cdd7e9f4db1"
+
+[[package]]
+name = "humantime"
+version = "2.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9a3a5bfb195931eeb336b2a7b4d761daec841b97f947d34394601737a7bba5e4"
+
+[[package]]
+name = "iana-time-zone"
+version = "0.1.56"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0722cd7114b7de04316e7ea5456a0bbb20e4adb46fd27a3697adb812cff0f37c"
+dependencies = [
+ "android_system_properties",
+ "core-foundation-sys",
+ "iana-time-zone-haiku",
+ "js-sys",
+ "wasm-bindgen",
+ "windows",
+]
+
+[[package]]
+name = "iana-time-zone-haiku"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0703ae284fc167426161c2e3f1da3ea71d94b21bedbcc9494e92b28e334e3dca"
+dependencies = [
+ "cxx",
+ "cxx-build",
+]
+
+[[package]]
+name = "ident_case"
+version = "1.0.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b9e0384b61958566e926dc50660321d12159025e767c18e043daf26b70104c39"
+
+[[package]]
+name = "image"
+version = "0.23.14"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "24ffcb7e7244a9bf19d35bf2883b9c080c4ced3c07a9895572178cdb8f13f6a1"
+dependencies = [
+ "bytemuck",
+ "byteorder",
+ "color_quant",
+ "num-iter",
+ "num-rational",
+ "num-traits 0.2.15",
+ "png",
+ "tiff",
+]
+
+[[package]]
+name = "include_dir"
+version = "0.7.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "18762faeff7122e89e0857b02f7ce6fcc0d101d5e9ad2ad7846cc01d61b7f19e"
+dependencies = [
+ "include_dir_macros",
+]
+
+[[package]]
+name = "include_dir_macros"
+version = "0.7.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b139284b5cf57ecfa712bcc66950bb635b31aff41c188e8a4cfc758eca374a3f"
+dependencies = [
+ "proc-macro2",
+ "quote",
+]
+
+[[package]]
+name = "indexmap"
+version = "1.9.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bd070e393353796e801d209ad339e89596eb4c8d430d18ede6a1cced8fafbd99"
+dependencies = [
+ "autocfg 1.1.0",
+ "hashbrown",
+]
+
+[[package]]
+name = "instant"
+version = "0.1.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7a5bbe824c507c5da5956355e86a746d82e0e1464f65d862cc5e71da70e94b2c"
+dependencies = [
+ "cfg-if 1.0.0",
+]
+
+[[package]]
+name = "io-lifetimes"
+version = "1.0.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9c66c74d2ae7e79a5a8f7ac924adbe38ee42a859c6539ad869eb51f0b52dc220"
+dependencies = [
+ "hermit-abi 0.3.1",
+ "libc",
+ "windows-sys 0.48.0",
+]
+
+[[package]]
+name = "iovec"
+version = "0.1.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b2b3ea6ff95e175473f8ffe6a7eb7c00d054240321b84c57051175fe3c1e075e"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "itertools"
+version = "0.9.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "284f18f85651fe11e8a991b2adb42cb078325c996ed026d994719efcfca1d54b"
+dependencies = [
+ "either",
+]
+
+[[package]]
+name = "itoa"
+version = "0.3.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8324a32baf01e2ae060e9de58ed0bc2320c9a2833491ee36cd3b4c414de4db8c"
+
+[[package]]
+name = "itoa"
+version = "1.0.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "453ad9f582a441959e5f0d088b02ce04cfe8d51a8eaf077f12ac6d3e94164ca6"
+
+[[package]]
+name = "jni"
+version = "0.19.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c6df18c2e3db7e453d3c6ac5b3e9d5182664d28788126d39b91f2d1e22b017ec"
+dependencies = [
+ "cesu8",
+ "combine",
+ "jni-sys",
+ "log",
+ "thiserror",
+ "walkdir",
+]
+
+[[package]]
+name = "jni-sys"
+version = "0.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8eaf4bc02d17cbdd7ff4c7438cafcdf7fb9a4613313ad11b4f8fefe7d3fa0130"
+
+[[package]]
+name = "jobserver"
+version = "0.1.26"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "936cfd212a0155903bcbc060e316fb6cc7cbf2e1907329391ebadc1fe0ce77c2"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "jpeg-decoder"
+version = "0.1.22"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "229d53d58899083193af11e15917b5640cd40b29ff475a1fe4ef725deb02d0f2"
+
+[[package]]
+name = "js-sys"
+version = "0.3.61"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "445dde2150c55e483f3d8416706b97ec8e8237c307e5b7b4b8dd15e6af2a0730"
+dependencies = [
+ "wasm-bindgen",
+]
+
+[[package]]
+name = "kernel32-sys"
+version = "0.2.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7507624b29483431c0ba2d82aece8ca6cdba9382bff4ddd0f7490560c056098d"
+dependencies = [
+ "winapi 0.2.8",
+ "winapi-build",
+]
+
+[[package]]
+name = "lazy_static"
+version = "1.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e2abad23fbc42b3700f2f279844dc832adb2b2eb069b2df918f455c4e18cc646"
+
+[[package]]
+name = "lazycell"
+version = "1.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "830d08ce1d1d941e6b30645f1a0eb5643013d835ce3779a5fc208261dbe10f55"
+
+[[package]]
+name = "libappindicator"
+version = "0.6.1"
+source = "git+https://github.com/liyue201/libappindicator-rs#3763cfd629dd90050af1feafa643cbfca0bf487e"
+dependencies = [
+ "glib 0.10.3",
+ "gtk 0.9.2",
+ "gtk-sys 0.10.0",
+ "libappindicator-sys 0.6.1",
+ "log",
+]
+
+[[package]]
+name = "libappindicator"
+version = "0.7.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "db2d3cb96d092b4824cb306c9e544c856a4cb6210c1081945187f7f1924b47e8"
+dependencies = [
+ "glib 0.15.12",
+ "gtk 0.15.5",
+ "gtk-sys 0.15.3",
+ "libappindicator-sys 0.7.3",
+ "log",
+]
+
+[[package]]
+name = "libappindicator-sys"
+version = "0.6.1"
+source = "git+https://github.com/liyue201/libappindicator-rs#3763cfd629dd90050af1feafa643cbfca0bf487e"
+dependencies = [
+ "gtk-sys 0.10.0",
+ "pkg-config",
+]
+
+[[package]]
+name = "libappindicator-sys"
+version = "0.7.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f1b3b6681973cea8cc3bce7391e6d7d5502720b80a581c9a95c9cbaf592826aa"
+dependencies = [
+ "gtk-sys 0.15.3",
+ "libloading",
+ "once_cell",
+]
+
+[[package]]
+name = "libc"
+version = "0.2.141"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3304a64d199bb964be99741b7a14d26972741915b3649639149b2479bb46f4b5"
+
+[[package]]
+name = "libdbus-sys"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9f8d7ae751e1cb825c840ae5e682f59b098cdfd213c350ac268b61449a5f58a0"
+dependencies = [
+ "pkg-config",
+]
+
+[[package]]
+name = "libloading"
+version = "0.7.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b67380fd3b2fbe7527a606e18729d21c6f3951633d0500574c4dc22d2d638b9f"
+dependencies = [
+ "cfg-if 1.0.0",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "libpulse-binding"
+version = "2.27.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1745b20bfc194ac12ef828f144f0ec2d4a7fe993281fa3567a0bd4969aee6890"
+dependencies = [
+ "bitflags",
+ "libc",
+ "libpulse-sys",
+ "num-derive",
+ "num-traits 0.2.15",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "libpulse-simple-binding"
+version = "2.27.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5ced94199e6e44133431374e4043f34e1f0697ebfb7b7d6c244a65bfaedf0e31"
+dependencies = [
+ "libpulse-binding",
+ "libpulse-simple-sys",
+ "libpulse-sys",
+]
+
+[[package]]
+name = "libpulse-simple-sys"
+version = "1.20.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "84e423d9c619c908ce9b4916080e65ab586ca55b8c4939379f15e6e72fb43842"
+dependencies = [
+ "libpulse-sys",
+ "pkg-config",
+]
+
+[[package]]
+name = "libpulse-sys"
+version = "1.20.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2191e6880818d1df4cf72eac8e91dce7a5a52ba0da4b2a5cdafabc22b937eadb"
+dependencies = [
+ "libc",
+ "num-derive",
+ "num-traits 0.2.15",
+ "pkg-config",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "libsamplerate-sys"
+version = "0.1.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "28853b399f78f8281cd88d333b54a63170c4275f6faea66726a2bea5cca72e0d"
+dependencies = [
+ "cmake",
+]
+
+[[package]]
+name = "libsodium-sys"
+version = "0.2.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6b779387cd56adfbc02ea4a668e704f729be8d6a6abd2c27ca5ee537849a92fd"
+dependencies = [
+ "cc",
+ "libc",
+ "pkg-config",
+ "walkdir",
+]
+
+[[package]]
+name = "link-cplusplus"
+version = "1.0.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ecd207c9c713c34f95a097a5b029ac2ce6010530c7b49d7fea24d977dede04f5"
+dependencies = [
+ "cc",
+]
+
+[[package]]
+name = "linux-raw-sys"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d59d8c75012853d2e872fb56bc8a2e53718e2cafe1a4c823143141c6d90c322f"
+
+[[package]]
+name = "lock_api"
+version = "0.4.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "435011366fe56583b16cf956f9df0095b405b82d76425bc8981c0e22e60ec4df"
+dependencies = [
+ "autocfg 1.1.0",
+ "scopeguard",
+]
+
+[[package]]
+name = "log"
+version = "0.4.17"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "abb12e687cfb44aa40f41fc3978ef76448f9b6038cad6aef4259d3c095a2382e"
+dependencies = [
+ "cfg-if 1.0.0",
+]
+
+[[package]]
+name = "mac_address"
+version = "1.1.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b238e3235c8382b7653c6408ed1b08dd379bdb9fdf990fb0bbae3db2cc0ae963"
+dependencies = [
+ "nix 0.23.2",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "mach"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b823e83b2affd8f40a9ee8c29dbc56404c1e34cd2710921f2801e2cf29527afa"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "machine-uid"
+version = "0.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1f1595709b0a7386bcd56ba34d250d626e5503917d05d32cdccddcd68603e212"
+dependencies = [
+ "winreg 0.6.2",
+]
+
+[[package]]
+name = "magnum-opus"
+version = "0.4.0"
+source = "git+https://github.com/TheRadioGuy/magnum-opus#171e1d021004626f7444d1e39b98f50bc3cb2604"
+dependencies = [
+ "libc",
+ "opusic-sys",
+]
+
+[[package]]
+name = "malloc_buf"
+version = "0.0.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "62bb907fe88d54d8d9ce32a3cceab4218ed2f6b7d35617cafe9adf84e43919cb"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "memchr"
+version = "2.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2dffe52ecf27772e601905b7522cb4ef790d2cc203488bbd0e2fe85fcb74566d"
+
+[[package]]
+name = "memoffset"
+version = "0.6.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5aa361d4faea93603064a027415f07bd8e1d5c88c9fbf68bf56a285428fd79ce"
+dependencies = [
+ "autocfg 1.1.0",
+]
+
+[[package]]
+name = "memoffset"
+version = "0.8.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d61c719bcfbcf5d62b3a09efa6088de8c54bc0bfcd3ea7ae39fcc186108b8de1"
+dependencies = [
+ "autocfg 1.1.0",
+]
+
+[[package]]
+name = "minimal-lexical"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a"
+
+[[package]]
+name = "miniz_oxide"
+version = "0.3.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "791daaae1ed6889560f8c4359194f56648355540573244a5448a83ba1ecc7435"
+dependencies = [
+ "adler32",
+]
+
+[[package]]
+name = "miniz_oxide"
+version = "0.4.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a92518e98c078586bc6c934028adcca4c92a53d6a958196de835170a01d84e4b"
+dependencies = [
+ "adler",
+ "autocfg 1.1.0",
+]
+
+[[package]]
+name = "miniz_oxide"
+version = "0.6.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b275950c28b37e794e8c55d88aeb5e139d0ce23fdbbeda68f8d7174abdf9e8fa"
+dependencies = [
+ "adler",
+]
+
+[[package]]
+name = "mio"
+version = "0.6.23"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4afd66f5b91bf2a3bc13fad0e21caedac168ca4c707504e75585648ae80e4cc4"
+dependencies = [
+ "cfg-if 0.1.10",
+ "fuchsia-zircon",
+ "fuchsia-zircon-sys",
+ "iovec",
+ "kernel32-sys",
+ "libc",
+ "log",
+ "miow 0.2.2",
+ "net2",
+ "slab",
+ "winapi 0.2.8",
+]
+
+[[package]]
+name = "mio"
+version = "0.8.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5b9d9a46eff5b4ff64b45a9e316a6d1e0bc719ef429cbec4dc630684212bfdf9"
+dependencies = [
+ "libc",
+ "log",
+ "wasi",
+ "windows-sys 0.45.0",
+]
+
+[[package]]
+name = "mio-named-pipes"
+version = "0.1.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0840c1c50fd55e521b247f949c241c9997709f23bd7f023b9762cd561e935656"
+dependencies = [
+ "log",
+ "mio 0.6.23",
+ "miow 0.3.7",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "miow"
+version = "0.2.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ebd808424166322d4a38da87083bfddd3ac4c131334ed55856112eb06d46944d"
+dependencies = [
+ "kernel32-sys",
+ "net2",
+ "winapi 0.2.8",
+ "ws2_32-sys",
+]
+
+[[package]]
+name = "miow"
+version = "0.3.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b9f1c5b025cda876f66ef43a113f91ebc9f4ccef34843000e0adf6ebbab84e21"
+dependencies = [
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "miow"
+version = "0.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a7377f7792b3afb6a3cba68daa54ca23c032137010460d667fda53a8d66be00e"
+dependencies = [
+ "windows-sys 0.28.0",
+]
+
+[[package]]
+name = "muldiv"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0419348c027fa7be448d2ae7ea0e4e04c2334c31dc4e74ab29f00a2a7ca69204"
+
+[[package]]
+name = "ndk"
+version = "0.6.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2032c77e030ddee34a6787a64166008da93f6a352b629261d0fee232b8742dd4"
+dependencies = [
+ "bitflags",
+ "jni-sys",
+ "ndk-sys",
+ "num_enum",
+ "thiserror",
+]
+
+[[package]]
+name = "ndk-context"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "27b02d87554356db9e9a873add8782d4ea6e3e58ea071a9adb9a2e8ddb884a8b"
+
+[[package]]
+name = "ndk-glue"
+version = "0.6.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0d0c4a7b83860226e6b4183edac21851f05d5a51756e97a1144b7f5a6b63e65f"
+dependencies = [
+ "lazy_static",
+ "libc",
+ "log",
+ "ndk",
+ "ndk-context",
+ "ndk-macro",
+ "ndk-sys",
+]
+
+[[package]]
+name = "ndk-macro"
+version = "0.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0df7ac00c4672f9d5aece54ee3347520b7e20f158656c7db2e6de01902eb7a6c"
+dependencies = [
+ "darling",
+ "proc-macro-crate 1.3.1",
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "ndk-sys"
+version = "0.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6e5a6ae77c8ee183dcbbba6150e2e6b9f3f4196a7666c02a715a95692ec1fa97"
+dependencies = [
+ "jni-sys",
+]
+
+[[package]]
+name = "net2"
+version = "0.2.38"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "74d0df99cfcd2530b2e694f6e17e7f37b8e26bb23983ac530c0c97408837c631"
+dependencies = [
+ "cfg-if 0.1.10",
+ "libc",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "nix"
+version = "0.22.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e4916f159ed8e5de0082076562152a76b7a1f64a01fd9d1e0fea002c37624faf"
+dependencies = [
+ "bitflags",
+ "cc",
+ "cfg-if 1.0.0",
+ "libc",
+ "memoffset 0.6.5",
+]
+
+[[package]]
+name = "nix"
+version = "0.23.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8f3790c00a0150112de0f4cd161e3d7fc4b2d8a5542ffc35f099a2562aecb35c"
+dependencies = [
+ "bitflags",
+ "cc",
+ "cfg-if 1.0.0",
+ "libc",
+ "memoffset 0.6.5",
+]
+
+[[package]]
+name = "nix"
+version = "0.26.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bfdda3d196821d6af13126e40375cdf7da646a96114af134d5f417a9a1dc8e1a"
+dependencies = [
+ "bitflags",
+ "cfg-if 1.0.0",
+ "libc",
+ "static_assertions",
+]
+
+[[package]]
+name = "nom"
+version = "7.1.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d273983c5a657a70a3e8f2a01329822f3b8c8172b73826411a55751e404a0a4a"
+dependencies = [
+ "memchr",
+ "minimal-lexical",
+]
+
+[[package]]
+name = "ntapi"
+version = "0.3.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c28774a7fd2fbb4f0babd8237ce554b73af68021b5f695a3cebd6c59bac0980f"
+dependencies = [
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "num-complex"
+version = "0.4.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "02e0d21255c828d6f128a1e41534206671e8c3ea0c62f32291e808dc82cff17d"
+dependencies = [
+ "num-traits 0.2.15",
+]
+
+[[package]]
+name = "num-derive"
+version = "0.3.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "876a53fff98e03a936a674b29568b0e605f06b29372c2489ff4de23f1949743d"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "num-integer"
+version = "0.1.45"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "225d3389fb3509a24c93f5c29eb6bde2586b98d9f016636dff58d7c6f7569cd9"
+dependencies = [
+ "autocfg 1.1.0",
+ "num-traits 0.2.15",
+]
+
+[[package]]
+name = "num-iter"
+version = "0.1.43"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7d03e6c028c5dc5cac6e2dec0efda81fc887605bb3d884578bb6d6bf7514e252"
+dependencies = [
+ "autocfg 1.1.0",
+ "num-integer",
+ "num-traits 0.2.15",
+]
+
+[[package]]
+name = "num-rational"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "12ac428b1cb17fce6f731001d307d351ec70a6d202fc2e60f7d4c5e42d8f4f07"
+dependencies = [
+ "autocfg 1.1.0",
+ "num-integer",
+ "num-traits 0.2.15",
+]
+
+[[package]]
+name = "num-traits"
+version = "0.1.43"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "92e5113e9fd4cc14ded8e499429f396a20f98c772a47cc8622a736e1ec843c31"
+dependencies = [
+ "num-traits 0.2.15",
+]
+
+[[package]]
+name = "num-traits"
+version = "0.2.15"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "578ede34cf02f8924ab9447f50c28075b4d3e5b269972345e7e0372b38c6cdcd"
+dependencies = [
+ "autocfg 1.1.0",
+]
+
+[[package]]
+name = "num_cpus"
+version = "1.15.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0fac9e2da13b5eb447a6ce3d392f23a29d8694bff781bf03a16cd9ac8697593b"
+dependencies = [
+ "hermit-abi 0.2.6",
+ "libc",
+]
+
+[[package]]
+name = "num_enum"
+version = "0.5.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1f646caf906c20226733ed5b1374287eb97e3c2a5c227ce668c1f2ce20ae57c9"
+dependencies = [
+ "num_enum_derive",
+]
+
+[[package]]
+name = "num_enum_derive"
+version = "0.5.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "dcbff9bc912032c62bf65ef1d5aea88983b420f4f839db1e9b0c281a25c9c799"
+dependencies = [
+ "proc-macro-crate 1.3.1",
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "num_threads"
+version = "0.1.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2819ce041d2ee131036f4fc9d6ae7ae125a3a40e97ba64d04fe799ad9dabbb44"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "objc"
+version = "0.2.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "915b1b472bc21c53464d6c8461c9d3af805ba1ef837e1cac254428f4a77177b1"
+dependencies = [
+ "malloc_buf",
+]
+
+[[package]]
+name = "objc-foundation"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1add1b659e36c9607c7aab864a76c7a4c2760cd0cd2e120f3fb8b952c7e22bf9"
+dependencies = [
+ "block",
+ "objc",
+ "objc_id",
+]
+
+[[package]]
+name = "objc_id"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c92d4ddb4bd7b50d730c215ff871754d0da6b2178849f8a2a2ab69712d0c073b"
+dependencies = [
+ "objc",
+]
+
+[[package]]
+name = "object"
+version = "0.30.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ea86265d3d3dcb6a27fc51bd29a4bf387fae9d2986b823079d4986af253eb439"
+dependencies = [
+ "memchr",
+]
+
+[[package]]
+name = "oboe"
+version = "0.4.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "27f63c358b4fa0fbcfefd7c8be5cfc39c08ce2389f5325687e7762a48d30a5c1"
+dependencies = [
+ "jni",
+ "ndk",
+ "ndk-context",
+ "num-derive",
+ "num-traits 0.2.15",
+ "oboe-sys",
+]
+
+[[package]]
+name = "oboe-sys"
+version = "0.4.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3370abb7372ed744232c12954d920d1a40f1c4686de9e79e800021ef492294bd"
+dependencies = [
+ "cc",
+]
+
+[[package]]
+name = "once_cell"
+version = "1.17.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b7e5500299e16ebb147ae15a00a942af264cf3688f47923b8fc2cd5858f23ad3"
+
+[[package]]
+name = "openssl-probe"
+version = "0.1.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ff011a302c396a5197692431fc1948019154afc178baf7d8e37367442a4601cf"
+
+[[package]]
+name = "opusic-sys"
+version = "0.3.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5eace752ce07a037241dba8f02c654799f051e431b27028056bcb480e83b54f5"
+dependencies = [
+ "cmake",
+ "libc",
+]
+
+[[package]]
+name = "os_str_bytes"
+version = "6.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ceedf44fb00f2d1984b0bc98102627ce622e083e49a5bacdb3e514fa4238e267"
+
+[[package]]
+name = "padlock"
+version = "0.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c10569378a1dacd9f30dbe7ae49e054d2c45dc2f8ee49899903e09c3924e8b6f"
+
+[[package]]
+name = "pango"
+version = "0.9.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9937068580bebd8ced19975938573803273ccbcbd598c58d4906efd4ac87c438"
+dependencies = [
+ "bitflags",
+ "glib 0.10.3",
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "once_cell",
+ "pango-sys 0.10.0",
+]
+
+[[package]]
+name = "pango"
+version = "0.15.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "22e4045548659aee5313bde6c582b0d83a627b7904dd20dc2d9ef0895d414e4f"
+dependencies = [
+ "bitflags",
+ "glib 0.15.12",
+ "libc",
+ "once_cell",
+ "pango-sys 0.15.10",
+]
+
+[[package]]
+name = "pango-sys"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "24d2650c8b62d116c020abd0cea26a4ed96526afda89b1c4ea567131fdefc890"
+dependencies = [
+ "glib-sys 0.10.1",
+ "gobject-sys 0.10.0",
+ "libc",
+ "system-deps 1.3.2",
+]
+
+[[package]]
+name = "pango-sys"
+version = "0.15.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d2a00081cde4661982ed91d80ef437c20eacaf6aa1a5962c0279ae194662c3aa"
+dependencies = [
+ "glib-sys 0.15.10",
+ "gobject-sys 0.15.10",
+ "libc",
+ "system-deps 6.0.4",
+]
+
+[[package]]
+name = "parity-tokio-ipc"
+version = "0.7.3-1"
+source = "git+https://github.com/open-trade/parity-tokio-ipc#20b2895910161605210657f3e751edd55321f698"
+dependencies = [
+ "futures",
+ "libc",
+ "log",
+ "mio-named-pipes",
+ "miow 0.4.0",
+ "rand 0.8.5",
+ "tokio",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "parking_lot"
+version = "0.11.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7d17b78036a60663b797adeaee46f5c9dfebb86948d1255007a1d6be0271ff99"
+dependencies = [
+ "instant",
+ "lock_api",
+ "parking_lot_core 0.8.6",
+]
+
+[[package]]
+name = "parking_lot"
+version = "0.12.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3742b2c103b9f06bc9fff0a37ff4912935851bee6d36f3c02bcc755bcfec228f"
+dependencies = [
+ "lock_api",
+ "parking_lot_core 0.9.7",
+]
+
+[[package]]
+name = "parking_lot_core"
+version = "0.8.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "60a2cfe6f0ad2bfc16aefa463b497d5c7a5ecd44a23efa72aa342d90177356dc"
+dependencies = [
+ "cfg-if 1.0.0",
+ "instant",
+ "libc",
+ "redox_syscall 0.2.16",
+ "smallvec",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "parking_lot_core"
+version = "0.9.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9069cbb9f99e3a5083476ccb29ceb1de18b9118cafa53e90c9551235de2b9521"
+dependencies = [
+ "cfg-if 1.0.0",
+ "libc",
+ "redox_syscall 0.2.16",
+ "smallvec",
+ "windows-sys 0.45.0",
+]
+
+[[package]]
+name = "paste"
+version = "1.0.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9f746c4065a8fa3fe23974dd82f15431cc8d40779821001404d10d2e79ca7d79"
+
+[[package]]
+name = "peeking_take_while"
+version = "0.1.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "19b17cddbe7ec3f8bc800887bab5e717348c95ea2ca0b1bf0837fb964dc67099"
+
+[[package]]
+name = "phf"
+version = "0.7.24"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b3da44b85f8e8dfaec21adae67f95d93244b2ecf6ad2a692320598dcc8e6dd18"
+dependencies = [
+ "phf_shared",
+]
+
+[[package]]
+name = "phf_codegen"
+version = "0.7.24"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b03e85129e324ad4166b06b2c7491ae27fe3ec353af72e72cd1654c7225d517e"
+dependencies = [
+ "phf_generator",
+ "phf_shared",
+]
+
+[[package]]
+name = "phf_generator"
+version = "0.7.24"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "09364cc93c159b8b06b1f4dd8a4398984503483891b0c26b867cf431fb132662"
+dependencies = [
+ "phf_shared",
+ "rand 0.6.5",
+]
+
+[[package]]
+name = "phf_shared"
+version = "0.7.24"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "234f71a15de2288bcb7e3b6515828d22af7ec8598ee6d24c3b526fa0a80b67a0"
+dependencies = [
+ "siphasher",
+]
+
+[[package]]
+name = "pin-project"
+version = "1.0.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ad29a609b6bcd67fee905812e544992d216af9d755757c05ed2d0e15a74c6ecc"
+dependencies = [
+ "pin-project-internal",
+]
+
+[[package]]
+name = "pin-project-internal"
+version = "1.0.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "069bdb1e05adc7a8990dce9cc75370895fbe4e3d58b9b73bf1aee56359344a55"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "pin-project-lite"
+version = "0.2.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e0a7ae3ac2f1173085d398531c705756c94a4c56843785df85a60c1a0afac116"
+
+[[package]]
+name = "pin-utils"
+version = "0.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8b870d8c151b6f2fb93e84a13146138f05d02ed11c7e7c54f8826aaaf7c9f184"
+
+[[package]]
+name = "pkg-config"
+version = "0.3.26"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6ac9a59f73473f1b8d852421e59e64809f025994837ef743615c6d0c5b305160"
+
+[[package]]
+name = "png"
+version = "0.16.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3c3287920cb847dee3de33d301c463fba14dda99db24214ddf93f83d3021f4c6"
+dependencies = [
+ "bitflags",
+ "crc32fast",
+ "deflate",
+ "miniz_oxide 0.3.7",
+]
+
+[[package]]
+name = "ppv-lite86"
+version = "0.2.17"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5b40af805b3121feab8a3c29f04d8ad262fa8e0561883e7653e024ae4479e6de"
+
+[[package]]
+name = "pretty-hex"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bc5c99d529f0d30937f6f4b8a86d988047327bb88d04d2c4afc356de74722131"
+
+[[package]]
+name = "primal-check"
+version = "0.3.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9df7f93fd637f083201473dab4fee2db4c429d32e55e3299980ab3957ab916a0"
+dependencies = [
+ "num-integer",
+]
+
+[[package]]
+name = "proc-macro-crate"
+version = "0.1.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1d6ea3c4595b96363c13943497db34af4460fb474a95c43f4446ad341b8c9785"
+dependencies = [
+ "toml 0.5.11",
+]
+
+[[package]]
+name = "proc-macro-crate"
+version = "1.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7f4c021e1093a56626774e81216a4ce732a735e5bad4868a03f3ed65ca0c3919"
+dependencies = [
+ "once_cell",
+ "toml_edit",
+]
+
+[[package]]
+name = "proc-macro-error"
+version = "1.0.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "da25490ff9892aab3fcf7c36f08cfb902dd3e71ca0f9f9517bea02a73a5ce38c"
+dependencies = [
+ "proc-macro-error-attr",
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+ "version_check",
+]
+
+[[package]]
+name = "proc-macro-error-attr"
+version = "1.0.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a1be40180e52ecc98ad80b184934baf3d0d29f979574e439af5a55274b35f869"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "version_check",
+]
+
+[[package]]
+name = "proc-macro2"
+version = "1.0.56"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2b63bdb0cd06f1f4dedf69b254734f9b45af66e4a031e42a7480257d9898b435"
+dependencies = [
+ "unicode-ident",
+]
+
+[[package]]
+name = "protobuf"
+version = "3.0.0-alpha.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9d5ef59c35c7472ce5e1b6c5924b87585143d1fc2cf39eae0009bba6c4df62f1"
+
+[[package]]
+name = "protobuf-codegen"
+version = "3.0.0-alpha.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "89100ee819f69b77a4cab389fec9dd155a305af4c615e6413ec1ef9341f333ef"
+dependencies = [
+ "anyhow",
+ "protobuf",
+ "protobuf-parse",
+ "thiserror",
+]
+
+[[package]]
+name = "protobuf-codegen-pure"
+version = "3.0.0-alpha.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "79453e74d08190551e821533ee42c447f9e21ca26f83520e120e6e8af27f6879"
+dependencies = [
+ "anyhow",
+ "protobuf",
+ "protobuf-codegen",
+ "protobuf-parse",
+ "thiserror",
+]
+
+[[package]]
+name = "protobuf-parse"
+version = "3.0.0-alpha.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c265ffc69976efc3056955b881641add3186ad0be893ef10622482d80d1d2b68"
+dependencies = [
+ "anyhow",
+ "protobuf",
+ "protoc",
+ "tempfile",
+ "thiserror",
+]
+
+[[package]]
+name = "protoc"
+version = "3.0.0-alpha.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1f1f8b318a54d18fbe542513331e058f4f8ce6502e542e057c50c7e5e803fdab"
+dependencies = [
+ "anyhow",
+ "log",
+ "thiserror",
+ "which 4.4.0",
+]
+
+[[package]]
+name = "quest"
+version = "0.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "556af5f5c953a2ee13f45753e581a38f9778e6551bc3ccc56d90b14628fe59d8"
+dependencies = [
+ "cfg-if 0.1.10",
+ "rpassword 2.1.0",
+ "tempfile",
+ "termios",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "quick-xml"
+version = "0.22.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8533f14c8382aaad0d592c812ac3b826162128b65662331e1127b45c3d18536b"
+dependencies = [
+ "memchr",
+]
+
+[[package]]
+name = "quinn"
+version = "0.8.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5b435e71d9bfa0d8889927231970c51fb89c58fa63bffcab117c9c7a41e5ef8f"
+dependencies = [
+ "bytes",
+ "futures-channel",
+ "futures-util",
+ "fxhash",
+ "quinn-proto",
+ "quinn-udp",
+ "rustls",
+ "thiserror",
+ "tokio",
+ "tracing",
+ "webpki",
+]
+
+[[package]]
+name = "quinn-proto"
+version = "0.8.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3fce546b9688f767a57530652488420d419a8b1f44a478b451c3d1ab6d992a55"
+dependencies = [
+ "bytes",
+ "fxhash",
+ "rand 0.8.5",
+ "ring",
+ "rustls",
+ "rustls-native-certs",
+ "rustls-pemfile 0.2.1",
+ "slab",
+ "thiserror",
+ "tinyvec",
+ "tracing",
+ "webpki",
+]
+
+[[package]]
+name = "quinn-udp"
+version = "0.1.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b07946277141531aea269befd949ed16b2c85a780ba1043244eda0969e538e54"
+dependencies = [
+ "futures-util",
+ "libc",
+ "quinn-proto",
+ "socket2 0.4.9",
+ "tokio",
+ "tracing",
+]
+
+[[package]]
+name = "quote"
+version = "1.0.26"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4424af4bf778aae2051a77b60283332f386554255d722233d09fbfc7e30da2fc"
+dependencies = [
+ "proc-macro2",
+]
+
+[[package]]
+name = "rand"
+version = "0.6.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6d71dacdc3c88c1fde3885a3be3fbab9f35724e6ce99467f7d9c5026132184ca"
+dependencies = [
+ "autocfg 0.1.8",
+ "libc",
+ "rand_chacha 0.1.1",
+ "rand_core 0.4.2",
+ "rand_hc",
+ "rand_isaac",
+ "rand_jitter",
+ "rand_os",
+ "rand_pcg",
+ "rand_xorshift",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "rand"
+version = "0.8.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "34af8d1a0e25924bc5b7c43c079c942339d8f0a8b57c39049bef581b46327404"
+dependencies = [
+ "libc",
+ "rand_chacha 0.3.1",
+ "rand_core 0.6.4",
+]
+
+[[package]]
+name = "rand_chacha"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "556d3a1ca6600bfcbab7c7c91ccb085ac7fbbcd70e008a98742e7847f4f7bcef"
+dependencies = [
+ "autocfg 0.1.8",
+ "rand_core 0.3.1",
+]
+
+[[package]]
+name = "rand_chacha"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e6c10a63a0fa32252be49d21e7709d4d4baf8d231c2dbce1eaa8141b9b127d88"
+dependencies = [
+ "ppv-lite86",
+ "rand_core 0.6.4",
+]
+
+[[package]]
+name = "rand_core"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7a6fdeb83b075e8266dcc8762c22776f6877a63111121f5f8c7411e5be7eed4b"
+dependencies = [
+ "rand_core 0.4.2",
+]
+
+[[package]]
+name = "rand_core"
+version = "0.4.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9c33a3c44ca05fa6f1807d8e6743f3824e8509beca625669633be0acbdf509dc"
+
+[[package]]
+name = "rand_core"
+version = "0.6.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ec0be4795e2f6a28069bec0b5ff3e2ac9bafc99e6a9a7dc3547996c5c816922c"
+dependencies = [
+ "getrandom",
+]
+
+[[package]]
+name = "rand_hc"
+version = "0.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7b40677c7be09ae76218dc623efbf7b18e34bced3f38883af07bb75630a21bc4"
+dependencies = [
+ "rand_core 0.3.1",
+]
+
+[[package]]
+name = "rand_isaac"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ded997c9d5f13925be2a6fd7e66bf1872597f759fd9dd93513dd7e92e5a5ee08"
+dependencies = [
+ "rand_core 0.3.1",
+]
+
+[[package]]
+name = "rand_jitter"
+version = "0.1.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1166d5c91dc97b88d1decc3285bb0a99ed84b05cfd0bc2341bdf2d43fc41e39b"
+dependencies = [
+ "libc",
+ "rand_core 0.4.2",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "rand_os"
+version = "0.1.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7b75f676a1e053fc562eafbb47838d67c84801e38fc1ba459e8f180deabd5071"
+dependencies = [
+ "cloudabi",
+ "fuchsia-cprng",
+ "libc",
+ "rand_core 0.4.2",
+ "rdrand",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "rand_pcg"
+version = "0.1.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "abf9b09b01790cfe0364f52bf32995ea3c39f4d2dd011eac241d2914146d0b44"
+dependencies = [
+ "autocfg 0.1.8",
+ "rand_core 0.4.2",
+]
+
+[[package]]
+name = "rand_xorshift"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "cbf7e9e623549b0e21f6e97cf8ecf247c1a8fd2e8a992ae265314300b2455d5c"
+dependencies = [
+ "rand_core 0.3.1",
+]
+
+[[package]]
+name = "rayon"
+version = "1.7.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1d2df5196e37bcc87abebc0053e20787d73847bb33134a69841207dd0a47f03b"
+dependencies = [
+ "either",
+ "rayon-core",
+]
+
+[[package]]
+name = "rayon-core"
+version = "1.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4b8f95bd6966f5c87776639160a66bd8ab9895d9d4ab01ddba9fc60661aebe8d"
+dependencies = [
+ "crossbeam-channel",
+ "crossbeam-deque",
+ "crossbeam-utils",
+ "num_cpus",
+]
+
+[[package]]
+name = "rdev"
+version = "0.5.0-2"
+source = "git+https://github.com/open-trade/rdev#0ad53987fa6f0e37a7bc000358f71c3802de4e7c"
+dependencies = [
+ "cocoa",
+ "core-foundation",
+ "core-foundation-sys",
+ "core-graphics",
+ "enum-map",
+ "lazy_static",
+ "libc",
+ "widestring 1.0.2",
+ "winapi 0.3.9",
+ "x11",
+]
+
+[[package]]
+name = "rdrand"
+version = "0.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "678054eb77286b51581ba43620cc911abf02758c91f93f479767aed0f90458b2"
+dependencies = [
+ "rand_core 0.3.1",
+]
+
+[[package]]
+name = "realfft"
+version = "3.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "93d6b8e8f0c6d2234aa58048d7290c60bf92cd36fd2888cd8331c66ad4f2e1d2"
+dependencies = [
+ "rustfft",
+]
+
+[[package]]
+name = "redox_syscall"
+version = "0.2.16"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "fb5a58c1855b4b6819d59012155603f0b22ad30cad752600aadfcb695265519a"
+dependencies = [
+ "bitflags",
+]
+
+[[package]]
+name = "redox_syscall"
+version = "0.3.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "567664f262709473930a4bf9e51bf2ebf3348f2e748ccc50dea20646858f8f29"
+dependencies = [
+ "bitflags",
+]
+
+[[package]]
+name = "redox_users"
+version = "0.4.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b033d837a7cf162d7993aded9304e30a83213c648b6e389db233191f891e5c2b"
+dependencies = [
+ "getrandom",
+ "redox_syscall 0.2.16",
+ "thiserror",
+]
+
+[[package]]
+name = "regex"
+version = "1.7.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8b1f693b24f6ac912f4893ef08244d70b6067480d2f1a46e950c9691e6749d1d"
+dependencies = [
+ "aho-corasick",
+ "memchr",
+ "regex-syntax",
+]
+
+[[package]]
+name = "regex-syntax"
+version = "0.6.29"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f162c6dd7b008981e4d40210aca20b4bd0f9b60ca9271061b07f78537722f2e1"
+
+[[package]]
+name = "repng"
+version = "0.2.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0dd57cd2cb5cc699b3eb4824d654e5a32f3bc013766da4966f71fe94805abbda"
+dependencies = [
+ "byteorder",
+ "flate2",
+]
+
+[[package]]
+name = "ring"
+version = "0.16.20"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3053cf52e236a3ed746dfc745aa9cacf1b791d846bdaf412f60a8d7d6e17c8fc"
+dependencies = [
+ "cc",
+ "libc",
+ "once_cell",
+ "spin",
+ "untrusted",
+ "web-sys",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "rpassword"
+version = "2.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d37473170aedbe66ffa3ad3726939ba677d83c646ad4fd99e5b4bc38712f45ec"
+dependencies = [
+ "kernel32-sys",
+ "libc",
+ "winapi 0.2.8",
+]
+
+[[package]]
+name = "rpassword"
+version = "6.0.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2bf099a1888612545b683d2661a1940089f6c2e5a8e38979b2159da876bfd956"
+dependencies = [
+ "libc",
+ "serde 1.0.159",
+ "serde_json 1.0.95",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "rubato"
+version = "0.12.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "cd70209c27d5b08f5528bdc779ea3ffb418954e28987f9f9775c6eac41003f9c"
+dependencies = [
+ "num-complex",
+ "num-integer",
+ "num-traits 0.2.15",
+ "realfft",
+]
+
+[[package]]
+name = "runas"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a620b0994a180cdfa25c0439e6d58c0628272571501880d626ffff58e96a0799"
+dependencies = [
+ "cc",
+ "which 3.1.1",
+]
+
+[[package]]
+name = "rust-pulsectl"
+version = "0.2.12"
+source = "git+https://github.com/open-trade/pulsectl#5e68f4c2b7c644fa321984688602d71e8ad0bba3"
+dependencies = [
+ "libpulse-binding",
+]
+
+[[package]]
+name = "rustc-demangle"
+version = "0.1.22"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d4a36c42d1873f9a77c53bde094f9664d9891bc604a45b4798fd2c389ed12e5b"
+
+[[package]]
+name = "rustc-hash"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2"
+
+[[package]]
+name = "rustc_version"
+version = "0.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bfa0f585226d2e68097d4f95d113b15b83a82e819ab25717ec0590d9584ef366"
+dependencies = [
+ "semver",
+]
+
+[[package]]
+name = "rustdesk"
+version = "1.1.9"
+dependencies = [
+ "android_logger",
+ "arboard",
+ "async-trait",
+ "base64 0.13.1",
+ "cc",
+ "cfg-if 1.0.0",
+ "clap 3.2.23",
+ "clipboard",
+ "clipboard-master",
+ "cocoa",
+ "core-foundation",
+ "core-graphics",
+ "cpal",
+ "ctrlc",
+ "dasp",
+ "dispatch",
+ "enigo",
+ "flexi_logger",
+ "hbb_common",
+ "hound",
+ "include_dir",
+ "lazy_static",
+ "libc",
+ "libpulse-binding",
+ "libpulse-simple-binding",
+ "mac_address",
+ "machine-uid",
+ "magnum-opus",
+ "num_cpus",
+ "objc",
+ "parity-tokio-ipc",
+ "rdev",
+ "repng",
+ "rpassword 6.0.1",
+ "rubato",
+ "runas",
+ "rust-pulsectl",
+ "samplerate",
+ "sciter-rs",
+ "scrap",
+ "serde 1.0.159",
+ "serde_derive",
+ "serde_json 1.0.95",
+ "sha2",
+ "sys-locale",
+ "sysinfo",
+ "systray",
+ "tray-item",
+ "uuid",
+ "virtual_display",
+ "whoami",
+ "winapi 0.3.9",
+ "windows-service",
+ "winreg 0.10.1",
+ "winres",
+]
+
+[[package]]
+name = "rustfft"
+version = "6.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e17d4f6cbdb180c9f4b2a26bbf01c4e647f1e1dea22fe8eb9db54198b32f9434"
+dependencies = [
+ "num-complex",
+ "num-integer",
+ "num-traits 0.2.15",
+ "primal-check",
+ "strength_reduce",
+ "transpose",
+ "version_check",
+]
+
+[[package]]
+name = "rustix"
+version = "0.37.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "85597d61f83914ddeba6a47b3b8ffe7365107221c2e557ed94426489fefb5f77"
+dependencies = [
+ "bitflags",
+ "errno",
+ "io-lifetimes",
+ "libc",
+ "linux-raw-sys",
+ "windows-sys 0.48.0",
+]
+
+[[package]]
+name = "rustls"
+version = "0.20.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "fff78fc74d175294f4e83b28343315ffcfb114b156f0185e9741cb5570f50e2f"
+dependencies = [
+ "ring",
+ "sct",
+ "webpki",
+]
+
+[[package]]
+name = "rustls-native-certs"
+version = "0.6.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0167bac7a9f490495f3c33013e7722b53cb087ecbe082fb0c6387c96f634ea50"
+dependencies = [
+ "openssl-probe",
+ "rustls-pemfile 1.0.2",
+ "schannel",
+ "security-framework",
+]
+
+[[package]]
+name = "rustls-pemfile"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5eebeaeb360c87bfb72e84abdb3447159c0eaececf1bef2aecd65a8be949d1c9"
+dependencies = [
+ "base64 0.13.1",
+]
+
+[[package]]
+name = "rustls-pemfile"
+version = "1.0.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d194b56d58803a43635bdc398cd17e383d6f71f9182b9a192c127ca42494a59b"
+dependencies = [
+ "base64 0.21.0",
+]
+
+[[package]]
+name = "rustversion"
+version = "1.0.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4f3208ce4d8448b3f3e7d168a73f5e0c43a61e32930de3bceeccedb388b6bf06"
+
+[[package]]
+name = "ryu"
+version = "1.0.13"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f91339c0467de62360649f8d3e185ca8de4224ff281f66000de5eb2a77a79041"
+
+[[package]]
+name = "same-file"
+version = "1.0.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "93fc1dc3aaa9bfed95e02e6eadabb4baf7e3078b0bd1b4d7b6b0b68378900502"
+dependencies = [
+ "winapi-util",
+]
+
+[[package]]
+name = "samplerate"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e032b2b24715c4f982f483ea3abdb3c9ba444d9f63e87b2843d6f998f5ba2698"
+dependencies = [
+ "libsamplerate-sys",
+]
+
+[[package]]
+name = "schannel"
+version = "0.1.21"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "713cfb06c7059f3588fb8044c0fad1d09e3c01d225e25b9220dbfdcf16dbb1b3"
+dependencies = [
+ "windows-sys 0.42.0",
+]
+
+[[package]]
+name = "sciter-rs"
+version = "0.5.57"
+source = "git+https://github.com/open-trade/rust-sciter?branch=dyn#82025b9ba77d5ae14543009444033036dbe25917"
+dependencies = [
+ "lazy_static",
+ "libc",
+ "objc",
+ "objc-foundation",
+]
+
+[[package]]
+name = "scopeguard"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d29ab0c6d3fc0ee92fe66e2d99f700eab17a8d57d1c1d3b748380fb20baa78cd"
+
+[[package]]
+name = "scrap"
+version = "0.5.0"
+dependencies = [
+ "bindgen 0.59.2",
+ "block",
+ "cfg-if 1.0.0",
+ "dbus",
+ "docopt",
+ "gstreamer",
+ "gstreamer-app",
+ "gstreamer-video",
+ "libc",
+ "num_cpus",
+ "quest",
+ "repng",
+ "serde 1.0.159",
+ "target_build_utils",
+ "tracing",
+ "vcpkg",
+ "webm",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "scratch"
+version = "1.0.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1792db035ce95be60c3f8853017b3999209281c24e2ba5bc8e59bf97a0c590c1"
+
+[[package]]
+name = "sct"
+version = "0.7.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d53dcdb7c9f8158937a7981b48accfd39a43af418591a5d008c7b22b5e1b7ca4"
+dependencies = [
+ "ring",
+ "untrusted",
+]
+
+[[package]]
+name = "security-framework"
+version = "2.8.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a332be01508d814fed64bf28f798a146d73792121129962fdf335bb3c49a4254"
+dependencies = [
+ "bitflags",
+ "core-foundation",
+ "core-foundation-sys",
+ "libc",
+ "security-framework-sys",
+]
+
+[[package]]
+name = "security-framework-sys"
+version = "2.8.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "31c9bb296072e961fcbd8853511dd39c2d8be2deb1e17c6860b1d30732b323b4"
+dependencies = [
+ "core-foundation-sys",
+ "libc",
+]
+
+[[package]]
+name = "semver"
+version = "1.0.17"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bebd363326d05ec3e2f532ab7660680f3b02130d780c299bca73469d521bc0ed"
+
+[[package]]
+name = "serde"
+version = "0.9.15"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "34b623917345a631dc9608d5194cc206b3fe6c3554cd1c75b937e55e285254af"
+
+[[package]]
+name = "serde"
+version = "1.0.159"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3c04e8343c3daeec41f58990b9d77068df31209f2af111e059e9fe9646693065"
+dependencies = [
+ "serde_derive",
+]
+
+[[package]]
+name = "serde_derive"
+version = "1.0.159"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4c614d17805b093df4b147b51339e7e44bf05ef59fba1e45d83500bcfb4d8585"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 2.0.13",
+]
+
+[[package]]
+name = "serde_json"
+version = "0.9.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ad8bcf487be7d2e15d3d543f04312de991d631cfe1b43ea0ade69e6a8a5b16a1"
+dependencies = [
+ "dtoa",
+ "itoa 0.3.4",
+ "num-traits 0.1.43",
+ "serde 0.9.15",
+]
+
+[[package]]
+name = "serde_json"
+version = "1.0.95"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d721eca97ac802aa7777b701877c8004d950fc142651367300d21c1cc0194744"
+dependencies = [
+ "itoa 1.0.6",
+ "ryu",
+ "serde 1.0.159",
+]
+
+[[package]]
+name = "serde_spanned"
+version = "0.6.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0efd8caf556a6cebd3b285caf480045fcc1ac04f6bd786b09a6f11af30c4fcf4"
+dependencies = [
+ "serde 1.0.159",
+]
+
+[[package]]
+name = "sha2"
+version = "0.10.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "82e6b795fe2e3b1e845bafcb27aa35405c4d47cdfc92af5fc8d3002f76cebdc0"
+dependencies = [
+ "cfg-if 1.0.0",
+ "cpufeatures",
+ "digest",
+]
+
+[[package]]
+name = "shlex"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "43b2853a4d09f215c24cc5489c992ce46052d359b5109343cbafbf26bc62f8a3"
+
+[[package]]
+name = "signal-hook-registry"
+version = "1.4.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d8229b473baa5980ac72ef434c4415e70c4b5e71b423043adb4ba059f89c99a1"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "signature"
+version = "1.6.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "74233d3b3b2f6d4b006dc19dee745e73e2a6bfb6f93607cd3b02bd5b00797d7c"
+
+[[package]]
+name = "siphasher"
+version = "0.2.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0b8de496cf83d4ed58b6be86c3a275b8602f6ffe98d3024a869e124147a9a3ac"
+
+[[package]]
+name = "slab"
+version = "0.4.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6528351c9bc8ab22353f9d776db39a20288e8d6c37ef8cfe3317cf875eecfc2d"
+dependencies = [
+ "autocfg 1.1.0",
+]
+
+[[package]]
+name = "smallvec"
+version = "1.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a507befe795404456341dfab10cef66ead4c041f62b8b11bbb92bffe5d0953e0"
+
+[[package]]
+name = "socket2"
+version = "0.3.19"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "122e570113d28d773067fab24266b66753f6ea915758651696b6e35e49f88d6e"
+dependencies = [
+ "cfg-if 1.0.0",
+ "libc",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "socket2"
+version = "0.4.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "64a4a911eed85daf18834cfaa86a79b7d266ff93ff5ba14005426219480ed662"
+dependencies = [
+ "libc",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "sodiumoxide"
+version = "0.2.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e26be3acb6c2d9a7aac28482586a7856436af4cfe7100031d219de2d2ecb0028"
+dependencies = [
+ "ed25519",
+ "libc",
+ "libsodium-sys",
+ "serde 1.0.159",
+]
+
+[[package]]
+name = "spin"
+version = "0.5.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6e63cff320ae2c57904679ba7cb63280a3dc4613885beafb148ee7bf9aa9042d"
+
+[[package]]
+name = "static_assertions"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f"
+
+[[package]]
+name = "stdweb"
+version = "0.1.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ef5430c8e36b713e13b48a9f709cc21e046723fe44ce34587b73a830203b533e"
+
+[[package]]
+name = "str-buf"
+version = "1.0.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9e08d8363704e6c71fc928674353e6b7c23dcea9d82d7012c8faf2a3a025f8d0"
+
+[[package]]
+name = "strength_reduce"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "fe895eb47f22e2ddd4dabc02bce419d2e643c8e3b585c78158b349195bc24d82"
+
+[[package]]
+name = "strsim"
+version = "0.8.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8ea5119cdb4c55b55d432abb513a0429384878c15dde60cc77b1c99de1a95a6a"
+
+[[package]]
+name = "strsim"
+version = "0.10.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "73473c0e59e6d5812c5dfe2a064a6444949f089e20eec9a2e5506596494e4623"
+
+[[package]]
+name = "strum"
+version = "0.18.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "57bd81eb48f4c437cadc685403cad539345bf703d78e63707418431cecd4522b"
+
+[[package]]
+name = "strum_macros"
+version = "0.18.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "87c85aa3f8ea653bfd3ddf25f7ee357ee4d204731f6aa9ad04002306f6e2774c"
+dependencies = [
+ "heck 0.3.3",
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "syn"
+version = "1.0.109"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "72b64191b275b66ffe2469e8af2c1cfe3bafa67b529ead792a6d0160888b4237"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "unicode-ident",
+]
+
+[[package]]
+name = "syn"
+version = "2.0.13"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4c9da457c5285ac1f936ebd076af6dac17a61cfe7826f2076b4d015cf47bc8ec"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "unicode-ident",
+]
+
+[[package]]
+name = "synstructure"
+version = "0.12.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f36bdaa60a83aca3921b5259d5400cbf5e90fc51931376a9bd4a0eb79aa7210f"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+ "unicode-xid",
+]
+
+[[package]]
+name = "sys-locale"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f8a11bd9c338fdba09f7881ab41551932ad42e405f61d01e8406baea71c07aee"
+dependencies = [
+ "js-sys",
+ "libc",
+ "wasm-bindgen",
+ "web-sys",
+ "windows-sys 0.45.0",
+]
+
+[[package]]
+name = "sysinfo"
+version = "0.23.13"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3977ec2e0520829be45c8a2df70db2bf364714d8a748316a10c3c35d4d2b01c9"
+dependencies = [
+ "cfg-if 1.0.0",
+ "core-foundation-sys",
+ "libc",
+ "ntapi",
+ "once_cell",
+ "rayon",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "system-deps"
+version = "1.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0f3ecc17269a19353b3558b313bba738b25d82993e30d62a18406a24aba4649b"
+dependencies = [
+ "heck 0.3.3",
+ "pkg-config",
+ "strum",
+ "strum_macros",
+ "thiserror",
+ "toml 0.5.11",
+ "version-compare 0.0.10",
+]
+
+[[package]]
+name = "system-deps"
+version = "6.0.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "555fc8147af6256f3931a36bb83ad0023240ce9cf2b319dec8236fd1f220b05f"
+dependencies = [
+ "cfg-expr",
+ "heck 0.4.1",
+ "pkg-config",
+ "toml 0.7.3",
+ "version-compare 0.1.1",
+]
+
+[[package]]
+name = "systray"
+version = "0.4.1"
+source = "git+https://github.com/liyue201/systray-rs#84cca4b4171661bc6c4d1ba5aaa2320ff8e085aa"
+dependencies = [
+ "glib 0.10.3",
+ "gtk 0.9.2",
+ "libappindicator 0.6.1",
+ "libc",
+ "log",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "target_build_utils"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "013d134ae4a25ee744ad6129db589018558f620ddfa44043887cdd45fa08e75c"
+dependencies = [
+ "phf",
+ "phf_codegen",
+ "serde_json 0.9.10",
+]
+
+[[package]]
+name = "tempfile"
+version = "3.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b9fbec84f381d5795b08656e4912bec604d162bff9291d6189a78f4c8ab87998"
+dependencies = [
+ "cfg-if 1.0.0",
+ "fastrand",
+ "redox_syscall 0.3.5",
+ "rustix",
+ "windows-sys 0.45.0",
+]
+
+[[package]]
+name = "termcolor"
+version = "1.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "be55cf8942feac5c765c2c993422806843c9a9a45d4d5c407ad6dd2ea95eb9b6"
+dependencies = [
+ "winapi-util",
+]
+
+[[package]]
+name = "termios"
+version = "0.3.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "411c5bf740737c7918b8b1fe232dca4dc9f8e754b8ad5e20966814001ed0ac6b"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "textwrap"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d326610f408c7a4eb6f51c37c330e496b08506c9457c9d34287ecc38809fb060"
+dependencies = [
+ "unicode-width",
+]
+
+[[package]]
+name = "textwrap"
+version = "0.16.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "222a222a5bfe1bba4a77b45ec488a741b3cb8872e5e499451fd7d0129c9c7c3d"
+
+[[package]]
+name = "thiserror"
+version = "1.0.40"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "978c9a314bd8dc99be594bc3c175faaa9794be04a5a5e153caba6915336cebac"
+dependencies = [
+ "thiserror-impl",
+]
+
+[[package]]
+name = "thiserror-impl"
+version = "1.0.40"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f9456a42c5b0d803c8cd86e73dd7cc9edd429499f37a3550d286d5e86720569f"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 2.0.13",
+]
+
+[[package]]
+name = "tiff"
+version = "0.6.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9a53f4706d65497df0c4349241deddf35f84cee19c87ed86ea8ca590f4464437"
+dependencies = [
+ "jpeg-decoder",
+ "miniz_oxide 0.4.4",
+ "weezl",
+]
+
+[[package]]
+name = "time"
+version = "0.3.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c2702e08a7a860f005826c6815dcac101b19b5eb330c27fe4a5928fec1d20ddd"
+dependencies = [
+ "itoa 1.0.6",
+ "libc",
+ "num_threads",
+ "time-macros",
+]
+
+[[package]]
+name = "time-macros"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "42657b1a6f4d817cda8e7a0ace261fe0cc946cf3a80314390b22cc61ae080792"
+
+[[package]]
+name = "tinyvec"
+version = "1.6.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "87cc5ceb3875bb20c2890005a4e226a4651264a5c75edb2421b52861a0a0cb50"
+dependencies = [
+ "tinyvec_macros",
+]
+
+[[package]]
+name = "tinyvec_macros"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1f3ccbac311fea05f86f61904b462b55fb3df8837a366dfc601a0161d0532f20"
+
+[[package]]
+name = "tokio"
+version = "1.27.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d0de47a4eecbe11f498978a9b29d792f0d2692d1dd003650c24c76510e3bc001"
+dependencies = [
+ "autocfg 1.1.0",
+ "bytes",
+ "libc",
+ "mio 0.8.6",
+ "num_cpus",
+ "parking_lot 0.12.1",
+ "pin-project-lite",
+ "signal-hook-registry",
+ "socket2 0.4.9",
+ "tokio-macros",
+ "windows-sys 0.45.0",
+]
+
+[[package]]
+name = "tokio-macros"
+version = "2.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "61a573bdc87985e9d6ddeed1b3d864e8a302c847e40d647746df2f1de209d1ce"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 2.0.13",
+]
+
+[[package]]
+name = "tokio-socks"
+version = "0.5.1-1"
+source = "git+https://github.com/open-trade/tokio-socks#7034e79263ce25c348be072808d7601d82cd892d"
+dependencies = [
+ "bytes",
+ "either",
+ "futures-core",
+ "futures-sink",
+ "futures-util",
+ "pin-project",
+ "thiserror",
+ "tokio",
+ "tokio-util 0.7.7",
+]
+
+[[package]]
+name = "tokio-util"
+version = "0.6.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "36943ee01a6d67977dd3f84a5a1d2efeb4ada3a1ae771cadfaa535d9d9fc6507"
+dependencies = [
+ "bytes",
+ "futures-core",
+ "futures-io",
+ "futures-sink",
+ "log",
+ "pin-project-lite",
+ "slab",
+ "tokio",
+]
+
+[[package]]
+name = "tokio-util"
+version = "0.7.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5427d89453009325de0d8f342c9490009f76e999cb7672d77e46267448f7e6b2"
+dependencies = [
+ "bytes",
+ "futures-core",
+ "futures-sink",
+ "pin-project-lite",
+ "tokio",
+ "tracing",
+]
+
+[[package]]
+name = "toml"
+version = "0.5.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f4f7f0dd8d50a853a531c426359045b1998f04219d88799810762cd4ad314234"
+dependencies = [
+ "serde 1.0.159",
+]
+
+[[package]]
+name = "toml"
+version = "0.7.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b403acf6f2bb0859c93c7f0d967cb4a75a7ac552100f9322faf64dc047669b21"
+dependencies = [
+ "serde 1.0.159",
+ "serde_spanned",
+ "toml_datetime",
+ "toml_edit",
+]
+
+[[package]]
+name = "toml_datetime"
+version = "0.6.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3ab8ed2edee10b50132aed5f331333428b011c99402b5a534154ed15746f9622"
+dependencies = [
+ "serde 1.0.159",
+]
+
+[[package]]
+name = "toml_edit"
+version = "0.19.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "239410c8609e8125456927e6707163a3b1fdb40561e4b803bc041f466ccfdc13"
+dependencies = [
+ "indexmap",
+ "serde 1.0.159",
+ "serde_spanned",
+ "toml_datetime",
+ "winnow",
+]
+
+[[package]]
+name = "tracing"
+version = "0.1.37"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8ce8c33a8d48bd45d624a6e523445fd21ec13d3653cd51f681abf67418f54eb8"
+dependencies = [
+ "cfg-if 1.0.0",
+ "pin-project-lite",
+ "tracing-attributes",
+ "tracing-core",
+]
+
+[[package]]
+name = "tracing-attributes"
+version = "0.1.23"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4017f8f45139870ca7e672686113917c71c7a6e02d4924eda67186083c03081a"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+]
+
+[[package]]
+name = "tracing-core"
+version = "0.1.30"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "24eb03ba0eab1fd845050058ce5e616558e8f8d8fca633e6b163fe25c797213a"
+dependencies = [
+ "once_cell",
+]
+
+[[package]]
+name = "transpose"
+version = "0.2.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e6522d49d03727ffb138ae4cbc1283d3774f0d10aa7f9bf52e6784c45daf9b23"
+dependencies = [
+ "num-integer",
+ "strength_reduce",
+]
+
+[[package]]
+name = "tray-item"
+version = "0.7.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0914b62e00e8f51241806cb9f9c4ea6b10c75d94cae02c89278de6f4b98c7d0f"
+dependencies = [
+ "cocoa",
+ "core-graphics",
+ "gtk 0.15.5",
+ "libappindicator 0.7.1",
+ "libc",
+ "objc",
+ "objc-foundation",
+ "objc_id",
+ "padlock",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "typenum"
+version = "1.16.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "497961ef93d974e23eb6f433eb5fe1b7930b659f06d12dec6fc44a8f554c0bba"
+
+[[package]]
+name = "unicode-ident"
+version = "1.0.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e5464a87b239f13a63a501f2701565754bae92d243d4bb7eb12f6d57d2269bf4"
+
+[[package]]
+name = "unicode-segmentation"
+version = "1.10.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1dd624098567895118886609431a7c3b8f516e41d30e0643f03d94592a147e36"
+
+[[package]]
+name = "unicode-width"
+version = "0.1.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c0edd1e5b14653f783770bce4a4dabb4a5108a5370a5f5d8cfe8710c361f6c8b"
+
+[[package]]
+name = "unicode-xid"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f962df74c8c05a667b5ee8bcf162993134c104e96440b663c8daa176dc772d8c"
+
+[[package]]
+name = "untrusted"
+version = "0.7.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a156c684c91ea7d62626509bce3cb4e1d9ed5c4d978f7b4352658f96a4c26b4a"
+
+[[package]]
+name = "uuid"
+version = "1.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1674845326ee10d37ca60470760d4288a6f80f304007d92e5c53bab78c9cfd79"
+dependencies = [
+ "getrandom",
+]
+
+[[package]]
+name = "vcpkg"
+version = "0.2.15"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "accd4ea62f7bb7a82fe23066fb0957d48ef677f6eeb8215f372f52e48bb32426"
+
+[[package]]
+name = "vec_map"
+version = "0.8.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f1bddf1187be692e79c5ffeab891132dfb0f236ed36a43c7ed39f1165ee20191"
+
+[[package]]
+name = "version-compare"
+version = "0.0.10"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d63556a25bae6ea31b52e640d7c41d1ab27faba4ccb600013837a3d0b3994ca1"
+
+[[package]]
+name = "version-compare"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "579a42fc0b8e0c63b76519a339be31bed574929511fa53c1a3acae26eb258f29"
+
+[[package]]
+name = "version_check"
+version = "0.9.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "49874b5167b65d7193b8aba1567f5c7d93d001cafc34600cee003eda787e483f"
+
+[[package]]
+name = "virtual_display"
+version = "0.1.0"
+dependencies = [
+ "cc",
+ "hbb_common",
+ "lazy_static",
+ "serde 1.0.159",
+ "serde_derive",
+ "thiserror",
+]
+
+[[package]]
+name = "walkdir"
+version = "2.3.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "36df944cda56c7d8d8b7496af378e6b16de9284591917d307c9b4d313c44e698"
+dependencies = [
+ "same-file",
+ "winapi-util",
+]
+
+[[package]]
+name = "wasi"
+version = "0.11.0+wasi-snapshot-preview1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9c8d87e72b64a3b4db28d11ce29237c246188f4f51057d65a7eab63b7987e423"
+
+[[package]]
+name = "wasm-bindgen"
+version = "0.2.84"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "31f8dcbc21f30d9b8f2ea926ecb58f6b91192c17e9d33594b3df58b2007ca53b"
+dependencies = [
+ "cfg-if 1.0.0",
+ "wasm-bindgen-macro",
+]
+
+[[package]]
+name = "wasm-bindgen-backend"
+version = "0.2.84"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "95ce90fd5bcc06af55a641a86428ee4229e44e07033963a2290a8e241607ccb9"
+dependencies = [
+ "bumpalo",
+ "log",
+ "once_cell",
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+ "wasm-bindgen-shared",
+]
+
+[[package]]
+name = "wasm-bindgen-macro"
+version = "0.2.84"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4c21f77c0bedc37fd5dc21f897894a5ca01e7bb159884559461862ae90c0b4c5"
+dependencies = [
+ "quote",
+ "wasm-bindgen-macro-support",
+]
+
+[[package]]
+name = "wasm-bindgen-macro-support"
+version = "0.2.84"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2aff81306fcac3c7515ad4e177f521b5c9a15f2b08f4e32d823066102f35a5f6"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 1.0.109",
+ "wasm-bindgen-backend",
+ "wasm-bindgen-shared",
+]
+
+[[package]]
+name = "wasm-bindgen-shared"
+version = "0.2.84"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0046fef7e28c3804e5e38bfa31ea2a0f73905319b677e57ebe37e49358989b5d"
+
+[[package]]
+name = "web-sys"
+version = "0.3.61"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e33b99f4b23ba3eec1a53ac264e35a755f00e966e0065077d6027c0f575b0b97"
+dependencies = [
+ "js-sys",
+ "wasm-bindgen",
+]
+
+[[package]]
+name = "webm"
+version = "1.0.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ecb047148a12ef1fd8ab26302bca7e82036f005c3073b48e17cc1b44ec577136"
+dependencies = [
+ "webm-sys",
+]
+
+[[package]]
+name = "webm-sys"
+version = "1.0.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0ded6ec82ccf51fe265b0b2b1579cac839574ed910c17baac58e807f8a9de7f3"
+dependencies = [
+ "cc",
+]
+
+[[package]]
+name = "webpki"
+version = "0.22.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f095d78192e208183081cc07bc5515ef55216397af48b873e5edcd72637fa1bd"
+dependencies = [
+ "ring",
+ "untrusted",
+]
+
+[[package]]
+name = "weezl"
+version = "0.1.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9193164d4de03a926d909d3bc7c30543cecb35400c02114792c2cae20d5e2dbb"
+
+[[package]]
+name = "which"
+version = "3.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d011071ae14a2f6671d0b74080ae0cd8ebf3a6f8c9589a2cd45f23126fe29724"
+dependencies = [
+ "failure",
+ "libc",
+]
+
+[[package]]
+name = "which"
+version = "4.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2441c784c52b289a054b7201fc93253e288f094e2f4be9058343127c4226a269"
+dependencies = [
+ "either",
+ "libc",
+ "once_cell",
+]
+
+[[package]]
+name = "whoami"
+version = "1.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2c70234412ca409cc04e864e89523cb0fc37f5e1344ebed5a3ebf4192b6b9f68"
+dependencies = [
+ "wasm-bindgen",
+ "web-sys",
+]
+
+[[package]]
+name = "widestring"
+version = "0.4.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c168940144dd21fd8046987c16a46a33d5fc84eec29ef9dcddc2ac9e31526b7c"
+
+[[package]]
+name = "widestring"
+version = "1.0.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "653f141f39ec16bba3c5abe400a0c60da7468261cc2cbf36805022876bc721a8"
+
+[[package]]
+name = "winapi"
+version = "0.2.8"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "167dc9d6949a9b857f3451275e911c3f44255842c1f7a76f33c55103a909087a"
+
+[[package]]
+name = "winapi"
+version = "0.3.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5c839a674fcd7a98952e593242ea400abe93992746761e38641405d28b00f419"
+dependencies = [
+ "winapi-i686-pc-windows-gnu",
+ "winapi-x86_64-pc-windows-gnu",
+]
+
+[[package]]
+name = "winapi-build"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2d315eee3b34aca4797b2da6b13ed88266e6d612562a0c46390af8299fc699bc"
+
+[[package]]
+name = "winapi-i686-pc-windows-gnu"
+version = "0.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ac3b87c63620426dd9b991e5ce0329eff545bccbbb34f3be09ff6fb6ab51b7b6"
+
+[[package]]
+name = "winapi-util"
+version = "0.1.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "70ec6ce85bb158151cae5e5c87f95a8e97d2c0c4b001223f33a334e3ce5de178"
+dependencies = [
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "winapi-wsapoll"
+version = "0.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "44c17110f57155602a80dca10be03852116403c9ff3cd25b079d666f2aa3df6e"
+dependencies = [
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "winapi-x86_64-pc-windows-gnu"
+version = "0.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "712e227841d057c1ee1cd2fb22fa7e5a5461ae8e48fa2ca79ec42cfc1931183f"
+
+[[package]]
+name = "windows"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e686886bc078bc1b0b600cac0147aadb815089b6e4da64016cbd754b6342700f"
+dependencies = [
+ "windows-targets 0.48.0",
+]
+
+[[package]]
+name = "windows-service"
+version = "0.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0c643e10139d127d30d6d753398c8a6f0a43532e8370f6c9d29ebbff29b984ab"
+dependencies = [
+ "bitflags",
+ "err-derive",
+ "widestring 0.4.3",
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "windows-sys"
+version = "0.28.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "82ca39602d5cbfa692c4b67e3bcbb2751477355141c1ed434c94da4186836ff6"
+dependencies = [
+ "windows_aarch64_msvc 0.28.0",
+ "windows_i686_gnu 0.28.0",
+ "windows_i686_msvc 0.28.0",
+ "windows_x86_64_gnu 0.28.0",
+ "windows_x86_64_msvc 0.28.0",
+]
+
+[[package]]
+name = "windows-sys"
+version = "0.42.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5a3e1820f08b8513f676f7ab6c1f99ff312fb97b553d30ff4dd86f9f15728aa7"
+dependencies = [
+ "windows_aarch64_gnullvm 0.42.2",
+ "windows_aarch64_msvc 0.42.2",
+ "windows_i686_gnu 0.42.2",
+ "windows_i686_msvc 0.42.2",
+ "windows_x86_64_gnu 0.42.2",
+ "windows_x86_64_gnullvm 0.42.2",
+ "windows_x86_64_msvc 0.42.2",
+]
+
+[[package]]
+name = "windows-sys"
+version = "0.45.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "75283be5efb2831d37ea142365f009c02ec203cd29a3ebecbc093d52315b66d0"
+dependencies = [
+ "windows-targets 0.42.2",
+]
+
+[[package]]
+name = "windows-sys"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "677d2418bec65e3338edb076e806bc1ec15693c5d0104683f2efe857f61056a9"
+dependencies = [
+ "windows-targets 0.48.0",
+]
+
+[[package]]
+name = "windows-targets"
+version = "0.42.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8e5180c00cd44c9b1c88adb3693291f1cd93605ded80c250a75d472756b4d071"
+dependencies = [
+ "windows_aarch64_gnullvm 0.42.2",
+ "windows_aarch64_msvc 0.42.2",
+ "windows_i686_gnu 0.42.2",
+ "windows_i686_msvc 0.42.2",
+ "windows_x86_64_gnu 0.42.2",
+ "windows_x86_64_gnullvm 0.42.2",
+ "windows_x86_64_msvc 0.42.2",
+]
+
+[[package]]
+name = "windows-targets"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7b1eb6f0cd7c80c79759c929114ef071b87354ce476d9d94271031c0497adfd5"
+dependencies = [
+ "windows_aarch64_gnullvm 0.48.0",
+ "windows_aarch64_msvc 0.48.0",
+ "windows_i686_gnu 0.48.0",
+ "windows_i686_msvc 0.48.0",
+ "windows_x86_64_gnu 0.48.0",
+ "windows_x86_64_gnullvm 0.48.0",
+ "windows_x86_64_msvc 0.48.0",
+]
+
+[[package]]
+name = "windows-win"
+version = "2.4.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8d4243ec23afe4e9b4e668b3c0a0e973f1b8265f6a46223cfcbc16fd267480c0"
+dependencies = [
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "windows_aarch64_gnullvm"
+version = "0.42.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "597a5118570b68bc08d8d59125332c54f1ba9d9adeedeef5b99b02ba2b0698f8"
+
+[[package]]
+name = "windows_aarch64_gnullvm"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "91ae572e1b79dba883e0d315474df7305d12f569b400fcf90581b06062f7e1bc"
+
+[[package]]
+name = "windows_aarch64_msvc"
+version = "0.28.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "52695a41e536859d5308cc613b4a022261a274390b25bd29dfff4bf08505f3c2"
+
+[[package]]
+name = "windows_aarch64_msvc"
+version = "0.42.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e08e8864a60f06ef0d0ff4ba04124db8b0fb3be5776a5cd47641e942e58c4d43"
+
+[[package]]
+name = "windows_aarch64_msvc"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b2ef27e0d7bdfcfc7b868b317c1d32c641a6fe4629c171b8928c7b08d98d7cf3"
+
+[[package]]
+name = "windows_i686_gnu"
+version = "0.28.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f54725ac23affef038fecb177de6c9bf065787c2f432f79e3c373da92f3e1d8a"
+
+[[package]]
+name = "windows_i686_gnu"
+version = "0.42.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c61d927d8da41da96a81f029489353e68739737d3beca43145c8afec9a31a84f"
+
+[[package]]
+name = "windows_i686_gnu"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "622a1962a7db830d6fd0a69683c80a18fda201879f0f447f065a3b7467daa241"
+
+[[package]]
+name = "windows_i686_msvc"
+version = "0.28.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "51d5158a43cc43623c0729d1ad6647e62fa384a3d135fd15108d37c683461f64"
+
+[[package]]
+name = "windows_i686_msvc"
+version = "0.42.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "44d840b6ec649f480a41c8d80f9c65108b92d89345dd94027bfe06ac444d1060"
+
+[[package]]
+name = "windows_i686_msvc"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4542c6e364ce21bf45d69fdd2a8e455fa38d316158cfd43b3ac1c5b1b19f8e00"
+
+[[package]]
+name = "windows_x86_64_gnu"
+version = "0.28.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bc31f409f565611535130cfe7ee8e6655d3fa99c1c61013981e491921b5ce954"
+
+[[package]]
+name = "windows_x86_64_gnu"
+version = "0.42.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8de912b8b8feb55c064867cf047dda097f92d51efad5b491dfb98f6bbb70cb36"
+
+[[package]]
+name = "windows_x86_64_gnu"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ca2b8a661f7628cbd23440e50b05d705db3686f894fc9580820623656af974b1"
+
+[[package]]
+name = "windows_x86_64_gnullvm"
+version = "0.42.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "26d41b46a36d453748aedef1486d5c7a85db22e56aff34643984ea85514e94a3"
+
+[[package]]
+name = "windows_x86_64_gnullvm"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7896dbc1f41e08872e9d5e8f8baa8fdd2677f29468c4e156210174edc7f7b953"
+
+[[package]]
+name = "windows_x86_64_msvc"
+version = "0.28.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3f2b8c7cbd3bfdddd9ab98769f9746a7fad1bca236554cd032b78d768bc0e89f"
+
+[[package]]
+name = "windows_x86_64_msvc"
+version = "0.42.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9aec5da331524158c6d1a4ac0ab1541149c0b9505fde06423b02f5ef0106b9f0"
+
+[[package]]
+name = "windows_x86_64_msvc"
+version = "0.48.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1a515f5799fe4961cb532f983ce2b23082366b898e52ffbce459c86f67c8378a"
+
+[[package]]
+name = "winnow"
+version = "0.4.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ae8970b36c66498d8ff1d66685dc86b91b29db0c7739899012f63a63814b4b28"
+dependencies = [
+ "memchr",
+]
+
+[[package]]
+name = "winreg"
+version = "0.6.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b2986deb581c4fe11b621998a5e53361efe6b48a151178d0cd9eeffa4dc6acc9"
+dependencies = [
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "winreg"
+version = "0.10.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "80d0f4e272c85def139476380b12f9ac60926689dd2e01d4923222f40580869d"
+dependencies = [
+ "winapi 0.3.9",
+]
+
+[[package]]
+name = "winres"
+version = "0.1.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b68db261ef59e9e52806f688020631e987592bd83619edccda9c47d42cde4f6c"
+dependencies = [
+ "toml 0.5.11",
+]
+
+[[package]]
+name = "ws2_32-sys"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d59cefebd0c892fa2dd6de581e937301d8552cb44489cdff035c6187cb63fa5e"
+dependencies = [
+ "winapi 0.2.8",
+ "winapi-build",
+]
+
+[[package]]
+name = "x11"
+version = "2.21.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "502da5464ccd04011667b11c435cb992822c2c0dbde1770c988480d312a0db2e"
+dependencies = [
+ "libc",
+ "pkg-config",
+]
+
+[[package]]
+name = "x11-clipboard"
+version = "0.5.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "473068b7b80ac86a18328824f1054e5e007898c47b5bbc281bd7abe32bc3653c"
+dependencies = [
+ "xcb",
+]
+
+[[package]]
+name = "x11rb"
+version = "0.9.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6e99be55648b3ae2a52342f9a870c0e138709a3493261ce9b469afe6e4df6d8a"
+dependencies = [
+ "gethostname",
+ "nix 0.22.3",
+ "winapi 0.3.9",
+ "winapi-wsapoll",
+]
+
+[[package]]
+name = "xcb"
+version = "0.10.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "771e2b996df720cd1c6dd9ff90f62d91698fd3610cc078388d0564bdd6622a9c"
+dependencies = [
+ "libc",
+ "log",
+ "quick-xml",
+]
+
+[[package]]
+name = "zstd"
+version = "0.9.2+zstd.1.5.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2390ea1bf6c038c39674f22d95f0564725fc06034a47129179810b2fc58caa54"
+dependencies = [
+ "zstd-safe",
+]
+
+[[package]]
+name = "zstd-safe"
+version = "4.1.3+zstd.1.5.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e99d81b99fb3c2c2c794e3fe56c305c63d5173a16a46b5850b07c935ffc7db79"
+dependencies = [
+ "libc",
+ "zstd-sys",
+]
+
+[[package]]
+name = "zstd-sys"
+version = "1.6.2+zstd.1.5.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2daf2f248d9ea44454bfcb2516534e8b8ad2fc91bf818a1885495fc42bc8ac9f"
+dependencies = [
+ "cc",
+ "libc",
+]
diff --git a/pkgs/applications/networking/remote/rustdesk/cargo.patch b/pkgs/applications/networking/remote/rustdesk/cargo.patch
index e2124f26171..eb8495f14ca 100644
--- a/pkgs/applications/networking/remote/rustdesk/cargo.patch
+++ b/pkgs/applications/networking/remote/rustdesk/cargo.patch
@@ -1,38 +1,3 @@
-diff --git a/Cargo.lock b/Cargo.lock
-index fb17c7e..ef157e5 100644
---- a/Cargo.lock
-+++ b/Cargo.lock
-@@ -2332,10 +2332,10 @@ dependencies = [
- [[package]]
- name = "magnum-opus"
- version = "0.4.0"
--source = "git+https://github.com/open-trade/magnum-opus#3c3d0b86ae95c84930bebffe4bcb03b3bd83342b"
-+source = "git+https://github.com/TheRadioGuy/magnum-opus#171e1d021004626f7444d1e39b98f50bc3cb2604"
- dependencies = [
-- "bindgen",
-- "target_build_utils",
-+ "libc",
-+ "opusic-sys",
- ]
- 
- [[package]]
-@@ -2796,6 +2796,16 @@ version = "0.1.5"
- source = "registry+https://github.com/rust-lang/crates.io-index"
- checksum = "ff011a302c396a5197692431fc1948019154afc178baf7d8e37367442a4601cf"
- 
-+[[package]]
-+name = "opusic-sys"
-+version = "0.3.6"
-+source = "registry+https://github.com/rust-lang/crates.io-index"
-+checksum = "5eace752ce07a037241dba8f02c654799f051e431b27028056bcb480e83b54f5"
-+dependencies = [
-+ "cmake",
-+ "libc",
-+]
-+
- [[package]]
- name = "os_str_bytes"
- version = "6.0.0"
 diff --git a/Cargo.toml b/Cargo.toml
 index 1b715bd..960e8da 100644
 --- a/Cargo.toml
diff --git a/pkgs/applications/networking/remote/rustdesk/default.nix b/pkgs/applications/networking/remote/rustdesk/default.nix
index fbbef2dc619..bb3b8ffed61 100644
--- a/pkgs/applications/networking/remote/rustdesk/default.nix
+++ b/pkgs/applications/networking/remote/rustdesk/default.nix
@@ -1,5 +1,4 @@
 { lib
-, stdenv
 , fetchFromGitHub
 , makeDesktopItem
 , copyDesktopItems
@@ -8,7 +7,6 @@
 , cmake
 , yasm
 , nasm
-, zip
 , pkg-config
 , clang
 , gtk3
@@ -43,9 +41,22 @@ rustPlatform.buildRustPackage rec {
     ./fix-for-rust-1.65.diff
   ];
 
-  cargoSha256 = "sha256-1OMWEk+DerltF7kwdo4d04rbgIFLHBRq3vZaL7jtrdE=";
+  LIBCLANG_PATH = "${llvmPackages.libclang.lib}/lib";
 
-  LIBCLANG_PATH="${llvmPackages.libclang.lib}/lib";
+  cargoLock = {
+    lockFile = ./Cargo.lock;
+    outputHashes = {
+      "confy-0.4.0" = "sha256-e91cvEixhpPzIthAxzTa3fDY6eCsHUy/eZQAqs7QTDo=";
+      "parity-tokio-ipc-0.7.3-1" = "sha256-eULJePtBu0iBI3It/bPH0h82Obsb1PJALgwYwrnCFYI=";
+      "rdev-0.5.0-2" = "sha256-7CEZ2wIM4QAPfY1tGKqXfHplTaxHnccVqFRPjY21Svo=";
+      "tokio-socks-0.5.1-1" = "sha256-45QQ6FrhGU9uEhbKXTKd/mY6MDumO6p46NmlakdyDQk=";
+      "libappindicator-0.6.1" = "sha256-JGnnZrcwbh8WJ6+/4bYhfD3HvgF2C7XaaGb6TaMRWdw=";
+      "magnum-opus-0.4.0" = "sha256-U5uuN4YolOYDnFNbtPpwYefcBDTUUyioui0UCcW8dyo=";
+      "rust-pulsectl-0.2.12" = "sha256-8jXTspWvjONFcvw9/Z8C43g4BuGZ3rsG32tvLMQbtbM=";
+      "sciter-rs-0.5.57" = "sha256-ZZnZDhMjK0LjgmK0da1yvB0uoKueLhhhQtzmjoN+1R0=";
+      "systray-0.4.1" = "sha256-p1PMr/8oS6zHx4+Ng4zCqt0xZ57cq3wAu6/agyWq5Jw=";
+    };
+  };
 
   # Change magnus-opus version to upstream so that it does not use
   # vcpkg for libopus since it does not work.
@@ -55,35 +66,37 @@ rustPlatform.buildRustPackage rec {
 
   # Manually simulate a vcpkg installation so that it can link the libaries
   # properly.
-  postUnpack = let
-    vcpkg_target = "x64-linux";
+  postUnpack =
+    let
+      vcpkg_target = "x64-linux";
 
-    updates_vcpkg_file = writeText "update_vcpkg_rustdesk"
+      updates_vcpkg_file = writeText "update_vcpkg_rustdesk"
+        ''
+          Package : libyuv
+          Architecture : ${vcpkg_target}
+          Version : 1.0
+          Status : is installed
+
+          Package : libvpx
+          Architecture : ${vcpkg_target}
+          Version : 1.0
+          Status : is installed
+        '';
+    in
     ''
-    Package : libyuv
-    Architecture : ${vcpkg_target}
-    Version : 1.0
-    Status : is installed
+      export VCPKG_ROOT="$TMP/vcpkg";
 
-    Package : libvpx
-    Architecture : ${vcpkg_target}
-    Version : 1.0
-    Status : is installed
+      mkdir -p $VCPKG_ROOT/.vcpkg-root
+      mkdir -p $VCPKG_ROOT/installed/${vcpkg_target}/lib
+      mkdir -p $VCPKG_ROOT/installed/vcpkg/updates
+      ln -s ${updates_vcpkg_file} $VCPKG_ROOT/installed/vcpkg/status
+      mkdir -p $VCPKG_ROOT/installed/vcpkg/info
+      touch $VCPKG_ROOT/installed/vcpkg/info/libyuv_1.0_${vcpkg_target}.list
+      touch $VCPKG_ROOT/installed/vcpkg/info/libvpx_1.0_${vcpkg_target}.list
+
+      ln -s ${libvpx.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/
+      ln -s ${libyuv.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/
     '';
-  in ''
-    export VCPKG_ROOT="$TMP/vcpkg";
-
-    mkdir -p $VCPKG_ROOT/.vcpkg-root
-    mkdir -p $VCPKG_ROOT/installed/${vcpkg_target}/lib
-    mkdir -p $VCPKG_ROOT/installed/vcpkg/updates
-    ln -s ${updates_vcpkg_file} $VCPKG_ROOT/installed/vcpkg/status
-    mkdir -p $VCPKG_ROOT/installed/vcpkg/info
-    touch $VCPKG_ROOT/installed/vcpkg/info/libyuv_1.0_${vcpkg_target}.list
-    touch $VCPKG_ROOT/installed/vcpkg/info/libvpx_1.0_${vcpkg_target}.list
-
-    ln -s ${libvpx.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/
-    ln -s ${libyuv.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/
-  '';
 
   nativeBuildInputs = [ pkg-config cmake makeWrapper copyDesktopItems yasm nasm clang wrapGAppsHook ];
   buildInputs = [ alsa-lib pulseaudio libXfixes libxcb xdotool gtk3 libvpx libopus libXtst libyuv ];
@@ -91,15 +104,22 @@ rustPlatform.buildRustPackage rec {
   # Checks require an active X display.
   doCheck = false;
 
-  desktopItems = [ (makeDesktopItem {
-    name = "rustdesk";
-    exec = meta.mainProgram;
-    icon = "rustdesk";
-    desktopName = "RustDesk";
-    comment = meta.description;
-    genericName = "Remote Desktop";
-    categories = ["Network"];
-  }) ];
+  desktopItems = [
+    (makeDesktopItem {
+      name = "rustdesk";
+      exec = meta.mainProgram;
+      icon = "rustdesk";
+      desktopName = "RustDesk";
+      comment = meta.description;
+      genericName = "Remote Desktop";
+      categories = [ "Network" ];
+    })
+  ];
+
+  postPatch = ''
+    rm Cargo.lock
+    ln -s ${./Cargo.lock} Cargo.lock
+  '';
 
   # Add static ui resources and libsciter to same folder as binary so that it
   # can find them.
@@ -122,7 +142,7 @@ rustPlatform.buildRustPackage rec {
     description = "Yet another remote desktop software";
     homepage = "https://rustdesk.com";
     license = licenses.gpl3Only;
-    maintainers = with maintainers; [ leixb ];
+    maintainers = with maintainers; [ ocfox leixb ];
     platforms = [ "x86_64-linux" ];
     mainProgram = "rustdesk";
   };
diff --git a/pkgs/applications/virtualization/docker/default.nix b/pkgs/applications/virtualization/docker/default.nix
index 17b3982a1f7..0d846e9a696 100644
--- a/pkgs/applications/virtualization/docker/default.nix
+++ b/pkgs/applications/virtualization/docker/default.nix
@@ -254,6 +254,11 @@ rec {
     meta = with lib; {
       homepage = "https://www.docker.com/";
       description = "An open source project to pack, ship and run any application as a lightweight container";
+      longDescription = ''
+        Docker is a platform designed to help developers build, share, and run modern applications.
+
+        To enable the docker daemon on NixOS, set the `virtualisation.docker.enable` option to `true`.
+      '';
       license = licenses.asl20;
       maintainers = with maintainers; [ offline tailhook vdemeester periklis mikroskeem maxeaubrey ];
     };
diff --git a/pkgs/build-support/fetchmavenartifact/default.nix b/pkgs/build-support/fetchmavenartifact/default.nix
index b1d8593e72d..0f3cd4e64dd 100644
--- a/pkgs/build-support/fetchmavenartifact/default.nix
+++ b/pkgs/build-support/fetchmavenartifact/default.nix
@@ -37,13 +37,8 @@ assert (url == "") || (urls == []);
 assert (repos != []) || (url != "") || (urls != []);
 
 let
-  name_ =
-    lib.concatStrings [
-      (lib.replaceStrings ["."] ["_"] groupId) "_"
-      (lib.replaceStrings ["."] ["_"] artifactId) "-"
-      version
-    ];
-  suffix = if isNull classifier then "" else "-${classifier}";
+  pname = (lib.replaceStrings [ "." ] [ "_" ] groupId) + "_" + (lib.replaceStrings [ "." ] [ "_" ] artifactId);
+  suffix = lib.optionalString (classifier != null) "-${classifier}";
   filename = "${artifactId}-${version}${suffix}.jar";
   mkJarUrl = repoUrl:
     lib.concatStringsSep "/" [
@@ -59,13 +54,13 @@ let
     else map mkJarUrl repos;
   jar =
     fetchurl (
-      builtins.removeAttrs args ["groupId" "artifactId" "version" "classifier" "repos" "url" ]
-        // { urls = urls_; name = "${name_}.jar"; }
+      builtins.removeAttrs args [ "groupId" "artifactId" "version" "classifier" "repos" "url" ]
+      // { urls = urls_; name = "${pname}-${version}.jar"; }
     );
 in
   stdenv.mkDerivation {
-    name = name_;
-    phases = "installPhase fixupPhase";
+    inherit pname version;
+    dontUnpack = true;
     # By moving the jar to $out/share/java we make it discoverable by java
     # packages packages that mention this derivation in their buildInputs.
     installPhase = ''
diff --git a/pkgs/development/compilers/llvm/rocm/llvm.nix b/pkgs/development/compilers/llvm/rocm/llvm.nix
index 1f1add5cf67..6092bc1a9fc 100644
--- a/pkgs/development/compilers/llvm/rocm/llvm.nix
+++ b/pkgs/development/compilers/llvm/rocm/llvm.nix
@@ -24,6 +24,8 @@
 , targetDir ? "llvm"
 , targetProjects ? [ ]
 , targetRuntimes ? [ ]
+# "NATIVE" resolves into x86 or aarch64 depending on stdenv
+, llvmTargetsToBuild ? [ "NATIVE" ]
 , extraPatches ? [ ]
 , extraNativeBuildInputs ? [ ]
 , extraBuildInputs ? [ ]
@@ -46,6 +48,8 @@ let
     if stdenv.isx86_64 then "X86"
     else if stdenv.isAarch64 then "AArch64"
     else throw "Unsupported ROCm LLVM platform";
+  inferNativeTarget = t: if t == "NATIVE" then llvmNativeTarget else t;
+  llvmTargetsToBuild' = [ "AMDGPU" ] ++ builtins.map inferNativeTarget llvmTargetsToBuild;
 in stdenv.mkDerivation (finalAttrs: {
   pname = "rocm-llvm-${targetName}";
   version = "5.4.4";
@@ -98,7 +102,7 @@ in stdenv.mkDerivation (finalAttrs: {
   sourceRoot = "${finalAttrs.src.name}/${targetDir}";
 
   cmakeFlags = [
-    "-DLLVM_TARGETS_TO_BUILD=AMDGPU;${llvmNativeTarget}"
+    "-DLLVM_TARGETS_TO_BUILD=${builtins.concatStringsSep ";" llvmTargetsToBuild'}"
   ] ++ lib.optionals (finalAttrs.passthru.isLLVM && targetProjects != [ ]) [
     "-DLLVM_ENABLE_PROJECTS=${lib.concatStringsSep ";" targetProjects}"
   ] ++ lib.optionals ((finalAttrs.passthru.isLLVM || targetDir == "runtimes") && targetRuntimes != [ ]) [
diff --git a/pkgs/development/libraries/libiio/default.nix b/pkgs/development/libraries/libiio/default.nix
index 815f23a0527..d217fbd8238 100644
--- a/pkgs/development/libraries/libiio/default.nix
+++ b/pkgs/development/libraries/libiio/default.nix
@@ -37,19 +37,19 @@ stdenv.mkDerivation rec {
     flex
     bison
     pkg-config
-  ];
+    python
+  ] ++ lib.optional python.isPy3k python.pkgs.setuptools;
 
   buildInputs = [
-    python
     libxml2
     libusb1
-  ] ++ lib.optional python.isPy3k python.pkgs.setuptools
-    ++ lib.optional avahiSupport avahi
+  ] ++ lib.optional avahiSupport avahi
     ++ lib.optional stdenv.isLinux libaio
     ++ lib.optionals stdenv.isDarwin [ CFNetwork CoreServices ];
 
   cmakeFlags = [
     "-DUDEV_RULES_INSTALL_DIR=${placeholder "out"}/lib/udev/rules.d"
+    "-DPython_EXECUTABLE=${python.pythonForBuild.interpreter}"
     "-DPYTHON_BINDINGS=on"
     # osx framework is disabled,
     # the linux-like directory structure is used for proper output splitting
diff --git a/pkgs/development/libraries/ndpi/default.nix b/pkgs/development/libraries/ndpi/default.nix
index b798320c691..d9d8f5895a2 100644
--- a/pkgs/development/libraries/ndpi/default.nix
+++ b/pkgs/development/libraries/ndpi/default.nix
@@ -1,31 +1,39 @@
 { lib
 , stdenv
-, fetchFromGitHub
-, which
 , autoconf
 , automake
-, libtool
-, libpcap
+, fetchFromGitHub
 , json_c
-, pkg-config }:
+, libpcap
+, libtool
+, pkg-config
+, which
+}:
 
 stdenv.mkDerivation rec {
   pname = "ndpi";
-  version = "4.2";
+  version = "4.6";
 
   src = fetchFromGitHub {
     owner = "ntop";
     repo = "nDPI";
-    rev = version;
-    sha256 = "sha256-ZWWuyPGl+hbrfXdtPvCBqMReuJ4FiGx+qiI7qCz6wtQ=";
+    rev = "refs/tags/${version}";
+    hash = "sha256-S0lVh5FZewPbYG/1ikI2RroCSC7OI8Xmfeq73hYCHnY=";
   };
 
   configureScript = "./autogen.sh";
 
-  nativeBuildInputs = [ which autoconf automake libtool pkg-config ];
+  nativeBuildInputs = [
+    autoconf
+    automake
+    libtool
+    pkg-config
+    which
+  ];
+
   buildInputs = [
-    libpcap
     json_c
+    libpcap
   ];
 
   meta = with lib; {
@@ -34,6 +42,7 @@ stdenv.mkDerivation rec {
       nDPI is a library for deep-packet inspection based on OpenDPI.
     '';
     homepage = "https://www.ntop.org/products/deep-packet-inspection/ndpi/";
+    changelog = "https://github.com/ntop/nDPI/blob/${version}/CHANGELOG.md";
     license = with licenses; [ lgpl3Plus bsd3 ];
     maintainers = with maintainers; [ takikawa ];
     mainProgram = "ndpiReader";
diff --git a/pkgs/development/libraries/rapidfuzz-cpp/default.nix b/pkgs/development/libraries/rapidfuzz-cpp/default.nix
index 9cde99b69f0..3050d64a4fd 100644
--- a/pkgs/development/libraries/rapidfuzz-cpp/default.nix
+++ b/pkgs/development/libraries/rapidfuzz-cpp/default.nix
@@ -5,14 +5,14 @@
 , catch2_3
 }:
 
-stdenv.mkDerivation rec {
+stdenv.mkDerivation (finalAttrs: {
   pname = "rapidfuzz-cpp";
   version = "1.10.4";
 
   src = fetchFromGitHub {
     owner = "maxbachmann";
     repo = "rapidfuzz-cpp";
-    rev = "v${version}";
+    rev = "v${finalAttrs.version}";
     hash = "sha256-I7MdeLs+J5a57ypgUJIW0/pSFPzK4nZA4ZrVRdKX7J4=";
   };
 
@@ -20,7 +20,7 @@ stdenv.mkDerivation rec {
     cmake
   ];
 
-  cmakeFlags = lib.optionals doCheck [
+  cmakeFlags = lib.optionals finalAttrs.finalPackage.doCheck [
     "-DRAPIDFUZZ_BUILD_TESTING=ON"
   ];
 
@@ -33,14 +33,12 @@ stdenv.mkDerivation rec {
     catch2_3
   ];
 
-  doCheck = stdenv.buildPlatform.canExecute stdenv.hostPlatform;
-
   meta = {
     description = "Rapid fuzzy string matching in C++ using the Levenshtein Distance";
     homepage = "https://github.com/maxbachmann/rapidfuzz-cpp";
-    changelog = "https://github.com/maxbachmann/rapidfuzz-cpp/blob/${src.rev}/CHANGELOG.md";
+    changelog = "https://github.com/maxbachmann/rapidfuzz-cpp/blob/${finalAttrs.src.rev}/CHANGELOG.md";
     license = lib.licenses.mit;
     maintainers = with lib.maintainers; [ dotlambda ];
     platforms = lib.platforms.unix;
   };
-}
+})
diff --git a/pkgs/development/libraries/science/math/tensorrt/extension.nix b/pkgs/development/libraries/science/math/tensorrt/extension.nix
index b4f47a8969c..507dc0cfd50 100644
--- a/pkgs/development/libraries/science/math/tensorrt/extension.nix
+++ b/pkgs/development/libraries/science/math/tensorrt/extension.nix
@@ -27,7 +27,9 @@ final: prev: let
     defaultBuild = { "tensorrt" = if allBuilds ? ${computeName tensorRTDefaultVersion}
       then allBuilds.${computeName tensorRTDefaultVersion}
       else throw "tensorrt-${tensorRTDefaultVersion} does not support your cuda version ${cudaVersion}"; };
-  in allBuilds // defaultBuild;
+  in {
+    inherit buildTensorRTPackage;
+  } // allBuilds // defaultBuild;
 
   tensorRTVersions = {
     "8.4.0" = [
diff --git a/pkgs/development/ocaml-modules/torch/default.nix b/pkgs/development/ocaml-modules/torch/default.nix
index e46374ee1c2..9ba356fa939 100644
--- a/pkgs/development/ocaml-modules/torch/default.nix
+++ b/pkgs/development/ocaml-modules/torch/default.nix
@@ -2,6 +2,7 @@
 , stdenv
 , buildDunePackage
 , fetchFromGitHub
+, fetchpatch
 , cmdliner
 , ctypes
 , dune-configurator
@@ -24,11 +25,19 @@ buildDunePackage rec {
 
   src = fetchFromGitHub {
     owner = "LaurentMazare";
-    repo   = "ocaml-${pname}";
-    rev    = version;
+    repo = "ocaml-${pname}";
+    rev = version;
     hash = "sha256-z/9NUBjeFWE63Z/e8OyzDiy8hrn6qzjaiBH8G9MPeos=";
   };
 
+  patches = [
+    # Pytorch 2.0 support. Drop when it reaches a release
+    (fetchpatch {
+      url = "https://github.com/LaurentMazare/ocaml-torch/commit/ef7ef30cafecb09e45ec1ed8ce4bedae5947cfa5.patch";
+      hash = "sha256-smdwKy40iIISp/25L2J4az6KmqFS1soeChBElUyhl5A=";
+    })
+  ];
+
   buildInputs = [ dune-configurator ];
 
   propagatedBuildInputs = [
diff --git a/pkgs/development/python-modules/bottle/default.nix b/pkgs/development/python-modules/bottle/default.nix
index 3d12ba7e3ae..7a97b999023 100644
--- a/pkgs/development/python-modules/bottle/default.nix
+++ b/pkgs/development/python-modules/bottle/default.nix
@@ -27,6 +27,8 @@ buildPythonPackage rec {
     "test_delete_cookie"
     "test_error"
     "test_error_in_generator_callback"
+    # timing sensitive
+    "test_ims"
   ];
 
   __darwinAllowLocalNetworking = true;
diff --git a/pkgs/development/python-modules/openai-triton/default.nix b/pkgs/development/python-modules/openai-triton/default.nix
new file mode 100644
index 00000000000..0e10642f069
--- /dev/null
+++ b/pkgs/development/python-modules/openai-triton/default.nix
@@ -0,0 +1,253 @@
+{ lib
+, buildPythonPackage
+, python
+, fetchpatch
+, fetchFromGitHub
+, addOpenGLRunpath
+, cmake
+, cudaPackages
+, llvmPackages
+, pybind11
+, gtest
+, zlib
+, ncurses
+, libxml2
+, lit
+, filelock
+, torchWithRocm
+, pytest
+, pytestCheckHook
+, pythonRelaxDepsHook
+, pkgsTargetTarget
+}:
+
+let
+  pname = "triton";
+  version = "2.0.0";
+
+  inherit (cudaPackages) cuda_cudart backendStdenv;
+
+  # A time may come we'll want to be cross-friendly
+  #
+  # Short explanation: we need pkgsTargetTarget, because we use string
+  # interpolation instead of buildInputs.
+  #
+  # Long explanation: OpenAI/triton downloads and vendors a copy of NVidia's
+  # ptxas compiler. We're not running this ptxas on the build machine, but on
+  # the user's machine, i.e. our Target platform. The second "Target" in
+  # pkgsTargetTarget maybe doesn't matter, because ptxas compiles programs to
+  # be executed on the GPU.
+  # Cf. https://nixos.org/manual/nixpkgs/unstable/#sec-cross-infra
+  ptxas = "${pkgsTargetTarget.cudaPackages.cuda_nvcc}/bin/ptxas";
+
+  llvm = (llvmPackages.llvm.override {
+    llvmTargetsToBuild = [ "NATIVE" "NVPTX" ];
+    # Upstream CI sets these too:
+    # targetProjects = [ "mlir" ];
+    extraCMakeFlags = [
+      "-DLLVM_INSTALL_UTILS=ON"
+    ];
+  });
+in
+buildPythonPackage {
+  inherit pname version;
+
+  format = "setuptools";
+
+  src = fetchFromGitHub {
+    owner = "openai";
+    repo = pname;
+    rev = "v${version}";
+    hash = "sha256-9GZzugab+Pdt74Dj6zjlEzjj4BcJ69rzMJmqcVMxsKU=";
+  };
+
+  patches = [
+    # Prerequisite for llvm15 patch
+    (fetchpatch {
+      url = "https://github.com/openai/triton/commit/2aba985daaa70234823ea8f1161da938477d3e02.patch";
+      hash = "sha256-LGv0+Ut2WYPC4Ksi4803Hwmhi3FyQOF9zElJc/JCobk=";
+    })
+    (fetchpatch {
+      url = "https://github.com/openai/triton/commit/e3941f9d09cdd31529ba4a41018cfc0096aafea6.patch";
+      hash = "sha256-A+Gor6qzFlGQhVVhiaaYOzqqx8yO2MdssnQS6TIfUWg=";
+    })
+
+    # Source: https://github.com/openai/triton/commit/fc7a8e35819bda632bdcf1cf75fd9abe4d4e077a.patch
+    # The original patch adds ptxas binary, so we include our own clean copy
+    # Drop with the next update
+    ./llvm15.patch
+
+    # TODO: there have been commits upstream aimed at removing the "torch"
+    # circular dependency, but the patches fail to apply on the release
+    # revision. Keeping the link for future reference
+    # Also cf. https://github.com/openai/triton/issues/1374
+
+    # (fetchpatch {
+    #   url = "https://github.com/openai/triton/commit/fc7c0b0e437a191e421faa61494b2ff4870850f1.patch";
+    #   hash = "sha256-f0shIqHJkVvuil2Yku7vuqWFn7VCRKFSFjYRlwx25ig=";
+    # })
+  ];
+
+  postPatch = ''
+    substituteInPlace python/setup.py \
+      --replace \
+        '= get_thirdparty_packages(triton_cache_path)' \
+        '= os.environ["cmakeFlags"].split()'
+  ''
+  # Wiring triton=2.0.0 with llcmPackages_rocm.llvm=5.4.3
+  # Revisit when updating either triton or llvm
+  + ''
+    substituteInPlace CMakeLists.txt \
+      --replace "nvptx" "NVPTX" \
+      --replace "LLVM 11" "LLVM"
+    sed -i '/AddMLIR/a set(MLIR_TABLEGEN_EXE "${llvmPackages.mlir}/bin/mlir-tblgen")' CMakeLists.txt
+    sed -i '/AddMLIR/a set(MLIR_INCLUDE_DIR ''${MLIR_INCLUDE_DIRS})' CMakeLists.txt
+    find -iname '*.td' -exec \
+      sed -i \
+      -e '\|include "mlir/IR/OpBase.td"|a include "mlir/IR/AttrTypeBase.td"' \
+      -e 's|include "mlir/Dialect/StandardOps/IR/Ops.td"|include "mlir/Dialect/Func/IR/FuncOps.td"|' \
+      '{}' ';'
+    substituteInPlace unittest/CMakeLists.txt --replace "include(GoogleTest)" "find_package(GTest REQUIRED)"
+    sed -i 's/^include.*$//' unittest/CMakeLists.txt
+    sed -i '/LINK_LIBS/i NVPTXInfo' lib/Target/PTX/CMakeLists.txt
+    sed -i '/LINK_LIBS/i NVPTXCodeGen' lib/Target/PTX/CMakeLists.txt
+  ''
+  # TritonMLIRIR already links MLIRIR. Not transitive?
+  # + ''
+  #   echo "target_link_libraries(TritonPTX PUBLIC MLIRIR)" >> lib/Target/PTX/CMakeLists.txt
+  # ''
+  # Already defined in llvm, when built with -DLLVM_INSTALL_UTILS
+  + ''
+    substituteInPlace bin/CMakeLists.txt \
+      --replace "add_subdirectory(FileCheck)" ""
+
+    rm cmake/FindLLVM.cmake
+  ''
+  +
+  (
+    let
+      # Bash was getting weird without linting,
+      # but basically upstream contains [cc, ..., "-lcuda", ...]
+      # and we replace it with [..., "-lcuda", "-L/run/opengl-driver/lib", "-L$stubs", ...]
+      old = [ "-lcuda" ];
+      new = [ "-lcuda" "-L${addOpenGLRunpath.driverLink}" "-L${cuda_cudart}/lib/stubs/" ];
+
+      quote = x: ''"${x}"'';
+      oldStr = lib.concatMapStringsSep ", " quote old;
+      newStr = lib.concatMapStringsSep ", " quote new;
+    in
+    ''
+      substituteInPlace python/triton/compiler.py \
+        --replace '${oldStr}' '${newStr}'
+    ''
+  )
+  # Triton seems to be looking up cuda.h
+  + ''
+    sed -i 's|cu_include_dir = os.path.join.*$|cu_include_dir = "${cuda_cudart}/include"|' python/triton/compiler.py
+  '';
+
+  nativeBuildInputs = [
+    cmake
+    pythonRelaxDepsHook
+
+    # Requires torch (circular dependency) and probably needs GPUs:
+    # pytestCheckHook
+
+    # Note for future:
+    # These *probably* should go in depsTargetTarget
+    # ...but we cannot test cross right now anyway
+    # because we only support cudaPackages on x86_64-linux atm
+    lit
+    llvm
+    llvmPackages.mlir
+  ];
+
+  buildInputs = [
+    gtest
+    libxml2.dev
+    ncurses
+    pybind11
+    zlib
+  ];
+
+  propagatedBuildInputs = [
+    filelock
+  ];
+
+  # Avoid GLIBCXX mismatch with other cuda-enabled python packages
+  preConfigure = ''
+    export CC="${backendStdenv.cc}/bin/cc";
+    export CXX="${backendStdenv.cc}/bin/c++";
+
+    # Upstream's setup.py tries to write cache somewhere in ~/
+    export HOME=$TMPDIR
+
+    # Upstream's github actions patch setup.cfg to write base-dir. May be redundant
+    echo "
+    [build_ext]
+    base-dir=$PWD" >> python/setup.cfg
+
+    # The rest (including buildPhase) is relative to ./python/
+    cd python/
+
+    # Work around download_and_copy_ptxas()
+    dst_cuda="$PWD/triton/third_party/cuda/bin"
+    mkdir -p "$dst_cuda"
+    ln -s "${ptxas}" "$dst_cuda/"
+  '';
+
+  # CMake is run by setup.py instead
+  dontUseCmakeConfigure = true;
+  cmakeFlags = [
+    "-DMLIR_DIR=${llvmPackages.mlir}/lib/cmake/mlir"
+  ];
+
+  postFixup =
+    let
+      ptxasDestination = "$out/${python.sitePackages}/triton/third_party/cuda/bin/ptxas";
+    in
+    # Setuptools (?) strips runpath and +x flags. Let's just restore the symlink
+    ''
+      rm -f ${ptxasDestination}
+      ln -s ${ptxas} ${ptxasDestination}
+    '';
+
+  checkInputs = [
+    cmake # ctest
+  ];
+  dontUseSetuptoolsCheck = true;
+  preCheck =
+    # build/temp* refers to build_ext.build_temp (looked up in the build logs)
+    ''
+      (cd /build/source/python/build/temp* ; ctest)
+    '' # For pytestCheckHook
+    + ''
+      cd test/unit
+    '';
+  pythonImportsCheck = [
+    # Circular dependency on torch
+    # "triton"
+    # "triton.language"
+  ];
+
+  # Ultimately, torch is our test suite:
+  passthru.tests = {
+    inherit torchWithRocm;
+  };
+
+  pythonRemoveDeps = [
+    # Circular dependency, cf. https://github.com/openai/triton/issues/1374
+    "torch"
+
+    # CLI tools without dist-info
+    "cmake"
+    "lit"
+  ];
+  meta = with lib; {
+    description = "Development repository for the Triton language and compiler";
+    homepage = "https://github.com/openai/triton/";
+    platforms = lib.platforms.unix;
+    license = licenses.mit;
+    maintainers = with maintainers; [ SomeoneSerge ];
+  };
+}
diff --git a/pkgs/development/python-modules/openai-triton/llvm15.patch b/pkgs/development/python-modules/openai-triton/llvm15.patch
new file mode 100644
index 00000000000..3e20cce2380
--- /dev/null
+++ b/pkgs/development/python-modules/openai-triton/llvm15.patch
@@ -0,0 +1,4617 @@
+From fc7a8e35819bda632bdcf1cf75fd9abe4d4e077a Mon Sep 17 00:00:00 2001
+From: Christian Sigg <chsigg@users.noreply.github.com>
+Date: Thu, 16 Feb 2023 15:40:53 +0100
+Subject: [PATCH] Rebase Triton to LLVM-15. (#1070)
+
+This PR rebases Triton from LLVM-14 to LLVM-15. Most changes are
+mechanical, except for the analysis framework changes.
+---
+ CMakeLists.txt                                |   6 +-
+ bin/CMakeLists.txt                            |   2 +-
+ bin/FileCheck/FileCheck.cpp                   |   3 +
+ bin/triton-opt.cpp                            |   6 +-
+ bin/triton-translate.cpp                      |   7 +-
+ include/triton/Analysis/Alias.h               |  21 +-
+ include/triton/Analysis/Allocation.h          |   2 +
+ include/triton/Analysis/AxisInfo.h            |  56 ++-
+ include/triton/Analysis/Utility.h             |   6 +-
+ include/triton/Conversion/Passes.td           |   4 +-
+ include/triton/Dialect/Triton/IR/Dialect.h    |   7 +-
+ .../triton/Dialect/Triton/IR/TritonDialect.td |   8 +-
+ include/triton/Dialect/Triton/IR/TritonOps.td |  12 +-
+ .../triton/Dialect/Triton/IR/TritonTypes.td   |   2 +
+ .../Dialect/Triton/Transforms/Passes.td       |   3 +-
+ include/triton/Dialect/TritonGPU/IR/Dialect.h |   4 +-
+ .../Dialect/TritonGPU/IR/TritonGPUAttrDefs.td |   7 +
+ .../Dialect/TritonGPU/IR/TritonGPUDialect.td  |   2 +-
+ .../Dialect/TritonGPU/IR/TritonGPUOps.td      |  13 +-
+ lib/Analysis/Alias.cpp                        |  14 +-
+ lib/Analysis/Allocation.cpp                   |  30 +-
+ lib/Analysis/AxisInfo.cpp                     |  79 ++--
+ lib/Analysis/CMakeLists.txt                   |   2 +-
+ lib/Analysis/Membar.cpp                       |   2 +-
+ lib/Analysis/Utility.cpp                      |  54 +++
+ .../TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp |   3 -
+ lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h |  10 +-
+ .../TritonGPUToLLVM/DotOpToLLVM.cpp           |   5 -
+ .../TritonGPUToLLVM/ElementwiseOpToLLVM.cpp   |   2 -
+ .../TritonGPUToLLVM/LoadStoreOpToLLVM.cpp     |   5 +-
+ .../TritonGPUToLLVM/ReduceOpToLLVM.cpp        |   2 -
+ .../TritonGPUToLLVM/TritonGPUToLLVM.cpp       |   7 +-
+ .../TritonGPUToLLVM/TritonGPUToLLVMBase.h     |  26 +-
+ .../TritonGPUToLLVM/TritonGPUToLLVMPass.cpp   |  52 +--
+ lib/Conversion/TritonGPUToLLVM/Utility.h      |   5 +-
+ .../TritonToTritonGPUPass.cpp                 |  69 ++--
+ lib/Dialect/Triton/IR/CMakeLists.txt          |  10 +-
+ lib/Dialect/Triton/IR/Ops.cpp                 |  34 +-
+ lib/Dialect/Triton/Transforms/Combine.cpp     |   6 +-
+ lib/Dialect/Triton/Transforms/Combine.td      |   2 +-
+ lib/Dialect/TritonGPU/IR/Dialect.cpp          |  27 +-
+ lib/Dialect/TritonGPU/Transforms/Coalesce.cpp |  20 +-
+ lib/Dialect/TritonGPU/Transforms/Combine.cpp  |   2 +-
+ lib/Dialect/TritonGPU/Transforms/Combine.td   |   1 +
+ .../Transforms/DecomposeConversions.cpp       |   2 +-
+ lib/Dialect/TritonGPU/Transforms/Pipeline.cpp |  10 +-
+ .../Transforms/ReorderInstructions.cpp        |   2 +-
+ .../Transforms/TritonGPUConversion.cpp        |  12 +-
+ .../Transforms/UpdateMmaForVolta.cpp          |   6 +-
+ lib/Dialect/TritonGPU/Transforms/Utility.cpp  |   2 +-
+ lib/Target/LLVMIR/CMakeLists.txt              |   3 +-
+ lib/Target/PTX/PTXTranslation.cpp             |   3 +
+ python/setup.py                               |  15 +-
+ python/src/triton.cc                          |  85 +++--
+ python/test/unit/language/test_core.py        |   2 +-
+ python/triton/compiler.py                     |   4 +-
+ test/Analysis/test-alias.mlir                 |  24 +-
+ test/Analysis/test-alignment.mlir             | 344 +++++++++---------
+ test/Analysis/test-allocation.mlir            |  32 +-
+ test/Analysis/test-membar.mlir                |  38 +-
+ test/Conversion/triton_ops.mlir               |  10 +-
+ test/Conversion/triton_to_tritongpu.mlir      |   6 +-
+ test/Conversion/tritongpu_to_llvm.mlir        |  94 ++---
+ test/Target/tritongpu_to_llvmir.mlir          |   4 +-
+ test/Target/tritongpu_to_ptx.mlir             |   2 +-
+ test/Triton/combine.mlir                      |  40 +-
+ test/Triton/vecadd.mlir                       |   4 +-
+ test/TritonGPU/coalesce.mlir                  |   2 +-
+ test/TritonGPU/combine.mlir                   |  38 +-
+ test/TritonGPU/loop-pipeline.mlir             |  22 +-
+ test/TritonGPU/matmul.mlir                    |   4 +-
+ test/TritonGPU/prefetch.mlir                  |   4 +-
+ test/TritonGPU/update-mma-for-volta.mlir      |   4 +-
+ test/lib/Analysis/TestAlias.cpp               |  29 +-
+ test/lib/Analysis/TestAllocation.cpp          |   5 +-
+ test/lib/Analysis/TestAxisInfo.cpp            |  51 +--
+ test/lib/Analysis/TestMembar.cpp              |   7 +-
+ 78 files changed, 808 insertions(+), 742 deletions(-)
+
+diff --git a/CMakeLists.txt b/CMakeLists.txt
+index d0d361fc7c..b281a28400 100644
+--- a/CMakeLists.txt
++++ b/CMakeLists.txt
+@@ -1,4 +1,7 @@
+ cmake_minimum_required(VERSION 3.6)
++
++cmake_policy(SET CMP0116 OLD)
++
+ include(ExternalProject)
+ 
+ set(CMAKE_CXX_STANDARD 17)
+@@ -155,7 +158,6 @@ if(TRITON_BUILD_PYTHON_MODULE)
+     endif()
+ endif()
+ 
+-
+ # # Triton
+ # file(GLOB_RECURSE LIBTRITON_SRC lib/*.cc)
+ # if (WIN32 AND TRITON_BUILD_PYTHON_MODULE)
+@@ -212,7 +214,7 @@ if(TRITON_BUILD_PYTHON_MODULE)
+     # optimizations
+     MLIRPass
+     MLIRTransforms
+-    MLIRLLVMIR
++    MLIRLLVMDialect
+     MLIRSupport
+     MLIRTargetLLVMIRExport
+     MLIRExecutionEngine
+diff --git a/bin/CMakeLists.txt b/bin/CMakeLists.txt
+index 906f635f8b..695b3479fd 100644
+--- a/bin/CMakeLists.txt
++++ b/bin/CMakeLists.txt
+@@ -48,7 +48,7 @@ llvm_update_compile_flags(triton-translate)
+          # MLIR core
+          MLIROptLib
+          MLIRIR
+-         MLIRLLVMIR
++         MLIRLLVMDialect
+          MLIRPass
+          MLIRSupport
+          MLIRTransforms
+diff --git a/bin/FileCheck/FileCheck.cpp b/bin/FileCheck/FileCheck.cpp
+index 819efc3541..9ac6f1b277 100644
+--- a/bin/FileCheck/FileCheck.cpp
++++ b/bin/FileCheck/FileCheck.cpp
+@@ -19,6 +19,7 @@
+ #include "llvm/Support/CommandLine.h"
+ #include "llvm/Support/InitLLVM.h"
+ #include "llvm/Support/Process.h"
++#include "llvm/Support/SourceMgr.h"
+ #include "llvm/Support/WithColor.h"
+ #include "llvm/Support/raw_ostream.h"
+ #include <cmath>
+@@ -360,6 +361,8 @@ static std::string GetCheckTypeAbbreviation(Check::FileCheckType Ty) {
+     return "bad-not";
+   case Check::CheckBadCount:
+     return "bad-count";
++  case Check::CheckMisspelled:
++    return "misspelled";
+   case Check::CheckNone:
+     llvm_unreachable("invalid FileCheckType");
+   }
+diff --git a/bin/triton-opt.cpp b/bin/triton-opt.cpp
+index 9f3b53b7ae..f96232e1b0 100644
+--- a/bin/triton-opt.cpp
++++ b/bin/triton-opt.cpp
+@@ -8,7 +8,7 @@
+ 
+ #include "mlir/IR/Dialect.h"
+ #include "mlir/InitAllPasses.h"
+-#include "mlir/Support/MlirOptMain.h"
++#include "mlir/Tools/mlir-opt/MlirOptMain.h"
+ 
+ namespace mlir {
+ namespace test {
+@@ -33,8 +33,8 @@ int main(int argc, char **argv) {
+   // TODO: register Triton & TritonGPU passes
+   mlir::DialectRegistry registry;
+   registry.insert<mlir::triton::TritonDialect,
+-                  mlir::triton::gpu::TritonGPUDialect, mlir::math::MathDialect,
+-                  mlir::arith::ArithmeticDialect, mlir::StandardOpsDialect,
++                  mlir::triton::gpu::TritonGPUDialect, mlir::func::FuncDialect,
++                  mlir::math::MathDialect, mlir::arith::ArithmeticDialect,
+                   mlir::scf::SCFDialect, mlir::gpu::GPUDialect>();
+ 
+   return mlir::asMainReturnCode(mlir::MlirOptMain(
+diff --git a/bin/triton-translate.cpp b/bin/triton-translate.cpp
+index 05ba15e453..56b5d65857 100644
+--- a/bin/triton-translate.cpp
++++ b/bin/triton-translate.cpp
+@@ -3,7 +3,7 @@
+ #include "mlir/IR/AsmState.h"
+ #include "mlir/IR/BuiltinOps.h"
+ #include "mlir/IR/Dialect.h"
+-#include "mlir/Parser.h"
++#include "mlir/Parser/Parser.h"
+ #include "mlir/Pass/Pass.h"
+ #include "mlir/Pass/PassManager.h"
+ #include "mlir/Support/FileUtilities.h"
+@@ -38,7 +38,7 @@ OwningOpRef<ModuleOp> loadMLIRModule(llvm::StringRef inputFilename,
+   mlir::DialectRegistry registry;
+   registry.insert<TritonDialect, triton::gpu::TritonGPUDialect,
+                   mlir::math::MathDialect, arith::ArithmeticDialect,
+-                  StandardOpsDialect, scf::SCFDialect>();
++                  scf::SCFDialect>();
+ 
+   context.appendDialectRegistry(registry);
+ 
+@@ -50,7 +50,8 @@ OwningOpRef<ModuleOp> loadMLIRModule(llvm::StringRef inputFilename,
+     context.loadAllAvailableDialects();
+     context.allowUnregisteredDialects();
+ 
+-    OwningOpRef<ModuleOp> module(parseSourceFile(sourceMgr, &context));
++    OwningOpRef<ModuleOp> module =
++        parseSourceFile<ModuleOp>(sourceMgr, &context);
+     if (!module) {
+       llvm::errs() << "Parse MLIR file failed.";
+       return nullptr;
+diff --git a/include/triton/Analysis/Alias.h b/include/triton/Analysis/Alias.h
+index fa6b906fc9..631df518bc 100644
+--- a/include/triton/Analysis/Alias.h
++++ b/include/triton/Analysis/Alias.h
+@@ -2,7 +2,7 @@
+ #define TRITON_ANALYSIS_ALIAS_H
+ 
+ #include "mlir/Analysis/AliasAnalysis.h"
+-#include "mlir/Analysis/DataFlowAnalysis.h"
++#include "mlir/Analysis/DataFlow/SparseAnalysis.h"
+ #include "llvm/ADT/DenseSet.h"
+ 
+ namespace mlir {
+@@ -21,7 +21,7 @@ class AliasInfo {
+   }
+ 
+   /// The pessimistic value state of a value without alias
+-  static AliasInfo getPessimisticValueState(MLIRContext *context) {
++  static AliasInfo getPessimisticValueState(MLIRContext *context = nullptr) {
+     return AliasInfo();
+   }
+   static AliasInfo getPessimisticValueState(Value value) { return AliasInfo(); }
+@@ -29,6 +29,10 @@ class AliasInfo {
+   /// The union of both arguments
+   static AliasInfo join(const AliasInfo &lhs, const AliasInfo &rhs);
+ 
++  void print(raw_ostream &os) const {
++    llvm::interleaveComma(allocs, os, [&](Value alloc) { alloc.print(os); });
++  }
++
+ private:
+   /// The set of allocated values that are aliased by this lattice.
+   /// For now, we only consider aliased value produced by the following
+@@ -58,9 +62,13 @@ class AliasInfo {
+ //===----------------------------------------------------------------------===//
+ // Shared Memory Alias Analysis
+ //===----------------------------------------------------------------------===//
+-class SharedMemoryAliasAnalysis : public ForwardDataFlowAnalysis<AliasInfo> {
++class SharedMemoryAliasAnalysis
++    : public dataflow::SparseDataFlowAnalysis<dataflow::Lattice<AliasInfo>> {
+ public:
+-  using ForwardDataFlowAnalysis<AliasInfo>::ForwardDataFlowAnalysis;
++  using dataflow::SparseDataFlowAnalysis<
++      dataflow::Lattice<AliasInfo>>::SparseDataFlowAnalysis;
++  using dataflow::SparseDataFlowAnalysis<
++      dataflow::Lattice<AliasInfo>>::getLatticeElement;
+ 
+   /// XXX(Keren): Compatible interface with MLIR AliasAnalysis for future use.
+   /// Given two values, returns their aliasing behavior.
+@@ -70,9 +78,10 @@ class SharedMemoryAliasAnalysis : public ForwardDataFlowAnalysis<AliasInfo> {
+   ModRefResult getModRef(Operation *op, Value location);
+ 
+   /// Computes if the alloc set of the results are changed.
+-  ChangeResult
++  void
+   visitOperation(Operation *op,
+-                 ArrayRef<LatticeElement<AliasInfo> *> operands) override;
++                 ArrayRef<const dataflow::Lattice<AliasInfo> *> operands,
++                 ArrayRef<dataflow::Lattice<AliasInfo> *> results) override;
+ };
+ 
+ } // namespace mlir
+diff --git a/include/triton/Analysis/Allocation.h b/include/triton/Analysis/Allocation.h
+index b7c136d602..89b77034cc 100644
+--- a/include/triton/Analysis/Allocation.h
++++ b/include/triton/Analysis/Allocation.h
+@@ -188,6 +188,8 @@ class Allocation {
+   friend class triton::AllocationAnalysis;
+ };
+ 
++template <typename T> Interval(T, T) -> Interval<T>;
++
+ } // namespace mlir
+ 
+ #endif // TRITON_ANALYSIS_ALLOCATION_H
+diff --git a/include/triton/Analysis/AxisInfo.h b/include/triton/Analysis/AxisInfo.h
+index fdfbd8fbb3..7083b9c43b 100644
+--- a/include/triton/Analysis/AxisInfo.h
++++ b/include/triton/Analysis/AxisInfo.h
+@@ -1,9 +1,10 @@
+ #ifndef TRITON_ANALYSIS_AXISINFO_H
+ #define TRITON_ANALYSIS_AXISINFO_H
+ 
+-#include "mlir/Analysis/DataFlowAnalysis.h"
++#include "mlir/Analysis/DataFlow/SparseAnalysis.h"
+ #include "llvm/Support/raw_ostream.h"
+ 
++#include "mlir/Support/LLVM.h"
+ #include "triton/Analysis/Utility.h"
+ #include "triton/Dialect/Triton/IR/Dialect.h"
+ #include "triton/Dialect/TritonGPU/IR/Dialect.h"
+@@ -62,7 +63,7 @@ class AxisInfo {
+   }
+ 
+   /// The pessimistic value state of the contiguity is unknown.
+-  static AxisInfo getPessimisticValueState(MLIRContext *context) {
++  static AxisInfo getPessimisticValueState(MLIRContext *context = nullptr) {
+     return AxisInfo();
+   }
+   static AxisInfo getPessimisticValueState(Value value);
+@@ -70,6 +71,22 @@ class AxisInfo {
+   /// The gcd of both arguments for each dimension
+   static AxisInfo join(const AxisInfo &lhs, const AxisInfo &rhs);
+ 
++  void print(raw_ostream &os) const {
++    auto print = [&](StringRef name, DimVectorT vec) {
++      os << name << " = [";
++      llvm::interleaveComma(vec, os);
++      os << "]";
++    };
++    print("contiguity", contiguity);
++    print(", divisibility", divisibility);
++    print(", constancy", constancy);
++    os << ", constant_value = ";
++    if (constantValue)
++      os << *constantValue;
++    else
++      os << "<none>";
++  }
++
+ private:
+   /// The _contiguity_ information maps the `d`-th
+   /// dimension to the length of the shortest
+@@ -147,7 +164,8 @@ class AxisInfoVisitor {
+   }
+ 
+   virtual AxisInfo
+-  getAxisInfo(Operation *op, ArrayRef<LatticeElement<AxisInfo> *> operands) = 0;
++  getAxisInfo(Operation *op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) = 0;
+ 
+   virtual bool match(Operation *op) = 0;
+ };
+@@ -157,15 +175,16 @@ template <typename OpTy> class AxisInfoVisitorImpl : public AxisInfoVisitor {
+ public:
+   using AxisInfoVisitor::AxisInfoVisitor;
+ 
+-  AxisInfo getAxisInfo(Operation *op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) final {
++  AxisInfo
++  getAxisInfo(Operation *op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) final {
+     return getAxisInfo(cast<OpTy>(op), operands);
+   }
+ 
+   bool match(Operation *op) final { return isa<OpTy>(op); }
+ 
+-  virtual AxisInfo getAxisInfo(OpTy op,
+-                               ArrayRef<LatticeElement<AxisInfo> *> operands) {
++  virtual AxisInfo
++  getAxisInfo(OpTy op, ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) {
+     llvm_unreachable("Unimplemented getAxisInfo");
+   }
+ };
+@@ -176,8 +195,9 @@ class BinaryOpVisitorImpl : public AxisInfoVisitorImpl<OpTy> {
+ public:
+   using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(OpTy op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(OpTy op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     auto lhsInfo = operands[0]->getValue();
+     auto rhsInfo = operands[1]->getValue();
+     auto rank = lhsInfo.getRank();
+@@ -230,7 +250,8 @@ class AxisInfoVisitorList {
+     (visitors.emplace_back(std::make_unique<Ts>()), ...);
+   }
+ 
+-  AxisInfo apply(Operation *op, ArrayRef<LatticeElement<AxisInfo> *> operands) {
++  AxisInfo apply(Operation *op,
++                 ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) {
+     for (auto &visitor : visitors)
+       if (visitor->match(op))
+         return visitor->getAxisInfo(op, operands);
+@@ -241,16 +262,19 @@ class AxisInfoVisitorList {
+   std::vector<std::unique_ptr<AxisInfoVisitor>> visitors;
+ };
+ 
+-class AxisInfoAnalysis : public ForwardDataFlowAnalysis<AxisInfo> {
++class AxisInfoAnalysis
++    : public dataflow::SparseDataFlowAnalysis<dataflow::Lattice<AxisInfo>> {
+ private:
+   AxisInfoVisitorList visitors;
+ 
+ public:
+-  AxisInfoAnalysis(MLIRContext *context);
++  AxisInfoAnalysis(DataFlowSolver &solver);
++  using dataflow::SparseDataFlowAnalysis<
++      dataflow::Lattice<AxisInfo>>::getLatticeElement;
+ 
+-  ChangeResult
+-  visitOperation(Operation *op,
+-                 ArrayRef<LatticeElement<AxisInfo> *> operands) override;
++  void visitOperation(Operation *op,
++                      ArrayRef<const dataflow::Lattice<AxisInfo> *> operands,
++                      ArrayRef<dataflow::Lattice<AxisInfo> *> results) override;
+ 
+   unsigned getPtrContiguity(Value ptr);
+ 
+@@ -261,4 +285,4 @@ class AxisInfoAnalysis : public ForwardDataFlowAnalysis<AxisInfo> {
+ 
+ } // namespace mlir
+ 
+-#endif
+\ No newline at end of file
++#endif
+diff --git a/include/triton/Analysis/Utility.h b/include/triton/Analysis/Utility.h
+index c5ac137dc1..ee7fadb59d 100644
+--- a/include/triton/Analysis/Utility.h
++++ b/include/triton/Analysis/Utility.h
+@@ -1,6 +1,7 @@
+ #ifndef TRITON_ANALYSIS_UTILITY_H
+ #define TRITON_ANALYSIS_UTILITY_H
+ 
++#include "mlir/Analysis/DataFlowFramework.h"
+ #include "mlir/Analysis/SliceAnalysis.h"
+ #include "triton/Dialect/TritonGPU/IR/Dialect.h"
+ #include <algorithm>
+@@ -12,7 +13,7 @@ namespace mlir {
+ class ReduceOpHelper {
+ public:
+   explicit ReduceOpHelper(triton::ReduceOp op) : op(op) {
+-    srcTy = op.operand().getType().cast<RankedTensorType>();
++    srcTy = op.getOperand().getType().cast<RankedTensorType>();
+   }
+ 
+   ArrayRef<int64_t> getSrcShape() { return srcTy.getShape(); }
+@@ -103,6 +104,9 @@ SetVector<Operation *>
+ multiRootGetSlice(Operation *op, TransitiveFilter backwardFilter = nullptr,
+                   TransitiveFilter forwardFilter = nullptr);
+ 
++// Create a basic DataFlowSolver with constant and dead code analysis included.
++std::unique_ptr<DataFlowSolver> createDataFlowSolver();
++
+ } // namespace mlir
+ 
+ #endif // TRITON_ANALYSIS_UTILITY_H
+diff --git a/include/triton/Conversion/Passes.td b/include/triton/Conversion/Passes.td
+index 70bb20b78e..be00eb2dac 100644
+--- a/include/triton/Conversion/Passes.td
++++ b/include/triton/Conversion/Passes.td
+@@ -12,7 +12,6 @@ def ConvertTritonToTritonGPU: Pass<"convert-triton-to-tritongpu", "mlir::ModuleO
+ 
+     let dependentDialects = ["mlir::arith::ArithmeticDialect",
+                              "mlir::math::MathDialect",
+-                             "mlir::StandardOpsDialect",
+                              // TODO: Does this pass depend on SCF?
+                              "mlir::scf::SCFDialect",
+                              "mlir::triton::TritonDialect",
+@@ -41,8 +40,7 @@ def ConvertTritonGPUToLLVM : Pass<"convert-triton-gpu-to-llvm", "mlir::ModuleOp"
+                              "mlir::tensor::TensorDialect",
+                              "mlir::triton::TritonDialect",
+                              "mlir::triton::gpu::TritonGPUDialect",
+-                             "mlir::NVVM::NVVMDialect",
+-                             "mlir::StandardOpsDialect"];
++                             "mlir::NVVM::NVVMDialect"];
+ 
+     let options = [
+         Option<"computeCapability", "compute-capability",
+diff --git a/include/triton/Dialect/Triton/IR/Dialect.h b/include/triton/Dialect/Triton/IR/Dialect.h
+index e8012a51df..15869e262e 100644
+--- a/include/triton/Dialect/Triton/IR/Dialect.h
++++ b/include/triton/Dialect/Triton/IR/Dialect.h
+@@ -1,14 +1,15 @@
+ #ifndef TRITON_DIALECT_TRITON_IR_DIALECT_H_
+ #define TRITON_DIALECT_TRITON_IR_DIALECT_H_
+ 
++#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
++#include "mlir/Dialect/ControlFlow/IR/ControlFlow.h"
++#include "mlir/Dialect/Func/IR/FuncOps.h"
+ #include "mlir/Dialect/Math/IR/Math.h"
+-#include "mlir/Dialect/SCF/SCF.h"
+-#include "mlir/Dialect/StandardOps/IR/Ops.h"
++#include "mlir/Dialect/SCF/IR/SCF.h"
+ #include "mlir/Dialect/Tensor/IR/Tensor.h"
+ #include "mlir/IR/BuiltinOps.h"
+ #include "mlir/IR/Dialect.h"
+ #include "mlir/Interfaces/ControlFlowInterfaces.h"
+-
+ #include "triton/Dialect/Triton/IR/Dialect.h.inc"
+ #include "triton/Dialect/Triton/IR/OpsEnums.h.inc"
+ #include "triton/Dialect/Triton/IR/Traits.h"
+diff --git a/include/triton/Dialect/Triton/IR/TritonDialect.td b/include/triton/Dialect/Triton/IR/TritonDialect.td
+index 07b069e14f..d98ce73884 100644
+--- a/include/triton/Dialect/Triton/IR/TritonDialect.td
++++ b/include/triton/Dialect/Triton/IR/TritonDialect.td
+@@ -25,12 +25,9 @@ def Triton_Dialect : Dialect {
+   let dependentDialects = [
+     "arith::ArithmeticDialect",
+     "math::MathDialect",
+-    "StandardOpsDialect",
+     "scf::SCFDialect",
+-
+-    // Since LLVM 15
+-    // "cf::ControlFlowDialect",
+-    // "func::FuncDialect"
++    "cf::ControlFlowDialect",
++    "func::FuncDialect"
+   ];
+ 
+   let extraClassDeclaration = [{
+@@ -38,6 +35,7 @@ def Triton_Dialect : Dialect {
+   }];
+ 
+   let hasConstantMaterializer = 1;
++  let useDefaultTypePrinterParser = 1;
+ }
+ 
+ include "triton/Dialect/Triton/IR/TritonTypes.td"
+diff --git a/include/triton/Dialect/Triton/IR/TritonOps.td b/include/triton/Dialect/Triton/IR/TritonOps.td
+index 779e0b648c..0a69211179 100644
+--- a/include/triton/Dialect/Triton/IR/TritonOps.td
++++ b/include/triton/Dialect/Triton/IR/TritonOps.td
+@@ -141,11 +141,7 @@ def TT_LoadOp : TT_Op<"load",
+                        "triton::EvictionPolicy":$evict, "bool":$isVolatile)>,
+     ];
+ 
+-    // let assemblyFormat = "operands attr-dict `:` type($result)";
+-    let parser = [{ return mlir::triton::parseLoadOp(parser, result); }];
+-
+-    let printer = [{ return mlir::triton::printLoadOp(p, *this); }];
+-
++    let hasCustomAssemblyFormat = 1;
+     let hasCanonicalizer = 1;
+ }
+ 
+@@ -170,11 +166,7 @@ def TT_StoreOp : TT_Op<"store",
+                        "triton::EvictionPolicy":$evict)>,
+     ];
+ 
+-    // let assemblyFormat = "operands attr-dict `:` type($value)";
+-    let parser = [{ return mlir::triton::parseStoreOp(parser, result); }];
+-
+-    let printer = [{ return mlir::triton::printStoreOp(p, *this); }];
+-
++    let hasCustomAssemblyFormat = 1;
+     let hasCanonicalizer = 1;
+ }
+ 
+diff --git a/include/triton/Dialect/Triton/IR/TritonTypes.td b/include/triton/Dialect/Triton/IR/TritonTypes.td
+index 66d2a7b9a9..2fe2fd077d 100644
+--- a/include/triton/Dialect/Triton/IR/TritonTypes.td
++++ b/include/triton/Dialect/Triton/IR/TritonTypes.td
+@@ -1,6 +1,7 @@
+ #ifndef TRITON_TYPES
+ #define TRITON_TYPES
+ 
++include "mlir/IR/AttrTypeBase.td"
+ include "triton/Dialect/Triton/IR/TritonDialect.td"
+ 
+ //
+@@ -58,6 +59,7 @@ def TT_Ptr : TritonTypeDef<"Pointer", "ptr"> {
+         }]>
+     ];
+ 
++    let hasCustomAssemblyFormat = 1;
+     let skipDefaultBuilders = 1;
+ }
+ def TT_PtrTensor : TensorOf<[TT_Ptr]>;
+diff --git a/include/triton/Dialect/Triton/Transforms/Passes.td b/include/triton/Dialect/Triton/Transforms/Passes.td
+index 8f77aed774..a25cdc5680 100644
+--- a/include/triton/Dialect/Triton/Transforms/Passes.td
++++ b/include/triton/Dialect/Triton/Transforms/Passes.td
+@@ -16,8 +16,7 @@ def TritonCombineOps : Pass</*cli-arg*/"triton-combine", /*Op*/"mlir::ModuleOp">
+ 
+   let constructor = "mlir::triton::createCombineOpsPass()";
+ 
+-  let dependentDialects = ["mlir::arith::ArithmeticDialect",
+-                           /*SelectOp*/"mlir::StandardOpsDialect"];
++  let dependentDialects = ["mlir::arith::ArithmeticDialect"];
+ }
+ 
+ #endif
+diff --git a/include/triton/Dialect/TritonGPU/IR/Dialect.h b/include/triton/Dialect/TritonGPU/IR/Dialect.h
+index b4c8daec7b..dfc5f53ab1 100644
+--- a/include/triton/Dialect/TritonGPU/IR/Dialect.h
++++ b/include/triton/Dialect/TritonGPU/IR/Dialect.h
+@@ -1,19 +1,17 @@
+ #ifndef TRITON_DIALECT_TRITONGPU_IR_DIALECT_H_
+ #define TRITON_DIALECT_TRITONGPU_IR_DIALECT_H_
+ 
+-#include "mlir/Dialect/GPU/GPUDialect.h"
++#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+ #include "mlir/Dialect/Tensor/IR/Tensor.h"
+ #include "mlir/IR/BuiltinOps.h"
+ #include "mlir/IR/Dialect.h"
+ 
+ // TritonGPU depends on Triton
+ #include "triton/Dialect/Triton/IR/Dialect.h"
+-
+ #include "triton/Dialect/TritonGPU/IR/Dialect.h.inc"
+ #include "triton/Dialect/TritonGPU/IR/Traits.h"
+ 
+ #define GET_ATTRDEF_CLASSES
+-#include "triton/Dialect/Triton/IR/AttrInterfaces.h.inc"
+ #include "triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.h.inc"
+ 
+ #define GET_OP_CLASSES
+diff --git a/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td b/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td
+index 0242c3cc17..af2aeb03a8 100644
+--- a/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td
++++ b/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td
+@@ -1,6 +1,7 @@
+ #ifndef TRITONGPU_ATTRDEFS
+ #define TRITONGPU_ATTRDEFS
+ 
++include "mlir/IR/AttrTypeBase.td"
+ include "triton/Dialect/TritonGPU/IR/TritonGPUDialect.td"
+ include "triton/Dialect/Triton/IR/TritonInterfaces.td"
+ 
+@@ -136,6 +137,7 @@ A_{3, 2}  A_{3, 3}  A_{3, 0}  A_{3, 1} ...   [phase 1] /
+   ];
+ 
+   let extraClassDeclaration = extraBaseClassDeclaration;
++  let hasCustomAssemblyFormat = 1;
+ }
+ 
+ //===----------------------------------------------------------------------===//
+@@ -273,6 +275,7 @@ for
+     // ArrayRefParameter<"unsigned">:$sizePerCTA
+   );
+ 
++  let hasCustomAssemblyFormat = 1;
+ }
+ 
+ //===----------------------------------------------------------------------===//
+@@ -422,6 +425,7 @@ For example, the matrix L corresponding to blockTileSize=[32,16] is:
+     static constexpr int numBitsToHoldMmaV1ID{5};
+   }];
+ 
++  let hasCustomAssemblyFormat = 1;
+ }
+ 
+ def SliceEncodingAttr : DistributedEncoding<"SliceEncoding"> {
+@@ -456,6 +460,8 @@ def SliceEncodingAttr : DistributedEncoding<"SliceEncoding"> {
+     template<class T>
+     SmallVector<T> paddedShape(ArrayRef<T> shape) const;
+   }];
++
++  let hasCustomAssemblyFormat = 1;
+ }
+ 
+ def DotOperandEncodingAttr : DistributedEncoding<"DotOperandEncoding"> {
+@@ -492,6 +498,7 @@ section 9.7.13.4.1 for more details.
+ 
+   ];
+ 
++  let hasCustomAssemblyFormat = 1;
+   let extraClassDeclaration = extraBaseClassDeclaration;
+ }
+ 
+diff --git a/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td b/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td
+index 87ec1d36c6..6489a721b4 100644
+--- a/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td
++++ b/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td
+@@ -30,7 +30,7 @@ def TritonGPU_Dialect : Dialect {
+     }
+   }];
+   
+-
++  let useDefaultAttributePrinterParser = 1;
+ }
+ 
+ #endif
+diff --git a/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td b/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td
+index 510f8d0183..7aba11dc75 100644
+--- a/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td
++++ b/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td
+@@ -59,7 +59,7 @@ def TTG_AsyncCommitGroupOp : TTG_Op<"async_commit_group"> {
+ // This is needed because these ops don't
+ // handle encodings
+ // e.g., https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/Dialect/Arith/IR/ArithOps.td#L111
+-def TTG_CmpIOp : TTG_Op<"cmpi", [NoSideEffect, Elementwise, 
++def TTG_CmpIOp : TTG_Op<"cmpi", [NoSideEffect, Elementwise,
+                                  SameOperandsAndResultShape, 
+                                  SameOperandsAndResultEncoding]> {
+   let summary = "integer comparison operation";
+@@ -73,7 +73,7 @@ def TTG_CmpIOp : TTG_Op<"cmpi", [NoSideEffect, Elementwise,
+   let results = (outs TT_BoolLike:$result);
+ }
+ 
+-def TTG_CmpFOp : TTG_Op<"cmpf", [NoSideEffect, Elementwise, 
++def TTG_CmpFOp : TTG_Op<"cmpf", [NoSideEffect, Elementwise,
+                                  SameOperandsAndResultShape, 
+                                  SameOperandsAndResultEncoding]> {
+   let summary = "floating-point comparison operation";
+@@ -88,8 +88,8 @@ def TTG_CmpFOp : TTG_Op<"cmpf", [NoSideEffect, Elementwise,
+ }
+ 
+ // TODO: migrate to arith::SelectOp on LLVM16
+-def TTG_SelectOp : TTG_Op<"select", [NoSideEffect, Elementwise, 
+-                                     SameOperandsAndResultShape, 
++def TTG_SelectOp : TTG_Op<"select", [NoSideEffect, Elementwise,
++                                     SameOperandsAndResultShape,
+                                      SameOperandsAndResultEncoding]> {
+   let summary = "select operation";
+ 
+@@ -188,10 +188,7 @@ def TTG_InsertSliceAsyncOp : TTG_Op<"insert_slice_async",
+     }
+   }];
+ 
+-  // The custom parser could be replaced with oilist in LLVM-16
+-  let parser = [{ return parseInsertSliceAsyncOp(parser, result); }];
+-
+-  let printer = [{ return printInsertSliceAsyncOp(p, *this); }];
++  let hasCustomAssemblyFormat = 1;
+ }
+ 
+ def TTG_AllocTensorOp : TTG_Op<"alloc_tensor", [MemoryEffects<[MemAlloc]>,  // Allocate shared memory
+diff --git a/lib/Analysis/Alias.cpp b/lib/Analysis/Alias.cpp
+index a39e4de9aa..208fdd4afc 100644
+--- a/lib/Analysis/Alias.cpp
++++ b/lib/Analysis/Alias.cpp
+@@ -18,8 +18,9 @@ AliasInfo AliasInfo::join(const AliasInfo &lhs, const AliasInfo &rhs) {
+   return ret;
+ }
+ 
+-ChangeResult SharedMemoryAliasAnalysis::visitOperation(
+-    Operation *op, ArrayRef<LatticeElement<AliasInfo> *> operands) {
++void SharedMemoryAliasAnalysis::visitOperation(
++    Operation *op, ArrayRef<const dataflow::Lattice<AliasInfo> *> operands,
++    ArrayRef<dataflow::Lattice<AliasInfo> *> results) {
+   AliasInfo aliasInfo;
+   bool pessimistic = true;
+   if (maybeSharedAllocationOp(op)) {
+@@ -44,14 +45,11 @@ ChangeResult SharedMemoryAliasAnalysis::visitOperation(
+   }
+ 
+   if (pessimistic) {
+-    return markAllPessimisticFixpoint(op->getResults());
++    return markAllPessimisticFixpoint(results);
+   }
+   // Join all lattice elements
+-  ChangeResult result = ChangeResult::NoChange;
+-  for (Value value : op->getResults()) {
+-    result |= getLatticeElement(value).join(aliasInfo);
+-  }
+-  return result;
++  for (auto *result : results)
++    propagateIfChanged(result, result->join(aliasInfo));
+ }
+ 
+ AliasResult SharedMemoryAliasAnalysis::alias(Value lhs, Value rhs) {
+diff --git a/lib/Analysis/Allocation.cpp b/lib/Analysis/Allocation.cpp
+index 712c08c475..b4de8dcd9d 100644
+--- a/lib/Analysis/Allocation.cpp
++++ b/lib/Analysis/Allocation.cpp
+@@ -1,4 +1,5 @@
+ #include "triton/Analysis/Allocation.h"
++#include "mlir/Analysis/DataFlowFramework.h"
+ #include "mlir/Analysis/Liveness.h"
+ #include "mlir/Analysis/SliceAnalysis.h"
+ #include "mlir/Dialect/Tensor/IR/Tensor.h"
+@@ -33,10 +34,8 @@ constexpr int kPtrBitWidth = 64;
+ 
+ static std::pair<SmallVector<unsigned>, SmallVector<unsigned>>
+ getCvtOrder(const Attribute &srcLayout, const Attribute &dstLayout) {
+-  auto srcBlockedLayout = srcLayout.dyn_cast<BlockedEncodingAttr>();
+   auto srcMmaLayout = srcLayout.dyn_cast<MmaEncodingAttr>();
+   auto srcDotLayout = srcLayout.dyn_cast<DotOperandEncodingAttr>();
+-  auto dstBlockedLayout = dstLayout.dyn_cast<BlockedEncodingAttr>();
+   auto dstMmaLayout = dstLayout.dyn_cast<MmaEncodingAttr>();
+   auto dstDotLayout = dstLayout.dyn_cast<DotOperandEncodingAttr>();
+   assert(!(srcMmaLayout && dstMmaLayout) &&
+@@ -224,14 +223,12 @@ class AllocationAnalysis {
+   }
+ 
+   void getValueAlias(Value value, SharedMemoryAliasAnalysis &analysis) {
+-    LatticeElement<AliasInfo> *latticeElement =
+-        analysis.lookupLatticeElement(value);
+-    if (latticeElement) {
+-      auto &info = latticeElement->getValue();
+-      if (!info.getAllocs().empty()) {
+-        for (auto alloc : info.getAllocs()) {
+-          allocation->addAlias(value, alloc);
+-        }
++    dataflow::Lattice<AliasInfo> *latticeElement =
++        analysis.getLatticeElement(value);
++    if (latticeElement && !latticeElement->isUninitialized()) {
++      AliasInfo &info = latticeElement->getValue();
++      for (auto alloc : info.getAllocs()) {
++        allocation->addAlias(value, alloc);
+       }
+     }
+   }
+@@ -244,14 +241,19 @@ class AllocationAnalysis {
+       getScratchValueSize(op);
+     });
+     // Get the alias values
+-    SharedMemoryAliasAnalysis aliasAnalysis(operation->getContext());
+-    aliasAnalysis.run(operation);
++    std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver();
++    SharedMemoryAliasAnalysis *aliasAnalysis =
++        solver->load<SharedMemoryAliasAnalysis>();
++    if (failed(solver->initializeAndRun(operation))) {
++      // TODO: return error instead of bailing out..
++      llvm_unreachable("failed to run SharedMemoryAliasAnalysis");
++    }
+     operation->walk<WalkOrder::PreOrder>([&](Operation *op) {
+       for (auto operand : op->getOperands()) {
+-        getValueAlias(operand, aliasAnalysis);
++        getValueAlias(operand, *aliasAnalysis);
+       }
+       for (auto value : op->getResults()) {
+-        getValueAlias(value, aliasAnalysis);
++        getValueAlias(value, *aliasAnalysis);
+       }
+     });
+   }
+diff --git a/lib/Analysis/AxisInfo.cpp b/lib/Analysis/AxisInfo.cpp
+index 0b7142b04d..4af46c3fbb 100644
+--- a/lib/Analysis/AxisInfo.cpp
++++ b/lib/Analysis/AxisInfo.cpp
+@@ -1,4 +1,4 @@
+-#include "mlir/Analysis/DataFlowAnalysis.h"
++#include "mlir/Analysis/DataFlowFramework.h"
+ #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+ #include "llvm/Support/raw_ostream.h"
+ 
+@@ -52,7 +52,7 @@ AxisInfo AxisInfo::getPessimisticValueState(Value value) {
+   BlockArgument blockArg = value.dyn_cast<BlockArgument>();
+   if (blockArg && blockArg.getOwner()->isEntryBlock()) {
+     Operation *op = blockArg.getOwner()->getParentOp();
+-    if (FuncOp fun = dyn_cast<FuncOp>(op)) {
++    if (func::FuncOp fun = dyn_cast<func::FuncOp>(op)) {
+       Attribute attr =
+           fun.getArgAttr(blockArg.getArgNumber(), "tt.divisibility");
+       if (attr)
+@@ -136,8 +136,9 @@ class CastOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> {
+ public:
+   using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(OpTy op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(OpTy op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     return operands[0]->getValue();
+   }
+ };
+@@ -147,8 +148,9 @@ class MakeRangeOpAxisInfoVisitor final
+ public:
+   using AxisInfoVisitorImpl<triton::MakeRangeOp>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(triton::MakeRangeOp op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(triton::MakeRangeOp op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     auto start = op.start();
+     auto end = op.end();
+     return AxisInfo(/*contiguity=*/{end - start},
+@@ -162,8 +164,9 @@ class ConstantOpAxisInfoVisitor final
+ public:
+   using AxisInfoVisitorImpl<arith::ConstantOp>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(arith::ConstantOp op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(arith::ConstantOp op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     auto intAttr = op.getValue().dyn_cast<IntegerAttr>();
+     auto boolAttr = op.getValue().dyn_cast<BoolAttr>();
+     if (intAttr || boolAttr) {
+@@ -416,8 +419,9 @@ class SplatOpAxisInfoVisitor final
+ public:
+   using AxisInfoVisitorImpl<triton::SplatOp>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(triton::SplatOp op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(triton::SplatOp op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     Type _retTy = *op->result_type_begin();
+     TensorType retTy = _retTy.cast<TensorType>();
+     AxisInfo opInfo = operands[0]->getValue();
+@@ -439,8 +443,9 @@ class ExpandDimsOpAxisInfoVisitor final
+ public:
+   using AxisInfoVisitorImpl<triton::ExpandDimsOp>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(triton::ExpandDimsOp op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(triton::ExpandDimsOp op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     AxisInfo opInfo = operands[0]->getValue();
+     AxisInfo::DimVectorT contiguity = opInfo.getContiguity();
+     AxisInfo::DimVectorT divisibility = opInfo.getDivisibility();
+@@ -458,8 +463,9 @@ class BroadcastOpAxisInfoVisitor final
+ public:
+   using AxisInfoVisitorImpl<triton::BroadcastOp>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(triton::BroadcastOp op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(triton::BroadcastOp op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     Type _retTy = *op->result_type_begin();
+     Type _opTy = *op->operand_type_begin();
+     TensorType retTy = _retTy.cast<TensorType>();
+@@ -486,8 +492,9 @@ class CmpOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> {
+ public:
+   using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(OpTy op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(OpTy op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     auto resTy = op.getResult().getType().template dyn_cast<RankedTensorType>();
+     if (!resTy)
+       return AxisInfo();
+@@ -596,8 +603,9 @@ class SelectOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> {
+ public:
+   using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(OpTy op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(OpTy op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     auto resTy = op.getResult().getType().template dyn_cast<RankedTensorType>();
+     if (!resTy)
+       return AxisInfo();
+@@ -757,8 +765,9 @@ class MaxMinOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> {
+ public:
+   using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl;
+ 
+-  AxisInfo getAxisInfo(OpTy op,
+-                       ArrayRef<LatticeElement<AxisInfo> *> operands) override {
++  AxisInfo
++  getAxisInfo(OpTy op,
++              ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override {
+     auto lhsInfo = operands[0]->getValue();
+     auto rhsInfo = operands[1]->getValue();
+     std::optional<int64_t> constantValue;
+@@ -786,8 +795,8 @@ class MaxMinOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> {
+ // AxisInfoAnalysis
+ //===----------------------------------------------------------------------===//
+ 
+-AxisInfoAnalysis::AxisInfoAnalysis(MLIRContext *context)
+-    : ForwardDataFlowAnalysis<AxisInfo>(context) {
++AxisInfoAnalysis::AxisInfoAnalysis(DataFlowSolver &solver)
++    : dataflow::SparseDataFlowAnalysis<dataflow::Lattice<AxisInfo>>(solver) {
+   // UnrealizedConversionCast:
+   // This is needed by TritonGPUToLLVM, to get AxisInfo when the graph is
+   // in the process of a PartialConversion, where UnrealizedConversionCast
+@@ -819,7 +828,7 @@ AxisInfoAnalysis::AxisInfoAnalysis(MLIRContext *context)
+   visitors.append<LogicalOpAxisInfoVisitor<arith::AndIOp>,
+                   LogicalOpAxisInfoVisitor<arith::OrIOp>,
+                   LogicalOpAxisInfoVisitor<arith::XOrIOp>>();
+-  visitors.append<SelectOpAxisInfoVisitor<mlir::SelectOp>,
++  visitors.append<SelectOpAxisInfoVisitor<mlir::arith::SelectOp>,
+                   SelectOpAxisInfoVisitor<triton::gpu::SelectOp>>();
+   visitors.append<ShLIOpAxisInfoVisitor, ShROpAxisInfoVisitor<arith::ShRUIOp>,
+                   ShROpAxisInfoVisitor<arith::ShRSIOp>>();
+@@ -829,11 +838,12 @@ AxisInfoAnalysis::AxisInfoAnalysis(MLIRContext *context)
+                   MaxMinOpAxisInfoVisitor<arith::MinUIOp>>();
+ }
+ 
+-ChangeResult AxisInfoAnalysis::visitOperation(
+-    Operation *op, ArrayRef<LatticeElement<AxisInfo> *> operands) {
++void AxisInfoAnalysis::visitOperation(
++    Operation *op, ArrayRef<const dataflow::Lattice<AxisInfo> *> operands,
++    ArrayRef<dataflow::Lattice<AxisInfo> *> results) {
+   AxisInfo curr = visitors.apply(op, operands);
+   if (curr.getRank() == 0) {
+-    return markAllPessimisticFixpoint(op->getResults());
++    return markAllPessimisticFixpoint(results);
+   }
+   // override with hint
+   auto newContiguity = curr.getContiguity();
+@@ -854,11 +864,8 @@ ChangeResult AxisInfoAnalysis::visitOperation(
+   curr = mlir::AxisInfo(newContiguity, newDivisibility, newConstancy,
+                         curr.getConstantValue());
+   // join all lattice elements
+-  ChangeResult result = ChangeResult::NoChange;
+-  for (Value value : op->getResults()) {
+-    result |= getLatticeElement(value).join(curr);
+-  }
+-  return result;
++  for (auto *result : results)
++    propagateIfChanged(result, result->join(curr));
+ }
+ 
+ unsigned AxisInfoAnalysis::getPtrContiguity(Value ptr) {
+@@ -884,7 +891,10 @@ unsigned AxisInfoAnalysis::getPtrAlignment(Value ptr) {
+   auto tensorTy = ptr.getType().dyn_cast<RankedTensorType>();
+   if (!tensorTy)
+     return 1;
+-  auto axisInfo = lookupLatticeElement(ptr)->getValue();
++  dataflow::Lattice<AxisInfo> *latticeElement = getLatticeElement(ptr);
++  if (!latticeElement || latticeElement->isUninitialized())
++    return 1;
++  auto axisInfo = latticeElement->getValue();
+   auto layout = tensorTy.getEncoding();
+   auto order = triton::gpu::getOrder(layout);
+   auto maxMultipleBytes = axisInfo.getDivisibility(order[0]);
+@@ -900,8 +910,11 @@ unsigned AxisInfoAnalysis::getMaskAlignment(Value mask) {
+   auto tensorTy = mask.getType().dyn_cast<RankedTensorType>();
+   if (!tensorTy)
+     return 1;
++  dataflow::Lattice<AxisInfo> *latticeElement = getLatticeElement(mask);
++  if (!latticeElement || latticeElement->isUninitialized())
++    return 1;
++  auto maskAxis = latticeElement->getValue();
+   auto maskOrder = triton::gpu::getOrder(tensorTy.getEncoding());
+-  auto maskAxis = lookupLatticeElement(mask)->getValue();
+   auto alignment = std::max<unsigned>(maskAxis.getConstancy(maskOrder[0]), 1);
+   return alignment;
+ }
+diff --git a/lib/Analysis/CMakeLists.txt b/lib/Analysis/CMakeLists.txt
+index afbc692510..1f761f845c 100644
+--- a/lib/Analysis/CMakeLists.txt
++++ b/lib/Analysis/CMakeLists.txt
+@@ -8,7 +8,7 @@ add_mlir_library(TritonAnalysis
+   DEPENDS
+   TritonTableGen
+   TritonGPUAttrDefsIncGen
+-  
++
+   LINK_LIBS PUBLIC
+   MLIRAnalysis
+ )
+diff --git a/lib/Analysis/Membar.cpp b/lib/Analysis/Membar.cpp
+index acc885e827..910274b2ac 100644
+--- a/lib/Analysis/Membar.cpp
++++ b/lib/Analysis/Membar.cpp
+@@ -2,7 +2,7 @@
+ #include "triton/Analysis/Alias.h"
+ #include "triton/Dialect/TritonGPU/IR/Dialect.h"
+ 
+-#include "mlir/Dialect/GPU/GPUDialect.h"
++#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+ #include "mlir/Dialect/Tensor/IR/Tensor.h"
+ 
+ namespace mlir {
+diff --git a/lib/Analysis/Utility.cpp b/lib/Analysis/Utility.cpp
+index d9e917e731..6ea52df272 100644
+--- a/lib/Analysis/Utility.cpp
++++ b/lib/Analysis/Utility.cpp
+@@ -1,5 +1,8 @@
+ #include "triton/Analysis/Utility.h"
++#include "mlir/Analysis/DataFlow/ConstantPropagationAnalysis.h"
++#include "mlir/Analysis/DataFlow/DeadCodeAnalysis.h"
+ #include "mlir/IR/Dialect.h"
++#include "mlir/IR/Matchers.h"
+ #include "triton/Dialect/Triton/IR/Dialect.h"
+ #include "triton/Dialect/TritonGPU/IR/Dialect.h"
+ #include <deque>
+@@ -325,4 +328,55 @@ SetVector<Operation *> multiRootGetSlice(Operation *op,
+   return multiRootTopologicalSort(slice);
+ }
+ 
++namespace {
++// Copied from TestDeadCodeAnalysis.cpp, because some dead code analysis
++// interacts with constant propagation, but SparseConstantPropagation
++// doesn't seem to be sufficient.
++struct ConstantAnalysis : public DataFlowAnalysis {
++  using DataFlowAnalysis::DataFlowAnalysis;
++
++  LogicalResult initialize(Operation *top) override {
++    WalkResult result = top->walk([&](Operation *op) {
++      if (failed(visit(op)))
++        return WalkResult::interrupt();
++      return WalkResult::advance();
++    });
++    return success(!result.wasInterrupted());
++  }
++
++  LogicalResult visit(ProgramPoint point) override {
++    Operation *op = point.get<Operation *>();
++    Attribute value;
++    if (matchPattern(op, m_Constant(&value))) {
++      auto *constant = getOrCreate<dataflow::Lattice<dataflow::ConstantValue>>(
++          op->getResult(0));
++      propagateIfChanged(constant, constant->join(dataflow::ConstantValue(
++                                       value, op->getDialect())));
++      return success();
++    }
++    setAllToUnknownConstants(op->getResults());
++    for (Region &region : op->getRegions())
++      setAllToUnknownConstants(region.getArguments());
++    return success();
++  }
++
++  /// Set all given values as not constants.
++  void setAllToUnknownConstants(ValueRange values) {
++    dataflow::ConstantValue unknownConstant(nullptr, nullptr);
++    for (Value value : values) {
++      auto *constant =
++          getOrCreate<dataflow::Lattice<dataflow::ConstantValue>>(value);
++      propagateIfChanged(constant, constant->join(unknownConstant));
++    }
++  }
++};
++} // namespace
++
++std::unique_ptr<DataFlowSolver> createDataFlowSolver() {
++  auto solver = std::make_unique<DataFlowSolver>();
++  solver->load<dataflow::DeadCodeAnalysis>();
++  solver->load<ConstantAnalysis>();
++  return solver;
++}
++
+ } // namespace mlir
+diff --git a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp
+index 6a46265bd7..e352eb3698 100644
+--- a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp
++++ b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp
+@@ -159,9 +159,6 @@ struct ConvertLayoutOpConversion
+                       Value smemBase) const {
+     auto accumNumCTAsEachRep = product<unsigned>(numCTAsEachRep);
+     auto layout = type.getEncoding();
+-    auto blockedLayout = layout.dyn_cast<BlockedEncodingAttr>();
+-    auto sliceLayout = layout.dyn_cast<SliceEncodingAttr>();
+-    auto mmaLayout = layout.dyn_cast<MmaEncodingAttr>();
+     auto rank = type.getRank();
+     auto sizePerThread = getSizePerThread(layout);
+     auto accumSizePerThread = product<unsigned>(sizePerThread);
+diff --git a/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h b/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h
+index 4b89965aa9..1d9e00519b 100644
+--- a/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h
++++ b/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h
+@@ -7,10 +7,8 @@
+ #include "mlir/Conversion/LLVMCommon/LoweringOptions.h"
+ #include "mlir/Conversion/LLVMCommon/Pattern.h"
+ #include "mlir/Conversion/MathToLLVM/MathToLLVM.h"
+-#include "mlir/Conversion/SCFToStandard/SCFToStandard.h"
+-#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
+ #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
+-#include "mlir/Dialect/GPU/GPUDialect.h"
++#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+ #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+ #include "mlir/Dialect/Tensor/IR/Tensor.h"
+ #include "mlir/IR/Matchers.h"
+@@ -422,9 +420,9 @@ struct MMA16816ConversionHelper {
+   MMA16816ConversionHelper(Type dotOperand, MmaEncodingAttr mmaLayout,
+                            Value thread, ConversionPatternRewriter &rewriter,
+                            TypeConverter *typeConverter, Location loc)
+-      : mmaLayout(mmaLayout), thread(thread), helper(mmaLayout),
+-        rewriter(rewriter), typeConverter(typeConverter), loc(loc),
+-        ctx(mmaLayout.getContext()), wpt(mmaLayout.getWarpsPerCTA()) {
++      : mmaLayout(mmaLayout), wpt(mmaLayout.getWarpsPerCTA()), thread(thread),
++        helper(mmaLayout), rewriter(rewriter), typeConverter(typeConverter),
++        loc(loc), ctx(mmaLayout.getContext()) {
+     helper.deduceMmaType(dotOperand);
+ 
+     Value _32 = i32_val(32);
+diff --git a/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp
+index 0f8070ca9f..e4bd47c411 100644
+--- a/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp
++++ b/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp
+@@ -115,8 +115,6 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern<triton::DotOp> {
+     auto DTensorTy = D.getType().cast<RankedTensorType>();
+     auto AShape = ATensorTy.getShape();
+     auto BShape = BTensorTy.getShape();
+-    auto DShape = DTensorTy.getShape();
+-    auto wpt = mmaLayout.getWarpsPerCTA();
+ 
+     bool isARow = ALayout.getIsMMAv1Row().cast<BoolAttr>().getValue();
+     bool isBRow = BLayout.getIsMMAv1Row().cast<BoolAttr>().getValue();
+@@ -221,7 +219,6 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern<triton::DotOp> {
+                               ConversionPatternRewriter &rewriter) const {
+     auto *ctx = rewriter.getContext();
+     auto loc = op.getLoc();
+-    auto threadId = getThreadId(rewriter, loc);
+ 
+     auto A = op.a();
+     auto B = op.b();
+@@ -230,12 +227,10 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern<triton::DotOp> {
+ 
+     auto aTensorTy = A.getType().cast<RankedTensorType>();
+     auto bTensorTy = B.getType().cast<RankedTensorType>();
+-    auto cTensorTy = C.getType().cast<RankedTensorType>();
+     auto dTensorTy = D.getType().cast<RankedTensorType>();
+ 
+     auto aShape = aTensorTy.getShape();
+     auto bShape = bTensorTy.getShape();
+-    auto cShape = cTensorTy.getShape();
+ 
+     BlockedEncodingAttr dLayout =
+         dTensorTy.getEncoding().cast<BlockedEncodingAttr>();
+diff --git a/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp
+index deb71b9597..0b9e67674b 100644
+--- a/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp
++++ b/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp
+@@ -61,7 +61,6 @@ struct FpToFpOpConversion
+   convertFp16x4ToFp8x4(Location loc, ConversionPatternRewriter &rewriter,
+                        const Value &v0, const Value &v1, const Value &v2,
+                        const Value &v3) {
+-    auto ctx = rewriter.getContext();
+     auto fp16x2VecTy = vec_ty(f16_ty, 2);
+     Value fp16x2Vec0 = undef(fp16x2VecTy);
+     Value fp16x2Vec1 = undef(fp16x2VecTy);
+@@ -153,7 +152,6 @@ struct FpToFpOpConversion
+   convertBf16x4ToFp8x4(Location loc, ConversionPatternRewriter &rewriter,
+                        const Value &v0, const Value &v1, const Value &v2,
+                        const Value &v3) {
+-    auto ctx = rewriter.getContext();
+     auto bf16x2VecTy = vec_ty(i16_ty, 2);
+     Value bf16x2Vec0 = undef(bf16x2VecTy);
+     Value bf16x2Vec1 = undef(bf16x2VecTy);
+diff --git a/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp
+index 9a8b4702bc..bae675f0cb 100644
+--- a/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp
++++ b/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp
+@@ -109,7 +109,8 @@ struct LoadOpConversion
+     DenseElementsAttr constAttr;
+     int64_t splatVal = 0;
+     if (other && valueElemTy.isa<IntegerType>() &&
+-        matchPattern(other, m_Constant(&constAttr)) && constAttr.isSplat()) {
++        matchPattern(other, m_Constant(&constAttr)) && constAttr.isSplat() &&
++        constAttr.getElementType().isa<IntegerType>()) {
+       otherIsSplatConstInt = true;
+       splatVal = constAttr.getSplatValue<APInt>().getSExtValue();
+     }
+@@ -333,7 +334,6 @@ struct StoreOpConversion
+             elem = rewriter.create<LLVM::SExtOp>(loc, type::i8Ty(ctx), elem);
+           elem = bitcast(elem, valueElemTy);
+ 
+-          Type u32Ty = typeConverter->convertType(type::u32Ty(ctx));
+           llWord = insert_element(wordTy, llWord, elem, i32_val(elemIdx));
+         }
+         llWord = bitcast(llWord, valArgTy);
+@@ -387,7 +387,6 @@ struct AtomicCASOpConversion
+                   ConversionPatternRewriter &rewriter) const override {
+     auto loc = op.getLoc();
+     MLIRContext *ctx = rewriter.getContext();
+-    Value ptr = op.ptr();
+ 
+     Value llPtr = adaptor.ptr();
+     Value llCmp = adaptor.cmp();
+diff --git a/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp
+index 69abd889be..1c973dc196 100644
+--- a/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp
++++ b/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp
+@@ -286,7 +286,6 @@ struct ReduceOpConversion
+     auto srcTy = op.operand().getType().cast<RankedTensorType>();
+     auto srcLayout = srcTy.getEncoding();
+     auto srcShape = srcTy.getShape();
+-    auto srcRank = srcTy.getRank();
+     auto order = getOrder(srcLayout);
+ 
+     auto threadsPerWarp = triton::gpu::getThreadsPerWarp(srcLayout);
+@@ -351,7 +350,6 @@ struct ReduceOpConversion
+ 
+     Value zero = i32_val(0);
+     Value laneZero = icmp_eq(laneIdAxis, zero);
+-    Value warpZero = icmp_eq(warpIdAxis, zero);
+ 
+     for (auto it : accs) {
+       const SmallVector<unsigned> &key = it.first;
+diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp
+index 5b77150b1a..78cfa076bd 100644
+--- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp
++++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp
+@@ -11,11 +11,11 @@ using ::mlir::LLVM::getStructFromElements;
+ using ::mlir::triton::gpu::getElemsPerThread;
+ using ::mlir::triton::gpu::SharedEncodingAttr;
+ 
+-struct ReturnOpConversion : public ConvertOpToLLVMPattern<::mlir::ReturnOp> {
+-  using ConvertOpToLLVMPattern<ReturnOp>::ConvertOpToLLVMPattern;
++struct ReturnOpConversion : public ConvertOpToLLVMPattern<func::ReturnOp> {
++  using ConvertOpToLLVMPattern<func::ReturnOp>::ConvertOpToLLVMPattern;
+ 
+   LogicalResult
+-  matchAndRewrite(ReturnOp op, OpAdaptor adaptor,
++  matchAndRewrite(func::ReturnOp op, OpAdaptor adaptor,
+                   ConversionPatternRewriter &rewriter) const override {
+     unsigned numArguments = op.getNumOperands();
+ 
+@@ -476,7 +476,6 @@ struct ExtractSliceOpConversion
+ 
+     auto llvmElemTy = getTypeConverter()->convertType(srcTy.getElementType());
+     auto elemPtrTy = ptr_ty(llvmElemTy, 3);
+-    auto resTy = op.getType().dyn_cast<RankedTensorType>();
+     smemObj = SharedMemoryObject(gep(elemPtrTy, smemObj.base, offset),
+                                  strideVals, offsetVals);
+     auto retVal = getStructFromSharedMemoryObject(loc, smemObj, rewriter);
+diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h
+index bb10d5b24a..00e399f848 100644
+--- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h
++++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h
+@@ -4,6 +4,7 @@
+ // TODO: refactor so that it doesn't fail if Allocation.h
+ // is included after utility.h (due to conflict in `store` macro
+ // and <atomic>
++#include "mlir/Dialect/Func/IR/FuncOps.h"
+ #include "triton/Analysis/Allocation.h"
+ 
+ //
+@@ -39,15 +40,15 @@ void vprintf_array(Value thread, ArrayRef<Value> arr, std::string info,
+ // TODO(Superjomn): remove the code when MLIR v15.0 is included.
+ // All the rights are reserved by the LLVM community.
+ 
+-struct FuncOpConversionBase : public ConvertOpToLLVMPattern<FuncOp> {
++struct FuncOpConversionBase : public ConvertOpToLLVMPattern<func::FuncOp> {
+ private:
+   /// Only retain those attributes that are not constructed by
+   /// `LLVMFuncOp::build`. If `filterArgAttrs` is set, also filter out argument
+   /// attributes.
+-  static void filterFuncAttributes(ArrayRef<NamedAttribute> attrs,
+-                                   bool filterArgAttrs,
++  static void filterFuncAttributes(func::FuncOp op, bool filterArgAttrs,
+                                    SmallVectorImpl<NamedAttribute> &result) {
+-    for (const auto &attr : attrs) {
++
++    for (const auto &attr : op->getAttrs()) {
+       if (attr.getName() == SymbolTable::getSymbolAttrName() ||
+           attr.getName() == FunctionOpInterface::getTypeAttrName() ||
+           attr.getName() == "std.varargs" ||
+@@ -65,27 +66,27 @@ struct FuncOpConversionBase : public ConvertOpToLLVMPattern<FuncOp> {
+   }
+ 
+ protected:
+-  using ConvertOpToLLVMPattern<FuncOp>::ConvertOpToLLVMPattern;
++  using ConvertOpToLLVMPattern<func::FuncOp>::ConvertOpToLLVMPattern;
+ 
+   // Convert input FuncOp to LLVMFuncOp by using the LLVMTypeConverter provided
+   // to this legalization pattern.
+   LLVM::LLVMFuncOp
+-  convertFuncOpToLLVMFuncOp(FuncOp funcOp,
++  convertFuncOpToLLVMFuncOp(func::FuncOp funcOp,
+                             ConversionPatternRewriter &rewriter) const {
+     // Convert the original function arguments. They are converted using the
+     // LLVMTypeConverter provided to this legalization pattern.
+     auto varargsAttr = funcOp->getAttrOfType<BoolAttr>("func.varargs");
+     TypeConverter::SignatureConversion result(funcOp.getNumArguments());
+     auto llvmType = getTypeConverter()->convertFunctionSignature(
+-        funcOp.getType(), varargsAttr && varargsAttr.getValue(), result);
++        funcOp.getFunctionType(), varargsAttr && varargsAttr.getValue(),
++        result);
+     if (!llvmType)
+       return nullptr;
+ 
+     // Propagate argument/result attributes to all converted arguments/result
+     // obtained after converting a given original argument/result.
+     SmallVector<NamedAttribute, 4> attributes;
+-    filterFuncAttributes(funcOp->getAttrs(), /*filterArgAttrs=*/true,
+-                         attributes);
++    filterFuncAttributes(funcOp, /*filterArgAttrs=*/true, attributes);
+     if (ArrayAttr resAttrDicts = funcOp.getAllResultAttrs()) {
+       assert(!resAttrDicts.empty() && "expected array to be non-empty");
+       auto newResAttrDicts =
+@@ -131,7 +132,7 @@ struct FuncOpConversionBase : public ConvertOpToLLVMPattern<FuncOp> {
+     }
+     auto newFuncOp = rewriter.create<LLVM::LLVMFuncOp>(
+         funcOp.getLoc(), funcOp.getName(), llvmType, linkage,
+-        /*dsoLocal*/ false, attributes);
++        /*dsoLocal*/ false, LLVM::CConv::C, attributes);
+     rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(),
+                                 newFuncOp.end());
+     if (failed(rewriter.convertRegionTypes(&newFuncOp.getBody(), *typeConverter,
+@@ -191,8 +192,8 @@ class ConvertTritonGPUOpToLLVMPatternBase {
+                                                const Allocation *allocation,
+                                                Value smem,
+                                                IndexCacheInfo indexCacheInfo)
+-      : converter(&typeConverter), indexCacheInfo(indexCacheInfo),
+-        allocation(allocation), smem(smem) {}
++      : converter(&typeConverter), allocation(allocation), smem(smem),
++        indexCacheInfo(indexCacheInfo) {}
+ 
+   LLVMTypeConverter *getTypeConverter() const { return converter; }
+ 
+@@ -861,7 +862,6 @@ class ConvertTritonGPUOpToLLVMPatternBase {
+                             ArrayRef<int64_t> shape) const {
+     auto parent = sliceLayout.getParent();
+     unsigned dim = sliceLayout.getDim();
+-    size_t rank = shape.size();
+     auto parentIndices =
+         emitIndices(loc, rewriter, parent, sliceLayout.paddedShape(shape));
+     unsigned numIndices = parentIndices.size();
+diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp
+index ff1af09835..6f66af4e34 100644
+--- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp
++++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp
+@@ -1,10 +1,11 @@
+ #include "triton/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.h"
+ 
++#include "mlir/Analysis/DataFlowFramework.h"
+ #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h"
++#include "mlir/Conversion/ControlFlowToLLVM//ControlFlowToLLVM.h"
+ #include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h"
+ #include "mlir/Conversion/MathToLLVM/MathToLLVM.h"
+-#include "mlir/Conversion/SCFToStandard/SCFToStandard.h"
+-#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
++#include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h"
+ #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+ #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
+ #include "mlir/Pass/Pass.h"
+@@ -40,7 +41,6 @@ class TritonLLVMConversionTarget : public ConversionTarget {
+     addIllegalDialect<triton::TritonDialect>();
+     addIllegalDialect<triton::gpu::TritonGPUDialect>();
+     addIllegalDialect<mlir::gpu::GPUDialect>();
+-    addIllegalDialect<mlir::StandardOpsDialect>();
+     addLegalOp<mlir::UnrealizedConversionCastOp>();
+   }
+ };
+@@ -51,7 +51,7 @@ class TritonLLVMFunctionConversionTarget : public ConversionTarget {
+       : ConversionTarget(ctx) {
+     addLegalDialect<LLVM::LLVMDialect>();
+     addLegalDialect<NVVM::NVVMDialect>();
+-    addIllegalOp<mlir::FuncOp>();
++    addIllegalOp<mlir::func::FuncOp>();
+     addLegalOp<mlir::UnrealizedConversionCastOp>();
+   }
+ };
+@@ -69,7 +69,7 @@ struct FuncOpConversion : public FuncOpConversionBase {
+       : FuncOpConversionBase(converter, benefit), numWarps(numWarps) {}
+ 
+   LogicalResult
+-  matchAndRewrite(FuncOp funcOp, OpAdaptor adaptor,
++  matchAndRewrite(func::FuncOp funcOp, OpAdaptor adaptor,
+                   ConversionPatternRewriter &rewriter) const override {
+     auto newFuncOp = convertFuncOpToLLVMFuncOp(funcOp, rewriter);
+     if (!newFuncOp)
+@@ -133,7 +133,8 @@ class ConvertTritonGPUToLLVM
+     decomposeBlockedToDotOperand(mod);
+ 
+     // Step 2
+-    decomposeInsertSliceAsyncOp(mod);
++    if (failed(decomposeInsertSliceAsyncOp(mod)))
++      return signalPassFailure();
+ 
+     // Step 3
+     Allocation allocation(mod);
+@@ -142,7 +143,7 @@ class ConvertTritonGPUToLLVM
+ 
+     // Step 4
+     RewritePatternSet scf_patterns(context);
+-    mlir::populateLoopToStdConversionPatterns(scf_patterns);
++    mlir::populateSCFToControlFlowConversionPatterns(scf_patterns);
+     mlir::ConversionTarget scf_target(*context);
+     scf_target.addIllegalOp<scf::ForOp, scf::IfOp, scf::ParallelOp,
+                             scf::WhileOp, scf::ExecuteRegionOp>();
+@@ -159,8 +160,10 @@ class ConvertTritonGPUToLLVM
+       return signalPassFailure();
+ 
+     // Step 6 - get axis and shared memory info
+-    AxisInfoAnalysis axisInfoAnalysis(mod.getContext());
+-    axisInfoAnalysis.run(mod);
++    std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver();
++    AxisInfoAnalysis *axisInfoAnalysis = solver->load<AxisInfoAnalysis>();
++    if (failed(solver->initializeAndRun(mod)))
++      return signalPassFailure();
+     initSharedMemory(allocation.getSharedMemorySize(), typeConverter);
+     mod->setAttr("triton_gpu.shared",
+                  mlir::IntegerAttr::get(mlir::IntegerType::get(context, 32),
+@@ -178,38 +181,39 @@ class ConvertTritonGPUToLLVM
+ 
+     // Normal conversions
+     populateTritonGPUToLLVMPatterns(typeConverter, patterns, numWarps,
+-                                    axisInfoAnalysis, &allocation, smem,
++                                    *axisInfoAnalysis, &allocation, smem,
+                                     indexCacheInfo, /*benefit=*/10);
+     // ConvertLayoutOp
+     populateConvertLayoutOpToLLVMPatterns(typeConverter, patterns, numWarps,
+-                                          axisInfoAnalysis, &allocation, smem,
++                                          *axisInfoAnalysis, &allocation, smem,
+                                           indexCacheInfo, /*benefit=*/10);
+     // DotOp
+     populateDotOpToLLVMPatterns(typeConverter, patterns, numWarps,
+-                                axisInfoAnalysis, &allocation, smem,
++                                *axisInfoAnalysis, &allocation, smem,
+                                 /*benefit=*/10);
+     // ElementwiseOp
+     populateElementwiseOpToLLVMPatterns(typeConverter, patterns, numWarps,
+-                                        axisInfoAnalysis, &allocation, smem,
++                                        *axisInfoAnalysis, &allocation, smem,
+                                         /*benefit=*/10);
+     // LoadStoreOp
+     populateLoadStoreOpToLLVMPatterns(typeConverter, patterns, numWarps,
+-                                      axisInfoAnalysis, &allocation, smem,
++                                      *axisInfoAnalysis, &allocation, smem,
+                                       indexCacheInfo, /*benefit=*/10);
+     // ReduceOp
+     populateReduceOpToLLVMPatterns(typeConverter, patterns, numWarps,
+-                                   axisInfoAnalysis, &allocation, smem,
++                                   *axisInfoAnalysis, &allocation, smem,
+                                    indexCacheInfo, /*benefit=*/10);
+     // ViewOp
+     populateViewOpToLLVMPatterns(typeConverter, patterns, numWarps,
+-                                 axisInfoAnalysis, &allocation, smem,
++                                 *axisInfoAnalysis, &allocation, smem,
+                                  /*benefit=*/10);
+ 
+     // Add arith/math's patterns to help convert scalar expression to LLVM.
+     mlir::arith::populateArithmeticToLLVMConversionPatterns(typeConverter,
+                                                             patterns);
+     mlir::populateMathToLLVMConversionPatterns(typeConverter, patterns);
+-    mlir::populateStdToLLVMConversionPatterns(typeConverter, patterns);
++    mlir::cf::populateControlFlowToLLVMConversionPatterns(typeConverter,
++                                                          patterns);
+     mlir::populateGpuToNVVMConversionPatterns(typeConverter, patterns);
+ 
+     if (failed(applyPartialConversion(mod, target, std::move(patterns))))
+@@ -306,9 +310,11 @@ class ConvertTritonGPUToLLVM
+     });
+   }
+ 
+-  void decomposeInsertSliceAsyncOp(ModuleOp mod) const {
+-    AxisInfoAnalysis axisInfoAnalysis(mod.getContext());
+-    axisInfoAnalysis.run(mod);
++  LogicalResult decomposeInsertSliceAsyncOp(ModuleOp mod) const {
++    std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver();
++    AxisInfoAnalysis *axisInfoAnalysis = solver->load<AxisInfoAnalysis>();
++    if (failed(solver->initializeAndRun(mod)))
++      return failure();
+     // TODO(Keren): This is a hacky knob that may cause performance regression
+     // when decomposition has been performed. We should remove this knob once we
+     // have thorough analysis on async wait. Currently, we decompose
+@@ -342,7 +348,7 @@ class ConvertTritonGPUToLLVM
+       auto resSharedLayout =
+           dstTy.getEncoding().dyn_cast<triton::gpu::SharedEncodingAttr>();
+       auto resElemTy = dstTy.getElementType();
+-      unsigned inVec = axisInfoAnalysis.getPtrContiguity(src);
++      unsigned inVec = axisInfoAnalysis->getPtrContiguity(src);
+       unsigned outVec = resSharedLayout.getVec();
+       unsigned minVec = std::min(outVec, inVec);
+       auto maxBitWidth =
+@@ -400,11 +406,11 @@ class ConvertTritonGPUToLLVM
+       } else if (decomposed) {
+         // Wait for all previous async ops
+         OpBuilder builder(asyncWaitOp);
+-        auto newAsyncWaitOp =
+-            builder.create<triton::gpu::AsyncWaitOp>(asyncWaitOp.getLoc(), 0);
++        builder.create<triton::gpu::AsyncWaitOp>(asyncWaitOp.getLoc(), 0);
+         asyncWaitOp.erase();
+       }
+     });
++    return success();
+   }
+ };
+ 
+diff --git a/lib/Conversion/TritonGPUToLLVM/Utility.h b/lib/Conversion/TritonGPUToLLVM/Utility.h
+index d35dac28c5..11976908cf 100644
+--- a/lib/Conversion/TritonGPUToLLVM/Utility.h
++++ b/lib/Conversion/TritonGPUToLLVM/Utility.h
+@@ -220,10 +220,7 @@ struct SharedMemoryObject {
+                      ConversionPatternRewriter &rewriter)
+       : base(base) {
+     strides = getStridesFromShapeAndOrder(shape, order, loc, rewriter);
+-
+-    for (auto idx : order) {
+-      offsets.emplace_back(i32_val(0));
+-    }
++    offsets.append(order.size(), i32_val(0));
+   }
+ 
+   SmallVector<Value> getElems() const {
+diff --git a/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp b/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp
+index fe42202c34..5f230f787f 100644
+--- a/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp
++++ b/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp
+@@ -1,10 +1,10 @@
+ #include "triton/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.h"
+ 
+ #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
+-#include "mlir/Dialect/GPU/GPUDialect.h"
++#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
++#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+ #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+ #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
+-#include "mlir/Dialect/StandardOps/IR/Ops.h"
+ #include "mlir/Pass/Pass.h"
+ #include "mlir/Transforms/DialectConversion.h"
+ #include "triton/Dialect/Triton/IR/Dialect.h"
+@@ -59,10 +59,13 @@ class ArithConstantPattern : public OpConversionPattern<arith::ConstantOp> {
+     Type retType = getTypeConverter()->convertType(op.getType());
+     auto value = adaptor.getValue().dyn_cast<DenseElementsAttr>();
+     assert(value);
+-    rewriter.replaceOpWithNewOp<arith::ConstantOp>(
+-        op, retType,
+-        value.reshape(retType) // This is a hack. We just want to add encoding
+-    );
++    if (value.getElementType().isInteger(1) && value.isSplat())
++      // Workaround until https://reviews.llvm.org/D133743 is included.
++      value = DenseElementsAttr::get(retType, value.getSplatValue<bool>());
++    else
++      // This is a hack. We just want to add encoding
++      value = value.reshape(retType);
++    rewriter.replaceOpWithNewOp<arith::ConstantOp>(op, retType, value);
+     return success();
+   }
+ };
+@@ -127,12 +130,12 @@ void populateArithmeticPatternsAndLegality(
+ }
+ 
+ // this shouldn't exist if mlir's SelectOp checked encodings properly
+-class StdSelectPattern : public OpConversionPattern<SelectOp> {
++class StdSelectPattern : public OpConversionPattern<arith::SelectOp> {
+ public:
+-  using OpConversionPattern<SelectOp>::OpConversionPattern;
++  using OpConversionPattern<arith::SelectOp>::OpConversionPattern;
+ 
+   LogicalResult
+-  matchAndRewrite(SelectOp op, typename SelectOp::Adaptor adaptor,
++  matchAndRewrite(arith::SelectOp op, OpAdaptor adaptor,
+                   ConversionPatternRewriter &rewriter) const override {
+     Type retType = this->getTypeConverter()->convertType(op.getType());
+     rewriter.replaceOpWithNewOp<triton::gpu::SelectOp>(
+@@ -148,8 +151,8 @@ void populateStdPatternsAndLegality(TritonGPUTypeConverter &typeConverter,
+   MLIRContext *context = patterns.getContext();
+   // Rewrite rule
+   patterns.add<StdSelectPattern>(typeConverter, context);
+-  target.addLegalOp<ReturnOp>(); // this is ok because all functions are inlined
+-                                 // by the frontend
++  target.addLegalOp<func::ReturnOp>(); // this is ok because all functions are
++                                       // inlined by the frontend
+ }
+ 
+ void populateMathPatternsAndLegality(TritonGPUTypeConverter &typeConverter,
+@@ -455,18 +458,19 @@ struct TritonPrintfPattern : public OpConversionPattern<triton::PrintfOp> {
+ void populateTritonPatterns(TritonGPUTypeConverter &typeConverter,
+                             RewritePatternSet &patterns) {
+   MLIRContext *context = patterns.getContext();
+-  patterns.add< // TODO: view should have custom pattern that views the layout
+-      TritonGenericPattern<triton::ViewOp>,
+-      TritonGenericPattern<triton::BitcastOp>,
+-      TritonGenericPattern<triton::FpToFpOp>,
+-      TritonGenericPattern<triton::IntToPtrOp>,
+-      TritonGenericPattern<triton::PtrToIntOp>,
+-      TritonGenericPattern<triton::SplatOp>, TritonBroadcastPattern,
+-      TritonGenericPattern<triton::AddPtrOp>, TritonCatPattern,
+-      TritonReducePattern, TritonTransPattern, TritonExpandDimsPattern,
+-      TritonMakeRangePattern, TritonDotPattern, TritonLoadPattern,
+-      TritonStorePattern, TritonExtElemwisePattern, TritonPrintfPattern,
+-      TritonAtomicRMWPattern>(typeConverter, context);
++  patterns
++      .insert< // TODO: view should have custom pattern that views the layout
++          TritonGenericPattern<triton::ViewOp>,
++          TritonGenericPattern<triton::BitcastOp>,
++          TritonGenericPattern<triton::FpToFpOp>,
++          TritonGenericPattern<triton::IntToPtrOp>,
++          TritonGenericPattern<triton::PtrToIntOp>,
++          TritonGenericPattern<triton::SplatOp>, TritonBroadcastPattern,
++          TritonGenericPattern<triton::AddPtrOp>, TritonCatPattern,
++          TritonReducePattern, TritonTransPattern, TritonExpandDimsPattern,
++          TritonMakeRangePattern, TritonDotPattern, TritonLoadPattern,
++          TritonStorePattern, TritonExtElemwisePattern, TritonPrintfPattern,
++          TritonAtomicRMWPattern>(typeConverter, context);
+ }
+ 
+ //
+@@ -623,29 +627,28 @@ void populateSCFPatterns(TritonGPUTypeConverter &typeConverter,
+ 
+ // CF
+ 
+-class CFBranchPattern : public OpConversionPattern<BranchOp> {
++class CFBranchPattern : public OpConversionPattern<cf::BranchOp> {
+ public:
+-  using OpConversionPattern<BranchOp>::OpConversionPattern;
++  using OpConversionPattern<cf::BranchOp>::OpConversionPattern;
+ 
+   LogicalResult
+-  matchAndRewrite(BranchOp op, BranchOp::Adaptor adaptor,
++  matchAndRewrite(cf::BranchOp op, cf::BranchOp::Adaptor adaptor,
+                   ConversionPatternRewriter &rewriter) const override {
+-    auto converter = getTypeConverter();
+-    auto newOp = rewriter.replaceOpWithNewOp<BranchOp>(op, op.getSuccessor(),
+-                                                       adaptor.getOperands());
++    auto newOp = rewriter.replaceOpWithNewOp<cf::BranchOp>(
++        op, op.getSuccessor(), adaptor.getOperands());
+     return success();
+   }
+ };
+ 
+-class CFCondBranchPattern : public OpConversionPattern<CondBranchOp> {
++class CFCondBranchPattern : public OpConversionPattern<cf::CondBranchOp> {
+ public:
+-  using OpConversionPattern<CondBranchOp>::OpConversionPattern;
++  using OpConversionPattern<cf::CondBranchOp>::OpConversionPattern;
+ 
+   LogicalResult
+-  matchAndRewrite(CondBranchOp op, CondBranchOp::Adaptor adaptor,
++  matchAndRewrite(cf::CondBranchOp op, cf::CondBranchOp::Adaptor adaptor,
+                   ConversionPatternRewriter &rewriter) const override {
+     auto converter = getTypeConverter();
+-    auto newOp = rewriter.replaceOpWithNewOp<CondBranchOp>(
++    auto newOp = rewriter.replaceOpWithNewOp<cf::CondBranchOp>(
+         op, adaptor.getCondition(), op.getTrueDest(),
+         adaptor.getTrueDestOperands(), op.getFalseDest(),
+         adaptor.getFalseDestOperands());
+diff --git a/lib/Dialect/Triton/IR/CMakeLists.txt b/lib/Dialect/Triton/IR/CMakeLists.txt
+index 2d679b21fd..705554ba6b 100644
+--- a/lib/Dialect/Triton/IR/CMakeLists.txt
++++ b/lib/Dialect/Triton/IR/CMakeLists.txt
+@@ -10,11 +10,7 @@ add_mlir_dialect_library(TritonIR
+ 
+   LINK_LIBS PUBLIC
+   MLIRIR
+-  MLIRArithmetic
+-  MLIRSCF
+-
+-  # Since LLVM 15
+-  # MLIRFunc
+-  # else
+-  MLIRStandard
++  MLIRArithmeticDialect
++  MLIRSCFDialect
++  MLIRFuncDialect
+ )
+diff --git a/lib/Dialect/Triton/IR/Ops.cpp b/lib/Dialect/Triton/IR/Ops.cpp
+index 3aadbfa0c0..86570359c5 100644
+--- a/lib/Dialect/Triton/IR/Ops.cpp
++++ b/lib/Dialect/Triton/IR/Ops.cpp
+@@ -1,10 +1,9 @@
+-#include "triton/Dialect/Triton/IR/Dialect.h"
+-#include "triton/Dialect/Triton/IR/Types.h"
+-
+ #include "mlir/IR/Builders.h"
+ #include "mlir/IR/BuiltinAttributes.h"
+ #include "mlir/IR/BuiltinTypes.h"
+ #include "mlir/IR/OperationSupport.h"
++#include "triton/Dialect/Triton/IR/Dialect.h"
++#include "triton/Dialect/Triton/IR/Types.h"
+ 
+ namespace mlir {
+ namespace triton {
+@@ -38,8 +37,8 @@ static Type getPointerTypeSameShape(Type type) {
+ }
+ 
+ // Parser & printer for assembly forms
+-ParseResult parseLoadOp(OpAsmParser &parser, OperationState &result) {
+-  SmallVector<OpAsmParser::OperandType, 4> allOperands;
++ParseResult LoadOp::parse(OpAsmParser &parser, OperationState &result) {
++  SmallVector<OpAsmParser::UnresolvedOperand, 4> allOperands;
+   Type resultTypes[1];
+   SMLoc allOperandLoc = parser.getCurrentLocation();
+   if (parser.parseOperandList(allOperands) ||
+@@ -73,18 +72,18 @@ ParseResult parseLoadOp(OpAsmParser &parser, OperationState &result) {
+   return success();
+ }
+ 
+-void printLoadOp(OpAsmPrinter &printer, LoadOp loadOp) {
++void LoadOp::print(OpAsmPrinter &printer) {
+   printer << " ";
+-  printer << loadOp.getOperation()->getOperands();
++  printer << getOperation()->getOperands();
+   // "operand_segment_sizes" can be deduced, so we don't print it.
+-  printer.printOptionalAttrDict(loadOp->getAttrs(),
+-                                {loadOp.operand_segment_sizesAttrName()});
++  printer.printOptionalAttrDict(getOperation()->getAttrs(),
++                                {operand_segment_sizesAttrName()});
+   printer << " : ";
+-  printer.printStrippedAttrOrType(loadOp.result().getType());
++  printer.printStrippedAttrOrType(getResult().getType());
+ }
+ 
+-ParseResult parseStoreOp(OpAsmParser &parser, OperationState &result) {
+-  SmallVector<OpAsmParser::OperandType, 4> allOperands;
++ParseResult StoreOp::parse(OpAsmParser &parser, OperationState &result) {
++  SmallVector<OpAsmParser::UnresolvedOperand, 4> allOperands;
+   Type valueType;
+   SMLoc allOperandLoc = parser.getCurrentLocation();
+   if (parser.parseOperandList(allOperands) ||
+@@ -104,12 +103,12 @@ ParseResult parseStoreOp(OpAsmParser &parser, OperationState &result) {
+   return success();
+ }
+ 
+-void printStoreOp(OpAsmPrinter &printer, StoreOp storeOp) {
++void StoreOp::print(OpAsmPrinter &printer) {
+   printer << " ";
+-  printer << storeOp.getOperation()->getOperands();
+-  printer.printOptionalAttrDict(storeOp->getAttrs(), /*elidedAttrs=*/{});
++  printer << getOperation()->getOperands();
++  printer.printOptionalAttrDict(getOperation()->getAttrs(), /*elidedAttrs=*/{});
+   printer << " : ";
+-  printer.printStrippedAttrOrType(storeOp.value().getType());
++  printer.printStrippedAttrOrType(value().getType());
+ }
+ 
+ } // namespace triton
+@@ -319,7 +318,8 @@ OpFoldResult SplatOp::fold(ArrayRef<Attribute> operands) {
+   if (!constOperand)
+     return {};
+   auto shapedType = getType().cast<ShapedType>();
+-  auto ret = SplatElementsAttr::get(shapedType, {constOperand.getValue()});
++  auto ret = SplatElementsAttr::get(
++      shapedType, ArrayRef<Attribute>(constOperand.getValue()));
+   return ret;
+ }
+ 
+diff --git a/lib/Dialect/Triton/Transforms/Combine.cpp b/lib/Dialect/Triton/Transforms/Combine.cpp
+index 2261472170..11570283d6 100644
+--- a/lib/Dialect/Triton/Transforms/Combine.cpp
++++ b/lib/Dialect/Triton/Transforms/Combine.cpp
+@@ -57,13 +57,13 @@ DenseElementsAttr getConstantValue(Builder &builder, Attribute value,
+ class CombineSelectMaskedLoadPattern : public mlir::RewritePattern {
+ public:
+   CombineSelectMaskedLoadPattern(mlir::MLIRContext *context)
+-      : mlir::RewritePattern(mlir::SelectOp::getOperationName(), 3, context,
+-                             {triton::LoadOp::getOperationName()}) {}
++      : mlir::RewritePattern(mlir::arith::SelectOp::getOperationName(), 3,
++                             context, {triton::LoadOp::getOperationName()}) {}
+ 
+   mlir::LogicalResult
+   matchAndRewrite(mlir::Operation *op,
+                   mlir::PatternRewriter &rewriter) const override {
+-    auto selectOp = llvm::dyn_cast<mlir::SelectOp>(op);
++    auto selectOp = llvm::dyn_cast<mlir::arith::SelectOp>(op);
+     if (!selectOp)
+       return mlir::failure();
+ 
+diff --git a/lib/Dialect/Triton/Transforms/Combine.td b/lib/Dialect/Triton/Transforms/Combine.td
+index 14f286b26e..ded0e346e6 100644
+--- a/lib/Dialect/Triton/Transforms/Combine.td
++++ b/lib/Dialect/Triton/Transforms/Combine.td
+@@ -1,9 +1,9 @@
+ #ifndef TRITON_PATTERNS
+ #define TRITON_PATTERNS
+ 
+-include "mlir/Dialect/StandardOps/IR/Ops.td"
+ include "mlir/Dialect/Arithmetic/IR/ArithmeticOps.td"
+ include "triton/Dialect/Triton/IR/TritonOps.td"
++include "mlir/IR/PatternBase.td"
+ 
+ 
+ // AddIOp(DotOp(a, b, c), d) and c==0 => DotOp(a, b, d)
+diff --git a/lib/Dialect/TritonGPU/IR/Dialect.cpp b/lib/Dialect/TritonGPU/IR/Dialect.cpp
+index 1fbc609e88..bfc3f3d3da 100644
+--- a/lib/Dialect/TritonGPU/IR/Dialect.cpp
++++ b/lib/Dialect/TritonGPU/IR/Dialect.cpp
+@@ -1,14 +1,14 @@
++#include "triton/Dialect/Triton/IR/Dialect.h"
++
+ #include <numeric>
+ 
+ #include "mlir/IR/DialectImplementation.h"
+ #include "mlir/IR/OpImplementation.h"
+ #include "triton/Analysis/Utility.h"
+-#include "triton/Dialect/Triton/IR/Dialect.h"
++#include "triton/Dialect/TritonGPU/IR/Dialect.cpp.inc"
+ #include "triton/Dialect/TritonGPU/IR/Dialect.h"
+ #include "llvm/ADT/TypeSwitch.h"
+ 
+-#include "triton/Dialect/TritonGPU/IR/Dialect.cpp.inc"
+-
+ using namespace mlir;
+ using namespace mlir::triton::gpu;
+ 
+@@ -366,7 +366,6 @@ template SmallVector<int64_t>
+ SliceEncodingAttr::paddedShape<int64_t>(ArrayRef<int64_t> shape) const;
+ 
+ unsigned SliceEncodingAttr::getElemsPerThread(ArrayRef<int64_t> shape) const {
+-  size_t rank = shape.size();
+   auto parent = getParent();
+   return ::getElemsPerThread(parent, paddedShape(shape));
+ }
+@@ -655,9 +654,9 @@ void DotOperandEncodingAttr::print(mlir::AsmPrinter &printer) const {
+ // InsertSliceAsyncOp
+ //===----------------------------------------------------------------------===//
+ 
+-ParseResult parseInsertSliceAsyncOp(OpAsmParser &parser,
+-                                    OperationState &result) {
+-  SmallVector<OpAsmParser::OperandType, 8> allOperands;
++ParseResult InsertSliceAsyncOp::parse(OpAsmParser &parser,
++                                      OperationState &result) {
++  SmallVector<OpAsmParser::UnresolvedOperand, 8> allOperands;
+   Type srcType, dstType;
+   SMLoc allOperandLoc = parser.getCurrentLocation();
+   if (parser.parseOperandList(allOperands) ||
+@@ -696,18 +695,16 @@ ParseResult parseInsertSliceAsyncOp(OpAsmParser &parser,
+   return success();
+ }
+ 
+-void printInsertSliceAsyncOp(OpAsmPrinter &printer,
+-                             InsertSliceAsyncOp insertSliceAsyncOp) {
++void InsertSliceAsyncOp::print(OpAsmPrinter &printer) {
+   printer << " ";
+-  printer << insertSliceAsyncOp.getOperation()->getOperands();
++  printer << getOperation()->getOperands();
+   // "operand_segment_sizes" can be deduced, so we don't print it.
+-  printer.printOptionalAttrDict(
+-      insertSliceAsyncOp->getAttrs(),
+-      {insertSliceAsyncOp.operand_segment_sizesAttrName()});
++  printer.printOptionalAttrDict(getOperation()->getAttrs(),
++                                {operand_segment_sizesAttrName()});
+   printer << " : ";
+-  printer.printStrippedAttrOrType(insertSliceAsyncOp.src().getType());
++  printer.printStrippedAttrOrType(src().getType());
+   printer << " -> ";
+-  printer.printStrippedAttrOrType(insertSliceAsyncOp.result().getType());
++  printer.printStrippedAttrOrType(result().getType());
+ }
+ 
+ //===----------------------------------------------------------------------===//
+diff --git a/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp b/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp
+index 82407980d3..ee6009f44a 100644
+--- a/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp
++++ b/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp
+@@ -27,7 +27,11 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> {
+     auto origType = ptr.getType().cast<RankedTensorType>();
+     // Get the shape of the tensor.
+     size_t rank = origType.getRank();
+-    AxisInfo info = axisInfo.lookupLatticeElement(ptr)->getValue();
++    dataflow::Lattice<AxisInfo> *latticeElement =
++        axisInfo.getLatticeElement(ptr);
++    AxisInfo info = latticeElement && !latticeElement->isUninitialized()
++                        ? latticeElement->getValue()
++                        : AxisInfo();
+     // Get the contiguity order of `ptr`
+     auto order = argSort(info.getContiguity());
+     // The desired divisibility is the maximum divisibility
+@@ -40,7 +44,7 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> {
+         for (Value val : op->getResults()) {
+           if (val.getType() != origType)
+             continue;
+-          auto valInfo = axisInfo.lookupLatticeElement(val);
++          auto valInfo = axisInfo.getLatticeElement(val);
+           auto currOrder = argSort(valInfo->getValue().getContiguity());
+           if (order == currOrder)
+             withSameOrder.insert(val);
+@@ -55,7 +59,7 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> {
+     unsigned elemNumBytes = std::max(elemNumBits / 8, 1u);
+     unsigned perThread = 1;
+     for (Value val : withSameOrder) {
+-      AxisInfo info = axisInfo.lookupLatticeElement(val)->getValue();
++      AxisInfo info = axisInfo.getLatticeElement(val)->getValue();
+       unsigned maxMultipleBytes = info.getDivisibility(order[0]);
+       unsigned maxMultiple = std::max(maxMultipleBytes / elemNumBytes, 1u);
+       unsigned maxContig = info.getContiguity(order[0]);
+@@ -123,8 +127,10 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> {
+   void runOnOperation() override {
+     Operation *op = getOperation();
+     // Run axis info analysis
+-    AxisInfoAnalysis axisInfo(&getContext());
+-    axisInfo.run(op);
++    std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver();
++    AxisInfoAnalysis *axisInfo = solver->load<AxisInfoAnalysis>();
++    if (failed(solver->initializeAndRun(op)))
++      return signalPassFailure();
+ 
+     // For each i/o operation, we determine what layout
+     // the pointers should have for best memory coalescing
+@@ -146,10 +152,10 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> {
+       RankedTensorType ty = ptr.getType().template dyn_cast<RankedTensorType>();
+       if (!ty || !ty.getElementType().isa<PointerType>())
+         return;
+-      AxisInfo info = axisInfo.lookupLatticeElement(ptr)->getValue();
++      AxisInfo info = axisInfo->getLatticeElement(ptr)->getValue();
+       auto mod = curr->getParentOfType<ModuleOp>();
+       int numWarps = triton::gpu::TritonGPUDialect::getNumWarps(mod);
+-      auto convertType = getTypeConverter(axisInfo, ptr, numWarps);
++      auto convertType = getTypeConverter(*axisInfo, ptr, numWarps);
+       layoutMap[ptr] = convertType;
+     });
+ 
+diff --git a/lib/Dialect/TritonGPU/Transforms/Combine.cpp b/lib/Dialect/TritonGPU/Transforms/Combine.cpp
+index efa37ff2dc..089ce3996c 100644
+--- a/lib/Dialect/TritonGPU/Transforms/Combine.cpp
++++ b/lib/Dialect/TritonGPU/Transforms/Combine.cpp
+@@ -1,6 +1,6 @@
+ #include "Utility.h"
+ #include "mlir/Analysis/SliceAnalysis.h"
+-#include "mlir/Dialect/SCF/SCF.h"
++#include "mlir/Dialect/SCF/IR/SCF.h"
+ #include "mlir/IR/BlockAndValueMapping.h"
+ #include "mlir/IR/BuiltinAttributes.h"
+ #include "mlir/IR/Matchers.h"
+diff --git a/lib/Dialect/TritonGPU/Transforms/Combine.td b/lib/Dialect/TritonGPU/Transforms/Combine.td
+index 6bf1b14866..6a7b10dbcb 100644
+--- a/lib/Dialect/TritonGPU/Transforms/Combine.td
++++ b/lib/Dialect/TritonGPU/Transforms/Combine.td
+@@ -3,5 +3,6 @@
+ 
+ include "triton/Dialect/TritonGPU/IR/TritonGPUOps.td"
+ include "triton/Dialect/Triton/IR/TritonOps.td"
++include "mlir/IR/PatternBase.td"
+ 
+ #endif
+diff --git a/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp b/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp
+index 4bd3bc76bf..b2f8defd81 100644
+--- a/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp
++++ b/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp
+@@ -1,5 +1,5 @@
+ #include "mlir/Analysis/SliceAnalysis.h"
+-#include "mlir/Dialect/SCF/SCF.h"
++#include "mlir/Dialect/SCF/IR/SCF.h"
+ #include "mlir/IR/BlockAndValueMapping.h"
+ #include "mlir/IR/BuiltinAttributes.h"
+ #include "mlir/IR/Matchers.h"
+diff --git a/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp b/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp
+index 9b2f42231e..85f746c1dc 100644
+--- a/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp
++++ b/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp
+@@ -2,6 +2,7 @@
+ #include "mlir/IR/BlockAndValueMapping.h"
+ #include "mlir/IR/TypeUtilities.h"
+ #include "triton/Analysis/AxisInfo.h"
++#include "triton/Analysis/Utility.h"
+ #include "triton/Dialect/TritonGPU/IR/Dialect.h"
+ #include "triton/Dialect/TritonGPU/Transforms/Passes.h"
+ 
+@@ -160,15 +161,18 @@ ttg::AllocTensorOp LoopPipeliner::allocateEmptyBuffer(Operation *op,
+ LogicalResult LoopPipeliner::initialize() {
+   Block *loop = forOp.getBody();
+ 
+-  AxisInfoAnalysis axisInfoAnalysis(forOp.getContext());
+-  axisInfoAnalysis.run(forOp->getParentOfType<ModuleOp>());
++  std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver();
++  AxisInfoAnalysis *axisInfoAnalysis = solver->load<AxisInfoAnalysis>();
++  if (failed(solver->initializeAndRun(forOp->getParentOfType<ModuleOp>()))) {
++    return failure();
++  }
+ 
+   // can we use forOp.walk(...) here?
+   SmallVector<triton::LoadOp, 2> allLoads;
+   for (Operation &op : *loop)
+     if (auto loadOp = dyn_cast<triton::LoadOp>(&op)) {
+       auto ptr = loadOp.ptr();
+-      unsigned vec = axisInfoAnalysis.getPtrContiguity(ptr);
++      unsigned vec = axisInfoAnalysis->getPtrContiguity(ptr);
+       auto tensorTy = ptr.getType().dyn_cast<RankedTensorType>();
+       if (!tensorTy)
+         continue;
+diff --git a/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp b/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp
+index 0e7dbe5264..b95a4f50a6 100644
+--- a/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp
++++ b/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp
+@@ -1,5 +1,5 @@
+ #include "mlir/Analysis/SliceAnalysis.h"
+-#include "mlir/Dialect/SCF/SCF.h"
++#include "mlir/Dialect/SCF/IR/SCF.h"
+ #include "mlir/IR/BlockAndValueMapping.h"
+ #include "mlir/IR/BuiltinAttributes.h"
+ #include "mlir/IR/Matchers.h"
+diff --git a/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp b/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp
+index 37ac710995..762e887f36 100644
+--- a/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp
++++ b/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp
+@@ -82,12 +82,12 @@ TritonGPUConversionTarget::TritonGPUConversionTarget(
+                scf::ReduceReturnOp>();
+ 
+   addDynamicallyLegalDialect<arith::ArithmeticDialect, math::MathDialect,
+-                             triton::TritonDialect, StandardOpsDialect,
+-                             scf::SCFDialect>([&](Operation *op) {
+-    if (typeConverter.isLegal(op))
+-      return true;
+-    return false;
+-  });
++                             triton::TritonDialect, scf::SCFDialect>(
++      [&](Operation *op) {
++        if (typeConverter.isLegal(op))
++          return true;
++        return false;
++      });
+ 
+   // We have requirements for the data layouts
+   addDynamicallyLegalOp<triton::DotOp>([](triton::DotOp dotOp) -> bool {
+diff --git a/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp b/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp
+index c229104286..c911fd4a5c 100644
+--- a/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp
++++ b/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp
+@@ -1,5 +1,5 @@
+ #include "Utility.h"
+-#include "mlir/Dialect/SCF/SCF.h"
++#include "mlir/Dialect/SCF/IR/SCF.h"
+ #include "mlir/IR/Matchers.h"
+ #include "mlir/IR/PatternMatch.h"
+ #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
+@@ -118,8 +118,8 @@ void setOpResultType(Operation *op, ArrayRef<Type> newTypes) {
+                     .get("value")
+                     .dyn_cast<mlir::DenseElementsAttr>();
+     if (attr) {
+-      auto newAttr = mlir::DenseElementsAttr::getFromRawBuffer(
+-          newType, attr.getRawData(), true);
++      auto newAttr =
++          mlir::DenseElementsAttr::getFromRawBuffer(newType, attr.getRawData());
+       op->setAttr("value", newAttr);
+     }
+   }
+diff --git a/lib/Dialect/TritonGPU/Transforms/Utility.cpp b/lib/Dialect/TritonGPU/Transforms/Utility.cpp
+index ed15f02f67..6400f1633a 100644
+--- a/lib/Dialect/TritonGPU/Transforms/Utility.cpp
++++ b/lib/Dialect/TritonGPU/Transforms/Utility.cpp
+@@ -1,5 +1,5 @@
+ #include "Utility.h"
+-#include "mlir/Dialect/SCF/SCF.h"
++#include "mlir/Dialect/SCF/IR/SCF.h"
+ #include "mlir/IR/BlockAndValueMapping.h"
+ #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
+ 
+diff --git a/lib/Target/LLVMIR/CMakeLists.txt b/lib/Target/LLVMIR/CMakeLists.txt
+index f1bbd0bf4e..ac8973ad19 100644
+--- a/lib/Target/LLVMIR/CMakeLists.txt
++++ b/lib/Target/LLVMIR/CMakeLists.txt
+@@ -6,8 +6,7 @@ add_mlir_translation_library(TritonLLVMIR
+ 
+         LINK_LIBS PUBLIC
+         MLIRIR
+-        MLIRLLVMIR
+-        MLIRSCFToStandard
++        MLIRLLVMDialect
+         MLIRSupport
+         MLIRTargetLLVMIRExport
+         )
+diff --git a/lib/Target/PTX/PTXTranslation.cpp b/lib/Target/PTX/PTXTranslation.cpp
+index 4cb0d8193c..6a5453a6e7 100644
+--- a/lib/Target/PTX/PTXTranslation.cpp
++++ b/lib/Target/PTX/PTXTranslation.cpp
+@@ -1,11 +1,14 @@
+ #include "triton/Target/PTX/PTXTranslation.h"
+ #include "triton/Target/LLVMIR/LLVMIRTranslation.h"
++#include <optional>
+ 
+ #include "llvm/IR/IRBuilder.h"
+ #include "llvm/IR/LegacyPassManager.h"
+ #include "llvm/IR/Module.h"
+ #include "llvm/IR/Verifier.h"
+ #include "llvm/MC/TargetRegistry.h"
++#include "llvm/Pass.h"
++#include "llvm/Support/CommandLine.h"
+ #include "llvm/Support/TargetSelect.h"
+ #include "llvm/Target/TargetMachine.h"
+ 
+diff --git a/python/setup.py b/python/setup.py
+index 2ac3accd25..4530b36714 100644
+--- a/python/setup.py
++++ b/python/setup.py
+@@ -57,19 +57,10 @@ def get_pybind11_package_info():
+ def get_llvm_package_info():
+     # download if nothing is installed
+     system = platform.system()
+-    if system == "Darwin":
+-        system_suffix = "apple-darwin"
+-    elif system == "Linux":
+-        vglibc = tuple(map(int, platform.libc_ver()[1].split('.')))
+-        vglibc = vglibc[0] * 100 + vglibc[1]
+-        linux_suffix = 'ubuntu-18.04' if vglibc > 217 else 'centos-7'
+-        system_suffix = f"linux-gnu-{linux_suffix}"
+-    else:
+-        raise RuntimeError(f"unsupported system: {system}")
++    system_suffix = {"Linux": "linux-gnu-ubuntu-18.04", "Darwin": "apple-darwin"}[system]
+     use_assert_enabled_llvm = check_env_flag("TRITON_USE_ASSERT_ENABLED_LLVM", "False")
+-    release_suffix = "assert" if use_assert_enabled_llvm else "release"
+-    name = f'llvm+mlir-14.0.6-x86_64-{system_suffix}-{release_suffix}'
+-    url = f"https://github.com/ptillet/triton-llvm-releases/releases/download/llvm-14.0.6-f28c006a5895/{name}.tar.xz"
++    name = 'llvm+mlir-15.0.7-x86_64-{}-{}'.format(system_suffix, "assert" if use_assert_enabled_llvm else "release")
++    url = "https://github.com/ptillet/triton-llvm-releases/releases/download/llvm-15.0.7-8dfdcc7b7bf6/{}.tar.xz".format(name)
+     return Package("llvm", name, url, "lib", "LLVM_INCLUDE_DIRS", "LLVM_LIBRARY_DIR", "LLVM_SYSPATH")
+ 
+ 
+diff --git a/python/src/triton.cc b/python/src/triton.cc
+index c40b117a55..f190eacc34 100644
+--- a/python/src/triton.cc
++++ b/python/src/triton.cc
+@@ -8,9 +8,10 @@
+ #include "mlir/Pass/PassManager.h"
+ #include "mlir/Transforms/Passes.h"
+ 
+-#include "mlir/Parser.h"
++#include "mlir/Parser/Parser.h"
+ #include "mlir/Support/FileUtilities.h"
+ 
++#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
+ #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+ #include "triton/Analysis/Allocation.h"
+ #include "triton/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.h"
+@@ -195,7 +196,7 @@ void init_triton_ir(py::module &&m) {
+                std::string attrName = name + "_arg" + std::to_string(id);
+                mlir::Block *owner = arg.getOwner();
+                if (owner->isEntryBlock() &&
+-                   !mlir::isa<mlir::FuncOp>(owner->getParentOp())) {
++                   !mlir::isa<mlir::func::FuncOp>(owner->getParentOp())) {
+                  owner->getParentOp()->setAttr(attrName, attr);
+                }
+              }
+@@ -348,7 +349,7 @@ void init_triton_ir(py::module &&m) {
+              return str;
+            })
+       .def("push_back",
+-           [](mlir::ModuleOp &self, mlir::FuncOp &funcOp) -> void {
++           [](mlir::ModuleOp &self, mlir::func::FuncOp &funcOp) -> void {
+              self.push_back(funcOp);
+            })
+       .def("has_function",
+@@ -358,16 +359,18 @@ void init_triton_ir(py::module &&m) {
+              return false;
+            })
+       .def("get_function",
+-           [](mlir::ModuleOp &self, std::string &funcName) -> mlir::FuncOp {
+-             return self.lookupSymbol<mlir::FuncOp>(funcName);
+-           })
+-      .def("get_single_function", [](mlir::ModuleOp &self) -> mlir::FuncOp {
+-        llvm::SmallVector<mlir::FuncOp> funcs;
+-        self.walk([&](mlir::FuncOp func) { funcs.push_back(func); });
+-        if (funcs.size() != 1)
+-          throw std::runtime_error("Expected a single function");
+-        return funcs[0];
+-      });
++           [](mlir::ModuleOp &self,
++              std::string &funcName) -> mlir::func::FuncOp {
++             return self.lookupSymbol<mlir::func::FuncOp>(funcName);
++           })
++      .def("get_single_function",
++           [](mlir::ModuleOp &self) -> mlir::func::FuncOp {
++             llvm::SmallVector<mlir::func::FuncOp> funcs;
++             self.walk([&](mlir::func::FuncOp func) { funcs.push_back(func); });
++             if (funcs.size() != 1)
++               throw std::runtime_error("Expected a single function");
++             return funcs[0];
++           });
+ 
+   m.def("make_attr",
+         [](const std::vector<int> &values, mlir::MLIRContext &context) {
+@@ -388,47 +391,48 @@ void init_triton_ir(py::module &&m) {
+         registry.insert<mlir::triton::TritonDialect,
+                         mlir::triton::gpu::TritonGPUDialect,
+                         mlir::math::MathDialect, mlir::arith::ArithmeticDialect,
+-                        mlir::StandardOpsDialect, mlir::scf::SCFDialect>();
++                        mlir::func::FuncDialect, mlir::scf::SCFDialect>();
+         context.appendDialectRegistry(registry);
+         context.loadAllAvailableDialects();
+ 
+         // parse module
+-        mlir::OwningOpRef<mlir::ModuleOp> module(
+-            mlir::parseSourceFile(inputFilename, &context));
++        mlir::OwningOpRef<mlir::ModuleOp> module =
++            mlir::parseSourceFile<mlir::ModuleOp>(inputFilename, &context);
++        if (!module)
++          throw std::runtime_error("Parse MLIR file failed.");
+         // locations are incompatible with ptx < 7.5 !
+         module->walk([](mlir::Operation *op) {
+           op->setLoc(mlir::UnknownLoc::get(op->getContext()));
+         });
+-        if (!module)
+-          throw std::runtime_error("Parse MLIR file failed.");
+ 
+         return module->clone();
+       },
+       ret::take_ownership);
+ 
+-  py::class_<mlir::FuncOp, mlir::OpState>(m, "function")
++  py::class_<mlir::func::FuncOp, mlir::OpState>(m, "function")
+       // .def_property_readonly("attrs", &ir::function::attrs)
+       // .def("add_attr", &ir::function::add_attr);
+       .def("args",
+-           [](mlir::FuncOp &self, unsigned idx) -> mlir::BlockArgument {
++           [](mlir::func::FuncOp &self, unsigned idx) -> mlir::BlockArgument {
+              return self.getArgument(idx);
+            })
+       .def(
+           "add_entry_block",
+-          [](mlir::FuncOp &self) -> mlir::Block * {
++          [](mlir::func::FuncOp &self) -> mlir::Block * {
+             return self.addEntryBlock();
+           },
+           ret::reference)
+       .def(
+           "set_arg_attr",
+-          [](mlir::FuncOp &self, int arg_no, const std::string &name, int val) {
++          [](mlir::func::FuncOp &self, int arg_no, const std::string &name,
++             int val) {
+             // set arg attributes "name" to value "val"
+             auto attrTy = mlir::IntegerType::get(self.getContext(), 32);
+             self.setArgAttr(arg_no, name, mlir::IntegerAttr::get(attrTy, val));
+           },
+           ret::reference)
+-      .def_property_readonly("type", &mlir::FuncOp::getType)
+-      .def("reset_type", &mlir::FuncOp::setType);
++      .def_property_readonly("type", &mlir::func::FuncOp::getFunctionType)
++      .def("reset_type", &mlir::func::FuncOp::setType);
+ 
+   py::class_<mlir::OpBuilder::InsertPoint>(m, "InsertPoint");
+ 
+@@ -445,13 +449,13 @@ void init_triton_ir(py::module &&m) {
+       .def("ret",
+            [](mlir::OpBuilder &self, std::vector<mlir::Value> &vals) -> void {
+              auto loc = self.getUnknownLoc();
+-             self.create<mlir::ReturnOp>(loc, vals);
++             self.create<mlir::func::ReturnOp>(loc, vals);
+            })
+       .def("call",
+-           [](mlir::OpBuilder &self, mlir::FuncOp &func,
++           [](mlir::OpBuilder &self, mlir::func::FuncOp &func,
+               std::vector<mlir::Value> &args) -> mlir::OpState {
+              auto loc = self.getUnknownLoc();
+-             return self.create<mlir::CallOp>(loc, func, args);
++             return self.create<mlir::func::CallOp>(loc, func, args);
+            })
+       // insertion block/point
+       .def("set_insertion_point_to_start",
+@@ -618,15 +622,16 @@ void init_triton_ir(py::module &&m) {
+       .def("get_or_insert_function",
+            [](mlir::OpBuilder &self, mlir::ModuleOp &module,
+               std::string &funcName, mlir::Type &funcType,
+-              std::string &visibility) -> mlir::FuncOp {
++              std::string &visibility) -> mlir::func::FuncOp {
+              if (mlir::Operation *funcOperation = module.lookupSymbol(funcName))
+-               return llvm::dyn_cast<mlir::FuncOp>(funcOperation);
++               return llvm::dyn_cast<mlir::func::FuncOp>(funcOperation);
+              auto loc = self.getUnknownLoc();
+              if (auto funcTy = funcType.dyn_cast<mlir::FunctionType>()) {
+                llvm::SmallVector<mlir::NamedAttribute> attrs = {
+                    mlir::NamedAttribute(self.getStringAttr("sym_visibility"),
+                                         self.getStringAttr(visibility))};
+-               return self.create<mlir::FuncOp>(loc, funcName, funcTy, attrs);
++               return self.create<mlir::func::FuncOp>(loc, funcName, funcTy,
++                                                      attrs);
+              }
+              throw std::runtime_error("invalid function type");
+            })
+@@ -658,15 +663,15 @@ void init_triton_ir(py::module &&m) {
+            [](mlir::OpBuilder &self, mlir::Value condition,
+               mlir::Block *trueDest, mlir::Block *falseDest) {
+              auto loc = self.getUnknownLoc();
+-             self.create<mlir::CondBranchOp>(loc, condition, trueDest,
+-                                             falseDest);
++             self.create<mlir::cf::CondBranchOp>(loc, condition, trueDest,
++                                                 falseDest);
+              return;
+            })
+       .def("create_branch",
+            [](mlir::OpBuilder &self, mlir::Block *dest,
+               std::vector<mlir::Value> &args) {
+              auto loc = self.getUnknownLoc();
+-             self.create<mlir::BranchOp>(loc, dest, args);
++             self.create<mlir::cf::BranchOp>(loc, dest, args);
+              return;
+            })
+       // Structured control flow
+@@ -792,14 +797,14 @@ void init_triton_ir(py::module &&m) {
+       .def("create_to_index",
+            [](mlir::OpBuilder &self, mlir::Value &input) -> mlir::Value {
+              auto loc = self.getUnknownLoc();
+-             return self.create<mlir::arith::IndexCastOp>(loc, input,
+-                                                          self.getIndexType());
++             return self.create<mlir::arith::IndexCastOp>(
++                 loc, self.getIndexType(), input);
+            })
+       .def("create_index_to_si",
+            [](mlir::OpBuilder &self, mlir::Value &input) -> mlir::Value {
+              auto loc = self.getUnknownLoc();
+-             return self.create<mlir::arith::IndexCastOp>(loc, input,
+-                                                          self.getI32Type());
++             return self.create<mlir::arith::IndexCastOp>(
++                 loc, self.getI32Type(), input);
+            })
+       .def("create_fmul",
+            [](mlir::OpBuilder &self, mlir::Value &lhs,
+@@ -1316,8 +1321,8 @@ void init_triton_ir(py::module &&m) {
+            [](mlir::OpBuilder &self, mlir::Value &condition,
+               mlir::Value &trueValue, mlir::Value &falseValue) -> mlir::Value {
+              auto loc = self.getUnknownLoc();
+-             return self.create<mlir::SelectOp>(loc, condition, trueValue,
+-                                                falseValue);
++             return self.create<mlir::arith::SelectOp>(loc, condition,
++                                                       trueValue, falseValue);
+            })
+       .def("create_printf",
+            [](mlir::OpBuilder &self, const std::string &prefix,
+@@ -1429,7 +1434,7 @@ void init_triton_ir(py::module &&m) {
+              self.addPass(mlir::triton::createConvertTritonGPUToLLVMPass());
+            })
+       .def("add_scf_to_cfg", [](mlir::PassManager &self) {
+-        self.addPass(mlir::createLowerToCFGPass());
++        self.addPass(mlir::createConvertSCFToCFPass());
+       });
+ }
+ 
+diff --git a/python/test/unit/language/test_core.py b/python/test/unit/language/test_core.py
+index 432544a8a4..018f544714 100644
+--- a/python/test/unit/language/test_core.py
++++ b/python/test/unit/language/test_core.py
+@@ -1918,7 +1918,7 @@ def test_convert2d(dtype, shape, src_layout, dst_layout, device='cuda'):
+ #dst = {dst_layout}
+ """ + """
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+-  func public @kernel_0d1d(%arg0: !tt.ptr<f16> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f16> {tt.divisibility = 16 : i32}) {
++  func.func public @kernel_0d1d(%arg0: !tt.ptr<f16> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f16> {tt.divisibility = 16 : i32}) {
+     %cst = arith.constant dense<128> : tensor<128x1xi32, #src>
+     %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #triton_gpu.slice<{dim = 1, parent = #src}>>
+     %1 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #triton_gpu.slice<{dim = 0, parent = #src}>>
+diff --git a/python/triton/compiler.py b/python/triton/compiler.py
+index 5d167634df..c36589037c 100644
+--- a/python/triton/compiler.py
++++ b/python/triton/compiler.py
+@@ -1514,14 +1514,14 @@ def make_hash(fn, **kwargs):
+     return hashlib.md5((Path(fn).read_text() + triton.runtime.jit.version_key()).encode("utf-8")).hexdigest()
+ 
+ 
+-# - ^\s*func\s+ : match the start of the string, any leading whitespace, the keyword func,
++# - ^\s*func\.func\s+ : match the start of the string, any leading whitespace, the keyword func,
+ #    and any following whitespace
+ # - (public\s+)? : optionally match the keyword public and any following whitespace
+ # - (@\w+) : match an @ symbol followed by one or more word characters
+ #   (letters, digits, or underscores), and capture it as group 1 (the function name)
+ # - (\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\)) : match a pair of parentheses enclosing
+ #   zero or more arguments separated by commas, and capture it as group 2 (the argument list)
+-mlir_prototype_pattern = r'^\s*func\s+(?:public\s+)?(@\w+)(\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\))\s*\{\s*$'
++mlir_prototype_pattern = r'^\s*func\.func\s+(?:public\s+)?(@\w+)(\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\))\s*\{\s*$'
+ ptx_prototype_pattern = r"\.(?:visible|extern)\s+\.(?:entry|func)\s+(\w+)\s*\(([^)]*)\)"
+ prototype_pattern = {
+     "ttir": mlir_prototype_pattern,
+diff --git a/test/Analysis/test-alias.mlir b/test/Analysis/test-alias.mlir
+index b3d5673f85..bb21615e68 100644
+--- a/test/Analysis/test-alias.mlir
++++ b/test/Analysis/test-alias.mlir
+@@ -11,7 +11,7 @@
+ 
+ // CHECK-LABEL: matmul_loop
+ // There shouldn't be any aliasing with the dot op encoding.
+-func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   %a_ptr_init = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL>
+   %b_ptr_init = tt.broadcast %B : (!tt.ptr<f16>) -> tensor<32x128x!tt.ptr<f16>, #BL>
+   %a_mask = arith.constant dense<true> : tensor<128x32xi1, #AL>
+@@ -36,7 +36,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B
+ }
+ 
+ // CHECK-LABEL: alloc
+-func @alloc(%A : !tt.ptr<f16>) {
++func.func @alloc(%A : !tt.ptr<f16>) {
+   // CHECK: %cst -> %cst
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   %cst1 = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #AL>
+@@ -46,7 +46,7 @@ func @alloc(%A : !tt.ptr<f16>) {
+ }
+ 
+ // CHECK-LABEL: convert
+-func @convert(%A : !tt.ptr<f16>) {
++func.func @convert(%A : !tt.ptr<f16>) {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL>
+   // CHECK: %0 -> %0
+   %cst1 = triton_gpu.convert_layout %cst0 : (tensor<16x16xf16, #AL>) -> tensor<16x16xf16, #A_SHARED>
+@@ -54,7 +54,7 @@ func @convert(%A : !tt.ptr<f16>) {
+ }
+ 
+ // CHECK-LABEL: trans
+-func @trans(%A : !tt.ptr<f16>) {
++func.func @trans(%A : !tt.ptr<f16>) {
+   // CHECK: %cst -> %cst
+   %tensor = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #A_SHARED>
+   // CHECK: %0 -> %cst
+@@ -63,7 +63,7 @@ func @trans(%A : !tt.ptr<f16>) {
+ }
+ 
+ // CHECK-LABEL: insert_slice_async
+-func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) {
++func.func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) {
+   %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL>
+   %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL>
+   %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL>
+@@ -76,7 +76,7 @@ func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) {
+ }
+ 
+ // CHECK-LABEL: insert_slice
+-func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) {
++func.func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) {
+   %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL>
+   %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL>
+   %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL>
+@@ -90,7 +90,7 @@ func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) {
+ }
+ 
+ // CHECK-LABEL: extract_slice
+-func @extract_slice(%A : !tt.ptr<f16>) {
++func.func @extract_slice(%A : !tt.ptr<f16>) {
+   // CHECK: %cst -> %cst
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<1x16x16xf16, #A_SHARED>
+   %index = arith.constant 0 : index
+@@ -100,7 +100,7 @@ func @extract_slice(%A : !tt.ptr<f16>) {
+ }
+ 
+ // CHECK-LABEL: if_cat
+-func @if_cat(%i1 : i1) {
++func.func @if_cat(%i1 : i1) {
+   // CHECK: %cst -> %cst
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   // CHECK: %cst_0 -> %cst_0
+@@ -119,7 +119,7 @@ func @if_cat(%i1 : i1) {
+ }
+ 
+ // CHECK-LABEL: if_alias
+-func @if_alias(%i1 : i1) {
++func.func @if_alias(%i1 : i1) {
+   // CHECK: %cst -> %cst
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   // CHECK-NEXT: %cst_0 -> %cst_0
+@@ -134,7 +134,7 @@ func @if_alias(%i1 : i1) {
+ }
+ 
+ // CHECK-LABEL: for
+-func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   // CHECK: %cst -> %cst
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: %cst_0 -> %cst_0
+@@ -154,7 +154,7 @@ func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.p
+ }
+ 
+ // CHECK-LABEL: for_if
+-func @for_if(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
++func.func @for_if(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
+   // CHECK: %cst -> %cst
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: %cst_0 -> %cst_0
+@@ -180,7 +180,7 @@ func @for_if(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !t
+ }
+ 
+ // CHECK-LABEL: for_if_for
+-func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
++func.func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
+   // CHECK: %cst -> %cst
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: %cst_0 -> %cst_0
+diff --git a/test/Analysis/test-alignment.mlir b/test/Analysis/test-alignment.mlir
+index 0ab34c7a78..af8ea6f856 100644
+--- a/test/Analysis/test-alignment.mlir
++++ b/test/Analysis/test-alignment.mlir
+@@ -1,288 +1,288 @@
+-// RUN: triton-opt %s -test-print-alignment -split-input-file 2>&1 | FileCheck %s
++// RUN: triton-opt %s -test-print-alignment -split-input-file -o %t 2>&1 | FileCheck %s
+ 
+-// CHECK-LABEL: cast
+-func @cast() {
+-  // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [1]
++// CHECK-LABEL: @cast
++func.func @cast() {
++  // CHECK: contiguity = [1], divisibility = [1], constancy = [1], constant_value = 1
+   %cst = arith.constant 1 : i32
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [1]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = 1
+   %0 = arith.extsi %cst : i32 to i64
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1
+   %cst_tensor = arith.constant dense<1> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1
+   %1 = tt.bitcast %cst_tensor : tensor<128xi32> -> tensor<128xi64>
+   return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: add
+-func @add() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++// CHECK-LABEL: @add
++func.func @add() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1
+   %1 = arith.constant dense<1> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [1], constancy = [1], constant_value = <none>
+   %2 = arith.addi %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [127]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 127
+   %3 = arith.constant dense<127> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128]
++  // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128
+   %4 = arith.addi %1, %3 : tensor<128xi32>
+   return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: sub
+-func @sub() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++// CHECK-LABEL: @sub
++func.func @sub() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1
+   %1 = arith.constant dense<1> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [1], constancy = [1], constant_value = <none>
+   %2 = arith.subi %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [129]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 129
+   %3 = arith.constant dense<129> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128]
++  // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128
+   %4 = arith.subi %3, %1 : tensor<128xi32>
+   return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: mul
+-func @mul() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++// CHECK-LABEL: @mul
++func.func @mul() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1
+   %1 = arith.constant dense<1> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %2 = arith.muli %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128]
++  // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128
+   %3 = arith.constant dense<128> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128]
++  // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128
+   %4 = arith.muli %3, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [2] ; Constancy: [128] ; ConstantValue: [2]
++  // CHECK-NEXT: contiguity = [1], divisibility = [2], constancy = [128], constant_value = 2
+   %5 = arith.constant dense<2> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [256] ; Constancy: [128] ; ConstantValue: [256]
++  // CHECK-NEXT: contiguity = [1], divisibility = [256], constancy = [128], constant_value = 256
+   %6 = arith.muli %4, %5 : tensor<128xi32>
+   return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: div
+-func @div() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++// CHECK-LABEL: @div
++func.func @div() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1
+   %1 = arith.constant dense<1> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %2 = arith.divsi %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %3 = arith.divui %1, %0 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64]
++  // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64
+   %4 = arith.constant dense<64> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16777216] ; Constancy: [64] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [16777216], constancy = [64], constant_value = <none>
+   %5 = arith.divsi %0, %4 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %6 = arith.divsi %4, %0 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64]
++  // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64
+   %7 = arith.divsi %4, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [2] ; Constancy: [128] ; ConstantValue: [66]
++  // CHECK-NEXT: contiguity = [1], divisibility = [2], constancy = [128], constant_value = 66
+   %8 = arith.constant dense<66> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [2] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [2], constant_value = <none>
+   %9 = arith.divui %0, %8 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [8192] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [8192], constancy = [1], constant_value = <none>
+   %10 = tt.make_range {end = 8320 : i32, start = 8192 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [64] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [64], constant_value = <none>
+   %11 = arith.divsi %10, %4 : tensor<128xi32>
+-  return 
++  return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: rem
+-func @rem() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++// CHECK-LABEL: @rem
++func.func @rem() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1
+   %1 = arith.constant dense<1> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0]
++  // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0
+   %2 = arith.remsi %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %3 = arith.remui %1, %0 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64]
++  // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64
+   %4 = arith.constant dense<64> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [64] ; Divisibility: [64] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [64], divisibility = [64], constancy = [1], constant_value = <none>
+   %5 = arith.remsi %0, %4 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [1], constant_value = <none>
+   %6 = arith.remsi %4, %0 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [2] ; Constancy: [128] ; ConstantValue: [66]
++  // CHECK-NEXT: contiguity = [1], divisibility = [2], constancy = [128], constant_value = 66
+   %7 = arith.constant dense<66> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [2] ; Divisibility: [2] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [2], divisibility = [2], constancy = [1], constant_value = <none>
+   %8 = arith.remui %0, %7 : tensor<128xi32>
+-  return 
++  return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: broadcast
+-func @broadcast() {
+-  // CHECK: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64]
++// CHECK-LABEL: @broadcast
++func.func @broadcast() {
++  // CHECK: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64
+   %0 = arith.constant dense<64> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [64, 1] ; Constancy: [128, 1] ; ConstantValue: [64]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [64, 1], constancy = [128, 1], constant_value = 64
+   %1 = tt.expand_dims %0 {axis = 1 : i32} : (tensor<128xi32>) -> tensor<128x1xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [64, 1] ; Constancy: [128, 128] ; ConstantValue: [64]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [64, 1], constancy = [128, 128], constant_value = 64
+   %2 = tt.broadcast %1 : (tensor<128x1xi32>) -> tensor<128x128xi32>
+   return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: splat
+-func @splat(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}) {
+-  // CHECK: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 128] ; ConstantValue: [None]
++// CHECK-LABEL: @splat
++func.func @splat(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}) {
++  // CHECK: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 128], constant_value = <none>
+   %0 = tt.splat %arg0 : (!tt.ptr<f32>) -> tensor<128x128x!tt.ptr<f32>>
+   return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: cmp
+-func @cmp() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++// CHECK-LABEL: @cmp
++func.func @cmp() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0]
++  // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0
+   %1 = arith.constant dense<0> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none>
+   %2 = arith.cmpi eq, %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none>
+   %3 = arith.cmpi slt, %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %4 = arith.cmpi sle, %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none>
+   %5 = arith.cmpi sge, %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8]
++  // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8
+   %6 = arith.constant dense<8> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = <none>
+   %7 = arith.cmpi sgt, %0, %6 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [0]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 0
+   %8 = arith.cmpi sgt, %1, %6 : tensor<128xi32>
+   return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: logic
+-func @logic() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++// CHECK-LABEL: @logic
++func.func @logic() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64]
++  // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64
+   %1 = arith.constant dense<64> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16777216] ; Constancy: [64] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [16777216], constancy = [64], constant_value = <none>
+   %2 = arith.divsi %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8]
++  // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8
+   %3 = arith.constant dense<8> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [134217728] ; Constancy: [8] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [134217728], constancy = [8], constant_value = <none>
+   %4 = arith.divsi %0, %3 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %5 = arith.andi %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %6 = arith.ori %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %7 = arith.xori %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = <none>
+   %8 = arith.andi %2, %4 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = <none>
+   %9 = arith.ori %2, %4 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = <none>
+   %10 = arith.xori %2, %4 : tensor<128xi32>
+   return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: select
+-func @select() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++// CHECK-LABEL: @select
++func.func @select() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0]
++  // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0
+   %1 = arith.constant dense<0> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none>
+   %2 = arith.cmpi eq, %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none>
+   %3 = arith.cmpi slt, %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [1] ; ConstantValue: [0]
++  // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [1], constant_value = 0
+   %4 = arith.constant 0 : i1
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0]
++  // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0
+   %7 = tt.splat %4 : (i1) -> tensor<128xi1>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0]
+-  %5 = select %4, %3, %7 : tensor<128xi1>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0
++  %5 = arith.select %4, %3, %7 : tensor<128xi1>
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none>
+   %8 = "triton_gpu.select"(%7, %3, %2) : (tensor<128xi1>, tensor<128xi1>, tensor<128xi1>) -> tensor<128xi1>
+   return
+ }
+ 
+ // -----
+ 
+-func @shift() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++func.func @shift() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8]
++  // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8
+   %1 = arith.constant dense<8> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4] ; Constancy: [128] ; ConstantValue: [4]
++  // CHECK-NEXT: contiguity = [1], divisibility = [4], constancy = [128], constant_value = 4
+   %2 = arith.constant dense<4> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [274877906944] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [274877906944], constancy = [1], constant_value = <none>
+   %3 = arith.shli %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [67108864] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [67108864], constancy = [1], constant_value = <none>
+   %4 = arith.shrsi %0, %2 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128]
++  // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128
+   %5 = arith.shli %1, %2 : tensor<128xi32>
+   return
+ }
+ 
+ // -----
+ 
+-func @max_min() {
+-  // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++func.func @max_min() {
++  // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [64] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [64], constancy = [1], constant_value = <none>
+   %1 = tt.make_range {end = 192 : i32, start = 64 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %2 = arith.maxsi %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %3 = arith.minsi %0, %1 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8]
++  // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8
+   %4 = arith.constant dense<8> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4] ; Constancy: [128] ; ConstantValue: [4]
++  // CHECK-NEXT: contiguity = [1], divisibility = [4], constancy = [128], constant_value = 4
+   %5 = arith.constant dense<4> : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [8]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = 8
+   %6 = arith.maxsi %4, %5 : tensor<128xi32>
+   return
+ }
+ 
+ // -----
+ 
+-// CHECK-LABEL: for
+-func @for() {
+-  // CHECK: Contiguity: [1, 1] ; Divisibility: [4611686018427387904, 4611686018427387904] ; Constancy: [128, 32] ; ConstantValue: [0]
++// CHECK-LABEL: @for
++func.func @for() {
++  // CHECK: contiguity = [1, 1], divisibility = [4611686018427387904, 4611686018427387904], constancy = [128, 32], constant_value = 0
+   %a_init = arith.constant dense<0> : tensor<128x32xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 32] ; ConstantValue: [1]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 32], constant_value = 1
+   %b_init = arith.constant dense<1> : tensor<128x32xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [4, 4] ; Constancy: [128, 32] ; ConstantValue: [4]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [4, 4], constancy = [128, 32], constant_value = 4
+   %c_init = arith.constant dense<4> : tensor<128x32xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [128]
++  // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [1], constant_value = 128
+   %ub = arith.constant 128 : index
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [1] ; ConstantValue: [0]
++  // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [1], constant_value = 0
+   %lb = arith.constant 0 : index
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [16]
++  // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [1], constant_value = 16
+   %step = arith.constant 16 : index
+   %a, %b, %c = scf.for %iv = %lb to %ub step %step iter_args(%a = %a_init, %b = %b_init, %c = %c_init) -> (tensor<128x32xi32>, tensor<128x32xi32>, tensor<128x32xi32>) {
+-    // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [None]
++    // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [1], constant_value = <none>
+     %t = arith.index_cast %iv : index to i32
+-    // CHECK: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 32] ; ConstantValue: [None]
+-    // CHECK: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 32] ; ConstantValue: [None]
+-    // CHECK: Contiguity: [1, 1] ; Divisibility: [4, 4] ; Constancy: [128, 32] ; ConstantValue: [4]
++    // CHECK: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 32], constant_value = <none>
++    // CHECK: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 32], constant_value = <none>
++    // CHECK: contiguity = [1, 1], divisibility = [4, 4], constancy = [128, 32], constant_value = 4
+     scf.yield %b, %a, %c : tensor<128x32xi32>, tensor<128x32xi32>, tensor<128x32xi32>
+   }
+   return
+@@ -290,53 +290,53 @@ func @for() {
+ 
+ // -----
+ 
+-// CHECK-LABEL: permute_2d
+-func @permute_2d(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32}) {
+-  // CHECK: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 128] ; ConstantValue: [1]
++// CHECK-LABEL: @permute_2d
++func.func @permute_2d(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32}) {
++  // CHECK: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 128], constant_value = 1
+   %cst = arith.constant dense<true> : tensor<128x128xi1>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [1, 1], constancy = [1, 1], constant_value = <none>
+   %cst_0 = arith.constant dense<0.000000e+00> : tensor<128x128xf32>
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %1 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [1073741824, 1] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128, 1], divisibility = [1073741824, 1], constancy = [1, 1], constant_value = <none>
+   %2 = tt.expand_dims %0 {axis = 1 : i32} : (tensor<128xi32>) -> tensor<128x1xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 1], constant_value = <none>
+   %3 = tt.splat %arg1 : (i32) -> tensor<128x1xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [17179869184, 16] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [17179869184, 16], constancy = [1, 1], constant_value = <none>
+   %4 = arith.muli %2, %3 : tensor<128x1xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 1], constant_value = <none>
+   %5 = tt.splat %arg0 : (!tt.ptr<f32>) -> tensor<128x1x!tt.ptr<f32>>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [1, 1], constant_value = <none>
+   %6 = tt.addptr %5, %4 : tensor<128x1x!tt.ptr<f32>>, tensor<128x1xi32>
+-  // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 1073741824] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 1073741824], constancy = [1, 1], constant_value = <none>
+   %7 = tt.expand_dims %1 {axis = 0 : i32}: (tensor<128xi32>) -> tensor<1x128xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [1, 128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [1, 128], constant_value = <none>
+   %8 = tt.broadcast %6 : (tensor<128x1x!tt.ptr<f32>>) -> tensor<128x128x!tt.ptr<f32>>
+-  // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 1073741824] ; Constancy: [128, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 1073741824], constancy = [128, 1], constant_value = <none>
+   %9 = tt.broadcast %7 : (tensor<1x128xi32>) -> tensor<128x128xi32>
+-  // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 16] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 16], constancy = [1, 1], constant_value = <none>
+   %10 = tt.addptr %8, %9 : tensor<128x128x!tt.ptr<f32>>, tensor<128x128xi32>
+-  // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [1073741824, 1] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128, 1], divisibility = [1073741824, 1], constancy = [1, 1], constant_value = <none>
+   %11 = tt.expand_dims %0 {axis = 1 : i32}: (tensor<128xi32>) -> tensor<128x1xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 1], constant_value = <none>
+   %12 = tt.splat %arg2 : (!tt.ptr<f32>) -> tensor<128x1x!tt.ptr<f32>>
+-  // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [16, 1] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128, 1], divisibility = [16, 1], constancy = [1, 1], constant_value = <none>
+   %13 = tt.addptr %12, %11 : tensor<128x1x!tt.ptr<f32>>, tensor<128x1xi32>
+-  // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 1073741824] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 1073741824], constancy = [1, 1], constant_value = <none>
+   %14 = tt.expand_dims %1 {axis = 0 : i32} : (tensor<128xi32>) -> tensor<1x128xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [1, 128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [1, 128], constant_value = <none>
+   %15 = tt.splat %arg3 : (i32) -> tensor<1x128xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 17179869184] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 17179869184], constancy = [1, 1], constant_value = <none>
+   %16 = arith.muli %14, %15 : tensor<1x128xi32>
+-  // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [16, 1] ; Constancy: [1, 128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128, 1], divisibility = [16, 1], constancy = [1, 128], constant_value = <none>
+   %17 = tt.broadcast %13 : (tensor<128x1x!tt.ptr<f32>>) -> tensor<128x128x!tt.ptr<f32>>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 17179869184] ; Constancy: [128, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 17179869184], constancy = [128, 1], constant_value = <none>
+   %18 = tt.broadcast %16 : (tensor<1x128xi32>) -> tensor<128x128xi32>
+-  // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [16, 1] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128, 1], divisibility = [16, 1], constancy = [1, 1], constant_value = <none>
+   %19 = tt.addptr %17, %18 : tensor<128x128x!tt.ptr<f32>>, tensor<128x128xi32>
+-  // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [1, 1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1, 1], divisibility = [1, 1], constancy = [1, 1], constant_value = <none>
+   %20 = tt.load %10, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<128x128xf32>
+   tt.store %19, %20, %cst : tensor<128x128xf32>
+   return
+@@ -347,29 +347,29 @@ func @permute_2d(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: i32 {t
+ module {
+ 
+ // This is a tiny test for verifying StoreOp-related alignment, It simply store a constant to a buffer.
+-// CHECK-LABEL: store_constant_align
+-func @store_constant_align(%addr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n: i32 {tt.divisibility = 16 : i32}) {
+-  // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++// CHECK-LABEL: @store_constant_align
++func.func @store_constant_align(%addr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n: i32 {tt.divisibility = 16 : i32}) {
++  // CHECK: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %pid = tt.get_program_id {axis = 0 : i32} : i32
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [128]
++  // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [1], constant_value = 128
+   %c128_i32 = arith.constant 128 : i32
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [1], constant_value = <none>
+   %1 = arith.muli %pid, %c128_i32 : i32
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none>
+   %2 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32>
+- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [None]
++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = <none>
+   %3 = tt.splat %1 : (i32) -> tensor<128xi32>
+- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [None]
++ // CHECK-NEXT: contiguity = [128], divisibility = [128], constancy = [1], constant_value = <none>
+   %4 = arith.addi %3, %2 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [128], constant_value = <none>
+   %5 = tt.splat %addr : (!tt.ptr<f32>) -> tensor<128x!tt.ptr<f32>>
+-  // CHECK-NEXT: Contiguity: [128] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [128], divisibility = [16], constancy = [1], constant_value = <none>
+   %6 = tt.addptr %5, %4 : tensor<128x!tt.ptr<f32>>, tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [128] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [128], constant_value = <none>
+   %9 = tt.splat %n : (i32) -> tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [16] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [16], constant_value = <none>
+   %mask = arith.cmpi slt, %4, %9 : tensor<128xi32>
+-  // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None]
++  // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %cst = arith.constant dense<0.0> : tensor<128xf32>
+   tt.store %5, %cst, %mask : tensor<128xf32>
+   return
+@@ -381,8 +381,8 @@ func @store_constant_align(%addr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n:
+ 
+ // This IR is dumped from vecadd test.
+ // Note, the hint {tt.divisibility = 16 : i32} for %n_elements affects the alignment of mask.
+-// CHECK-LABEL: vecadd_mask_align_16
+-func @vecadd_mask_align_16(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n_elements: i32 {tt.divisibility = 16 : i32}) {
++// CHECK-LABEL: @vecadd_mask_align_16
++func.func @vecadd_mask_align_16(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n_elements: i32 {tt.divisibility = 16 : i32}) {
+   %c64_i32 = arith.constant 64 : i32
+   %0 = tt.get_program_id {axis = 0 : i32} : i32
+   %1 = arith.muli %0, %c64_i32 : i32
+@@ -394,13 +394,13 @@ func @vecadd_mask_align_16(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %ar
+   %7 = tt.splat %arg1 : (!tt.ptr<f32>) -> tensor<64x!tt.ptr<f32>>
+   %8 = tt.addptr %7, %4 : tensor<64x!tt.ptr<f32>>, tensor<64xi32>
+   %9 = tt.splat %n_elements : (i32) -> tensor<64xi32>
+-  // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [16] ; ConstantValue: [None] ( %{{.*}} = arith.cmpi slt, %{{.*}}, %{{.*}} : tensor<64xi32> )
++  // CHECK: arith.cmpi slt, %{{.*}} => contiguity = [1], divisibility = [1], constancy = [16], constant_value = <none>
+   %mask = arith.cmpi slt, %4, %9 : tensor<64xi32>
+   %11 = tt.load %6, %mask {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32>
+   %12 = tt.load %8, %mask {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32>
+   %13 = arith.addf %11, %12 : tensor<64xf32>
+   %14 = tt.splat %arg2 : (!tt.ptr<f32>) -> tensor<64x!tt.ptr<f32>>
+-  // CHECK: Contiguity: [64] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [None] ( %{{.*}} = tt.addptr %{{.*}}, %{{.*}} : tensor<64x!tt.ptr<f32>>, tensor<64xi32> )
++  // CHECK: tt.addptr %{{.*}} => contiguity = [64], divisibility = [16], constancy = [1], constant_value = <none>
+   %15 = tt.addptr %14, %4 : tensor<64x!tt.ptr<f32>>, tensor<64xi32>
+   tt.store %15, %13, %mask : tensor<64xf32>
+   return
+@@ -410,8 +410,8 @@ func @vecadd_mask_align_16(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %ar
+ 
+ // This IR is dumped from vecadd test.
+ // Note, there is no divisibility hint for %n_elements, Triton should assume its divisibility to be 1 by default.
+-// CHECK-LABEL: vecadd_mask_align_1
+-func @vecadd_mask_align_1(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n_elements: i32) {
++// CHECK-LABEL: @vecadd_mask_align_1
++func.func @vecadd_mask_align_1(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n_elements: i32) {
+   %c64_i32 = arith.constant 64 : i32
+   %0 = tt.get_program_id {axis = 0 : i32} : i32
+   %1 = arith.muli %0, %c64_i32 : i32
+@@ -423,7 +423,7 @@ func @vecadd_mask_align_1(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg
+   %7 = tt.splat %arg1 : (!tt.ptr<f32>) -> tensor<64x!tt.ptr<f32>>
+   %8 = tt.addptr %7, %4 : tensor<64x!tt.ptr<f32>>, tensor<64xi32>
+   %9 = tt.splat %n_elements : (i32) -> tensor<64xi32>
+-  // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ( %{{.*}} = arith.cmpi slt, %{{.*}}, %{{.*}} : tensor<64xi32> )
++  // CHECK: arith.cmpi slt, %{{.*}} => contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none>
+   %10 = arith.cmpi slt, %4, %9 : tensor<64xi32>
+   %11 = tt.load %6, %10 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32>
+   %12 = tt.load %8, %10 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32>
+diff --git a/test/Analysis/test-allocation.mlir b/test/Analysis/test-allocation.mlir
+index efb00c404d..f79222aa7b 100644
+--- a/test/Analysis/test-allocation.mlir
++++ b/test/Analysis/test-allocation.mlir
+@@ -13,7 +13,7 @@
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ 
+ // CHECK-LABEL: matmul_loop
+-func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   %a_ptr_init = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL>
+   %b_ptr_init = tt.broadcast %B : (!tt.ptr<f16>) -> tensor<32x128x!tt.ptr<f16>, #BL>
+ 
+@@ -46,7 +46,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B
+ 
+ // Shared memory is available after a tensor's liveness range ends
+ // CHECK-LABEL: reusable
+-func @reusable(%A : !tt.ptr<f16>) {
++func.func @reusable(%A : !tt.ptr<f16>) {
+   %cst1 = arith.constant dense<true> : tensor<128x32xi1, #AL>
+   %cst2 = arith.constant dense<0.000000e+00> : tensor<128x32xf16, #AL>
+   %cst3 = arith.constant dense<true> : tensor<32x128xi1, #AL>
+@@ -78,7 +78,7 @@ func @reusable(%A : !tt.ptr<f16>) {
+ // %cst1->%cst4
+ // %cst3->%g->%h->%i
+ // CHECK-LABEL: preallocate
+-func @preallocate(%A : !tt.ptr<f16>) {
++func.func @preallocate(%A : !tt.ptr<f16>) {
+   // CHECK: offset = 0, size = 512
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   // CHECK-NEXT: offset = 1024, size = 512
+@@ -113,7 +113,7 @@ func @preallocate(%A : !tt.ptr<f16>) {
+ 
+ // Unused tensors are immediately released
+ // CHECK-LABEL: unused
+-func @unused(%A : !tt.ptr<f16>) {
++func.func @unused(%A : !tt.ptr<f16>) {
+   // CHECK: offset = 0, size = 1024
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<32x16xf16, #A_SHARED>
+   // CHECK-NEXT: offset = 0, size = 512
+@@ -128,7 +128,7 @@ func @unused(%A : !tt.ptr<f16>) {
+ 
+ // cst0 is alive through the entire function, it cannot be released before the end of the function
+ // CHECK-LABEL: longlive
+-func @longlive(%A : !tt.ptr<f16>) {
++func.func @longlive(%A : !tt.ptr<f16>) {
+   // CHECK: offset = 0, size = 512
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   // CHECK-NEXT: offset = 512, size = 512
+@@ -156,7 +156,7 @@ func @longlive(%A : !tt.ptr<f16>) {
+ }
+ 
+ // CHECK-LABEL: alloc
+-func @alloc(%A : !tt.ptr<f16>) {
++func.func @alloc(%A : !tt.ptr<f16>) {
+   // CHECK: offset = 0, size = 512
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   %cst1 = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #AL>
+@@ -167,7 +167,7 @@ func @alloc(%A : !tt.ptr<f16>) {
+ }
+ 
+ // CHECK-LABEL: scratch
+-func @scratch() {
++func.func @scratch() {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL>
+   // CHECK: scratch offset = 0, size = 512
+   %b = tt.reduce %cst0 {redOp = 1 : i32, axis = 0 : i32} : tensor<16x16xf16, #AL> -> tensor<16xf16, #sliceAd0>
+@@ -176,7 +176,7 @@ func @scratch() {
+ }
+ 
+ // CHECK-LABEL: trans
+-func @trans(%A : !tt.ptr<f16>) {
++func.func @trans(%A : !tt.ptr<f16>) {
+   // CHECK: offset = 0, size = 1024
+   %tensor = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #A_SHARED>
+   %b = tt.trans %tensor : (tensor<16x32xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED_T>
+@@ -184,7 +184,7 @@ func @trans(%A : !tt.ptr<f16>) {
+ }
+ 
+ // CHECK-LABEL: insert_slice_async
+-func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) {
++func.func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) {
+   %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL>
+   %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL>
+   %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL>
+@@ -197,7 +197,7 @@ func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) {
+ }
+ 
+ // CHECK-LABEL: extract_slice
+-func @extract_slice(%A : !tt.ptr<f16>) {
++func.func @extract_slice(%A : !tt.ptr<f16>) {
+   // CHECK: offset = 0, size = 512
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<1x16x16xf16, #A_SHARED>
+   %index = arith.constant 0 : index
+@@ -209,7 +209,7 @@ func @extract_slice(%A : !tt.ptr<f16>) {
+ // B0 -> (B1) -> B0
+ // Memory used by B1 can be reused by B0.
+ // CHECK-LABEL: if
+-func @if(%i1 : i1) {
++func.func @if(%i1 : i1) {
+   // CHECK: offset = 0, size = 512
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   // CHECK-NEXT: offset = 512, size = 512
+@@ -233,7 +233,7 @@ func @if(%i1 : i1) {
+ // B0 -> (B1) -> (B2) -> B0
+ // Memory used by B0 cannot be reused by B1 or B2.
+ // CHECK-LABEL: if_else
+-func @if_else(%i1 : i1) {
++func.func @if_else(%i1 : i1) {
+   // CHECK: offset = 0, size = 512
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   // CHECK-NEXT: offset = 512, size = 512
+@@ -260,7 +260,7 @@ func @if_else(%i1 : i1) {
+ // Block arguments and yields are memory aliases that do not trigger a new
+ // allocation.
+ // CHECK-LABEL: for
+-func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   // CHECK: offset = 0, size = 8192
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: offset = 8192, size = 8192
+@@ -275,7 +275,7 @@ func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.p
+ }
+ 
+ // CHECK-LABEL: for_if_slice
+-func @for_if_slice(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
++func.func @for_if_slice(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
+   // CHECK: offset = 0, size = 8192
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: offset = 8192, size = 8192
+@@ -296,7 +296,7 @@ func @for_if_slice(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %
+ 
+ // c0 cannot be released in the loop
+ // CHECK-LABEL: for_use_ancestor
+-func @for_use_ancestor(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
++func.func @for_use_ancestor(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
+   // CHECK: offset = 0, size = 8192
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: offset = 8192, size = 8192
+@@ -316,7 +316,7 @@ func @for_use_ancestor(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16
+ // a_shared_init, b_shared_init, and c_shared_init's liveness ranges are span over the entire function before cst2.
+ // So they cannot be reused by cst0 and cst1, but can be reused by cst2.
+ // CHECK-LABEL: for_if_for
+-func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
++func.func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) {
+   // CHECK: offset = 0, size = 8192
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: offset = 8192, size = 8192
+diff --git a/test/Analysis/test-membar.mlir b/test/Analysis/test-membar.mlir
+index 7199e5f53d..17880b2094 100644
+--- a/test/Analysis/test-membar.mlir
++++ b/test/Analysis/test-membar.mlir
+@@ -14,7 +14,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ 
+ // CHECK-LABEL: matmul_loop
+ // There shouldn't be any membar with the dot op encoding.
+-func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   %a_ptr_init = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL>
+   %b_ptr_init = tt.broadcast %B : (!tt.ptr<f16>) -> tensor<32x128x!tt.ptr<f16>, #BL>
+ 
+@@ -42,7 +42,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B
+ }
+ 
+ // CHECK-LABEL: raw_single_block
+-func @raw_single_block(%A : !tt.ptr<f16>) {
++func.func @raw_single_block(%A : !tt.ptr<f16>) {
+   %cst1 = arith.constant dense<true> : tensor<128x32xi1, #AL>
+   %cst2 = arith.constant dense<0.000000e+00> : tensor<128x32xf16, #AL>
+   %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL>
+@@ -54,7 +54,7 @@ func @raw_single_block(%A : !tt.ptr<f16>) {
+ }
+ 
+ // CHECK-LABEL: war_single_block
+-func @war_single_block(%A : !tt.ptr<f16>) {
++func.func @war_single_block(%A : !tt.ptr<f16>) {
+   %cst1 = arith.constant dense<true> : tensor<128x32xi1, #AL>
+   %cst2 = arith.constant dense<0.000000e+00> : tensor<128x32xf16, #AL>
+   %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL>
+@@ -70,7 +70,7 @@ func @war_single_block(%A : !tt.ptr<f16>) {
+ }
+ 
+ // CHECK-LABEL: scratch
+-func @scratch() {
++func.func @scratch() {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   // CHECK: Membar 1
+   %a = tt.cat %cst0, %cst0 {axis = 0} : (tensor<16x16xf16, #A_SHARED>, tensor<16x16xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED>
+@@ -81,7 +81,7 @@ func @scratch() {
+ }
+ 
+ // CHECK-LABEL: async_wait
+-func @async_wait() {
++func.func @async_wait() {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   // CHECK: Membar 1
+   %a = tt.cat %cst0, %cst0 {axis = 0} : (tensor<16x16xf16, #A_SHARED>, tensor<16x16xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED>
+@@ -92,7 +92,7 @@ func @async_wait() {
+ }
+ 
+ // CHECK-LABEL: alloc
+-func @alloc() {
++func.func @alloc() {
+   %cst0 = triton_gpu.alloc_tensor : tensor<16x16xf16, #A_SHARED>
+   %a = tt.cat %cst0, %cst0 {axis = 0} : (tensor<16x16xf16, #A_SHARED>, tensor<16x16xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED>
+   // CHECK: Membar 2
+@@ -101,7 +101,7 @@ func @alloc() {
+ }
+ 
+ // CHECK-LABEL: extract_slice
+-func @extract_slice() {
++func.func @extract_slice() {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<1x16x16xf16, #A_SHARED>
+   %index = arith.constant 0 : index
+   %cst1 = tensor.extract_slice %cst0[%index, 0, 0][1, 16, 16][1, 1, 1] : tensor<1x16x16xf16, #A_SHARED> to tensor<16x16xf16, #A_SHARED>
+@@ -113,14 +113,14 @@ func @extract_slice() {
+ }
+ 
+ // CHECK-LABEL: trans
+-func @trans() {
++func.func @trans() {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #A_SHARED>
+   %b = tt.trans %cst0 : (tensor<16x32xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED_T>
+   return
+ }
+ 
+ // CHECK-LABEL: insert_slice_async
+-func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) {
++func.func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) {
+   %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL>
+   %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL>
+   %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL>
+@@ -135,7 +135,7 @@ func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) {
+ }
+ 
+ // CHECK-LABEL: insert_slice
+-func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) {
++func.func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) {
+   %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL>
+   %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL>
+   %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL>
+@@ -153,7 +153,7 @@ func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) {
+ 
+ // If branch inserted a barrier for %cst0 and %cst1, but else didn't, then the barrier should be inserted in the parent region
+ // CHECK-LABEL: multi_blocks
+-func @multi_blocks(%i1 : i1) {
++func.func @multi_blocks(%i1 : i1) {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   scf.if %i1 {
+@@ -174,7 +174,7 @@ func @multi_blocks(%i1 : i1) {
+ 
+ // Both branches inserted a barrier for %cst0 and %cst1, then the barrier doesn't need to be inserted in the parent region
+ // CHECK-LABEL: multi_blocks_join_barrier
+-func @multi_blocks_join_barrier(%i1 : i1) {
++func.func @multi_blocks_join_barrier(%i1 : i1) {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   scf.if %i1 {
+@@ -192,7 +192,7 @@ func @multi_blocks_join_barrier(%i1 : i1) {
+ 
+ // Read yielded tensor requires a barrier
+ // CHECK-LABEL: multi_blocks_yield
+-func @multi_blocks_yield(%i1 : i1) {
++func.func @multi_blocks_yield(%i1 : i1) {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   %a = scf.if %i1 -> (tensor<32x16xf16, #A_SHARED>) {
+@@ -212,7 +212,7 @@ func @multi_blocks_yield(%i1 : i1) {
+ 
+ // Conservatively add a barrier as if the branch (%i1) is never taken
+ // CHECK-LABEL: multi_blocks_noelse
+-func @multi_blocks_noelse(%i1 : i1) {
++func.func @multi_blocks_noelse(%i1 : i1) {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   scf.if %i1 {
+@@ -226,7 +226,7 @@ func @multi_blocks_noelse(%i1 : i1) {
+ 
+ // Conservatively add a barrier as if the branch (%i2) is never taken
+ // CHECK-LABEL: multi_blocks_nested_scf
+-func @multi_blocks_nested_scf(%i1 : i1, %i2 : i1) {
++func.func @multi_blocks_nested_scf(%i1 : i1, %i2 : i1) {
+   %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED>
+   scf.if %i1 {
+@@ -247,7 +247,7 @@ func @multi_blocks_nested_scf(%i1 : i1, %i2 : i1) {
+ }
+ 
+ // CHECK-LABEL: for
+-func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   %c_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+@@ -262,7 +262,7 @@ func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.p
+ // Although a_shared and b_shared are synced before entering the loop,
+ // they are reassociated with aliases (c_shared) and thus require a barrier.
+ // CHECK-LABEL: for_alias
+-func @for_alias(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @for_alias(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: Membar 2
+@@ -282,7 +282,7 @@ func @for_alias(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B :
+ // Although cst2 is not an argument of scf.yield, its memory is reused by cst1.
+ // So we need a barrier both before and after cst1
+ // CHECK-LABEL: for_reuse
+-func @for_reuse(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @for_reuse(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: Membar 2
+@@ -302,7 +302,7 @@ func @for_reuse(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B :
+ 
+ 
+ // CHECK-LABEL: for_reuse_nested
+-func @for_reuse_nested(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @for_reuse_nested(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED>
+   // CHECK-NEXT: Membar 2
+diff --git a/test/Conversion/triton_ops.mlir b/test/Conversion/triton_ops.mlir
+index e9ee502435..0e979b148d 100644
+--- a/test/Conversion/triton_ops.mlir
++++ b/test/Conversion/triton_ops.mlir
+@@ -1,6 +1,6 @@
+ // RUN: triton-opt %s | FileCheck %s
+ 
+-func @cast_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_f32: f32, %scalar_i64: i64) {
++func.func @cast_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_f32: f32, %scalar_i64: i64) {
+   // scalar -> scalar
+   // CHECK:  i64 -> !tt.ptr<f32>
+   %0 = tt.int_to_ptr %scalar_i64 : i64 -> !tt.ptr<f32>
+@@ -35,7 +35,7 @@ func @cast_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_f32: f32, %scalar_i64: i64) {
+   return
+ }
+ 
+-func @addptr_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_i32: i32) {
++func.func @addptr_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_i32: i32) {
+   // scalar -> scalar
+   // CHECK: !tt.ptr<f32>
+   %0 = tt.addptr %scalar_ptr, %scalar_i32 : !tt.ptr<f32>, i32
+@@ -54,7 +54,7 @@ func @addptr_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_i32: i32) {
+   return
+ }
+ 
+-func @load_store_ops_scalar(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %mask : i1) {
++func.func @load_store_ops_scalar(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %mask : i1) {
+   // Test if Load/Store ops can handle scalar values
+   %other = arith.constant 0.0e+0 : f32
+ 
+@@ -76,7 +76,7 @@ func @load_store_ops_scalar(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %ma
+   return
+ }
+ 
+-func @reduce_ops_infer(%ptr: !tt.ptr<f32>, %v : tensor<1x2x4xf32>) {
++func.func @reduce_ops_infer(%ptr: !tt.ptr<f32>, %v : tensor<1x2x4xf32>) {
+   // Test if reduce ops infer types correctly
+ 
+   // CHECK: %{{.*}} = tt.reduce %{{.*}} -> tensor<2x4xf32>
+@@ -101,7 +101,7 @@ func @reduce_ops_infer(%ptr: !tt.ptr<f32>, %v : tensor<1x2x4xf32>) {
+   return
+ }
+ 
+-func @dot_ops_infer(%ptr: !tt.ptr<f32>, %v : f32) {
++func.func @dot_ops_infer(%ptr: !tt.ptr<f32>, %v : f32) {
+   // Test if reduce ops infer types correctly
+   %v128x32 = tt.splat %v : (f32) -> tensor<128x32xf32>
+   %v32x128 = tt.splat %v : (f32) -> tensor<32x128xf32>
+diff --git a/test/Conversion/triton_to_tritongpu.mlir b/test/Conversion/triton_to_tritongpu.mlir
+index a160bc8815..b461ca542f 100644
+--- a/test/Conversion/triton_to_tritongpu.mlir
++++ b/test/Conversion/triton_to_tritongpu.mlir
+@@ -1,6 +1,6 @@
+ // RUN: triton-opt %s -split-input-file -convert-triton-to-tritongpu=num-warps=2 | FileCheck %s
+ 
+-func @ops() {
++func.func @ops() {
+   // CHECK: module attributes {"triton_gpu.num-warps" = 2 : i32} {{.*}}
+   %a = arith.constant dense<1.00e+00> : tensor<128x32xf16>
+   %b = arith.constant dense<2.00e+00> : tensor<32x128xf16>
+@@ -11,7 +11,7 @@ func @ops() {
+ 
+ // -----
+ 
+-func @load_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) {
++func.func @load_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) {
+   // Test if LoadOp is lowered properly (see #771)
+   %ptrs = tt.splat %ptr : (!tt.ptr<f32>) -> tensor<128x!tt.ptr<f32>>
+   %mask = arith.constant dense<true> : tensor<128xi1>
+@@ -30,7 +30,7 @@ func @load_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) {
+ 
+ // -----
+ 
+-func @reduce_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) {
++func.func @reduce_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) {
+   // Test if the total number of threadsPerWarp is 32
+   // Test if the total number of warps is 2
+   // CHECK: #blocked0 = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [4, 8], warpsPerCTA = [1, 2], order = [0, 1]}>
+diff --git a/test/Conversion/tritongpu_to_llvm.mlir b/test/Conversion/tritongpu_to_llvm.mlir
+index e9e7d5a340..507b362c99 100644
+--- a/test/Conversion/tritongpu_to_llvm.mlir
++++ b/test/Conversion/tritongpu_to_llvm.mlir
+@@ -4,7 +4,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK: llvm.func @test_empty_kernel(%arg0: i32, %arg1: !llvm.ptr<f16, 1>)
+   // Here the 128 comes from the 4 in module attribute multiples 32
+   // CHECK:  attributes {nvvm.kernel = 1 : ui1, nvvm.maxntid = 128 : i32} {{.*}}
+-  func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) {
++  func.func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) {
+     // CHECK:  llvm.return
+     return
+   }
+@@ -15,7 +15,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_load
+-  func @basic_load(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) {
++  func.func @basic_load(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) {
+     // CHECK: llvm.inline_asm
+     // CHECK: llvm.inline_asm
+     %1 = tt.load %a_ptr_init, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked0>
+@@ -28,7 +28,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: vectorized_load
+-  func @vectorized_load(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) {
++  func.func @vectorized_load(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) {
+     // CHECK: llvm.inline_asm
+     // CHECK-SAME: ld.global.b32
+     // CHECK: llvm.inline_asm
+@@ -43,7 +43,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK-LABEL: vectorized_load_f16
+-  func @vectorized_load_f16(%a_ptr_init: tensor<256x!tt.ptr<f16>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf16, #blocked0>) {
++  func.func @vectorized_load_f16(%a_ptr_init: tensor<256x!tt.ptr<f16>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf16, #blocked0>) {
+     // CHECK: llvm.inline_asm
+     // CHECK-SAME: ld.global.b16
+     // CHECK: llvm.inline_asm
+@@ -59,7 +59,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [8], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: masked_load_const_other
+-  func @masked_load_const_other(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>) {
++  func.func @masked_load_const_other(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>) {
+     %cst_0 = arith.constant dense<0.000000e+00> : tensor<256xf32, #blocked0>
+     %1 = tt.load %a_ptr_init, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked0>
+     return
+@@ -72,7 +72,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [8], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: masked_load_const_other_vec
+-  func @masked_load_const_other_vec(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>) {
++  func.func @masked_load_const_other_vec(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>) {
+     %cst_0 = arith.constant dense<0.000000e+00> : tensor<256xf32, #blocked0>
+     %1 = tt.load %a_ptr_init, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked0>
+     return
+@@ -84,7 +84,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [2], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 2 : i32} {
+   // CHECK-LABEL: global_load_store_no_vec
+-  func @global_load_store_no_vec(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 4 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 4 : i32}, %arg3: i32) {
++  func.func @global_load_store_no_vec(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 4 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 4 : i32}, %arg3: i32) {
+     %c256_i32 = arith.constant 256 : i32
+     %0 = tt.get_program_id {axis = 0 : i32} : i32
+     %1 = arith.muli %0, %c256_i32 : i32
+@@ -128,7 +128,7 @@ module attributes {"triton_gpu.num-warps" = 2 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [2], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 2 : i32} {
+   // CHECK-LABEL: global_load_store_vec4
+-  func @global_load_store_vec4(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32) {
++  func.func @global_load_store_vec4(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32) {
+     %c256_i32 = arith.constant 256 : i32
+     %0 = tt.get_program_id {axis = 0 : i32} : i32
+     %1 = arith.muli %0, %c256_i32 : i32
+@@ -165,7 +165,7 @@ module attributes {"triton_gpu.num-warps" = 2 : i32} {
+ #blocked = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [2], order = [0]}>
+ // Note, the %n_elements doesn't have a "tt.divisibility" hint, so Triton assumes it's divisibility is 1, this should effect the mask's alignment and further restrict the load/store ops' vector width to be 1.
+ module attributes {"triton_gpu.num-warps" = 2 : i32} {
+-  func @vecadd_masked_vec1(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n_elements: i32) {
++  func.func @vecadd_masked_vec1(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n_elements: i32) {
+     %c64_i32 = arith.constant 64 : i32
+     %0 = tt.get_program_id {axis = 0 : i32} : i32
+     %1 = arith.muli %0, %c64_i32 : i32
+@@ -195,7 +195,7 @@ module attributes {"triton_gpu.num-warps" = 2 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK-LABEL: global_load_store_vec2
+-    func @global_load_store_vec2(%arg0: !tt.ptr<f32> {tt.divisibility = 8 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 8 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 8 : i32}, %arg3: i32) {
++    func.func @global_load_store_vec2(%arg0: !tt.ptr<f32> {tt.divisibility = 8 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 8 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 8 : i32}, %arg3: i32) {
+     %c256_i32 = arith.constant 256 : i32
+     %0 = tt.get_program_id {axis = 0 : i32} : i32
+     %1 = arith.muli %0, %c256_i32 : i32
+@@ -240,7 +240,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK-LABEL: global_load_store_vec8
+-    func @global_load_store_vec8(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32) {
++    func.func @global_load_store_vec8(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32) {
+     %c256_i32 = arith.constant 256 : i32
+     %0 = tt.get_program_id {axis = 0 : i32} : i32
+     %1 = arith.muli %0, %c256_i32 : i32
+@@ -283,7 +283,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ #blocked2 = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [4, 1], order = [0, 1]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_view_broadcast
+-  func @basic_view_broadcast(%arg : tensor<256xf32,#blocked0>) {
++  func.func @basic_view_broadcast(%arg : tensor<256xf32,#blocked0>) {
+     // CHECK: llvm.mlir.undef
+     // CHECK: %[[T0:.*]] = llvm.extractvalue
+     // CHECK: %[[T1:.*]] = llvm.extractvalue
+@@ -307,7 +307,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_make_range
+-  func @basic_make_range() {
++  func.func @basic_make_range() {
+     // CHECK: nvvm.read.ptx.sreg.tid.x
+     // CHECK: llvm.mlir.undef
+     // CHECK: llvm.insertvalue
+@@ -322,7 +322,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_addf
+-  func @basic_addf(%arg0 : tensor<256xf32,#blocked0>, %arg1 : tensor<256xf32,#blocked0>) {
++  func.func @basic_addf(%arg0 : tensor<256xf32,#blocked0>, %arg1 : tensor<256xf32,#blocked0>) {
+     // CHECK: llvm.fadd
+     // CHECK: llvm.fadd
+     %1 = arith.addf %arg0, %arg1 : tensor<256xf32,#blocked0>
+@@ -335,7 +335,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_addi
+-  func @basic_addi(%arg0 : tensor<256xi32,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) {
++  func.func @basic_addi(%arg0 : tensor<256xi32,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) {
+     // CHECK: llvm.add
+     // CHECK: llvm.add
+     %1 = arith.addi %arg0, %arg1 : tensor<256xi32,#blocked0>
+@@ -347,7 +347,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ 
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_program_id
+-  func @basic_program_id() {
++  func.func @basic_program_id() {
+     // CHECK: nvvm.read.ptx.sreg.ctaid.x : i32
+     %0 = tt.get_program_id {axis = 0 : i32} : i32
+     return
+@@ -359,7 +359,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_addptr
+-  func @basic_addptr(%arg0 : tensor<256x!tt.ptr<f32>,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) {
++  func.func @basic_addptr(%arg0 : tensor<256x!tt.ptr<f32>,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) {
+     // CHECK: llvm.getelementptr
+     // CHECK: llvm.getelementptr
+     %0 = tt.addptr %arg0, %arg1 : tensor<256x!tt.ptr<f32>, #blocked0>, tensor<256xi32, #blocked0>
+@@ -373,7 +373,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK: llvm.mlir.global external @global_smem
+   // CHECK-LABEL: basic_alloc_tensor
+-  func @basic_alloc_tensor() {
++  func.func @basic_alloc_tensor() {
+     // CHECK: llvm.mlir.addressof @global_smem
+     // CHECK-NEXT: llvm.bitcast
+     // CHECK-NEXT: llvm.mlir.constant
+@@ -390,7 +390,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK: llvm.mlir.global external @global_smem
+   // CHECK-LABEL: basic_extract_slice
+-  func @basic_extract_slice() {
++  func.func @basic_extract_slice() {
+     // CHECK: llvm.mlir.addressof @global_smem
+     // CHECK: llvm.extractvalue
+     // CHECK-NEXT: llvm.extractvalue
+@@ -423,7 +423,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ 
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_async_wait
+-  func @basic_async_wait() {
++  func.func @basic_async_wait() {
+     // CHECK: cp.async.wait_group 0x4
+     triton_gpu.async_wait {num = 4: i32}
+     return
+@@ -442,7 +442,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #A = #triton_gpu.shared<{vec = 8, perPhase = 1, maxPhase = 4, order = [1, 0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_insert_slice_async_fallback
+-  func @basic_insert_slice_async_fallback(%arg0: !tt.ptr<f16> {tt.divisibility = 1 : i32}) {
++  func.func @basic_insert_slice_async_fallback(%arg0: !tt.ptr<f16> {tt.divisibility = 1 : i32}) {
+     %off0_ = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #slice2d1>
+     %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<64xi32, #slice3d0>
+     %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<16xi32, #slice2d1>) -> tensor<16x1xi32, #block2>
+@@ -481,7 +481,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #A = #triton_gpu.shared<{vec = 8, perPhase = 1, maxPhase = 4, order = [1, 0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_insert_slice_async_v4
+-  func @basic_insert_slice_async_v4(%arg0: !tt.ptr<f32> {tt.divisibility = 32 : i32}) {
++  func.func @basic_insert_slice_async_v4(%arg0: !tt.ptr<f32> {tt.divisibility = 32 : i32}) {
+     %off0_ = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #slice2d1>
+     %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<64xi32, #slice3d0>
+     %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<16xi32, #slice2d1>) -> tensor<16x1xi32, #block2>
+@@ -523,7 +523,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #A = #triton_gpu.shared<{vec = 1, perPhase = 1, maxPhase = 4, order = [1, 0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_insert_slice_async_v1
+-  func @basic_insert_slice_async_v1(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}) {
++  func.func @basic_insert_slice_async_v1(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}) {
+     %off0_ = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #slice2d1>
+     %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #slice3d0>
+     %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<16xi32, #slice2d1>) -> tensor<16x1xi32, #block2>
+@@ -568,7 +568,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #A = #triton_gpu.shared<{vec = 1, perPhase = 1, maxPhase = 4, order = [1, 0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_insert_slice_async_v1_multictas
+-  func @basic_insert_slice_async_v1_multictas(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}) {
++  func.func @basic_insert_slice_async_v1_multictas(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}) {
+     %off0_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #slice2d1>
+     %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #slice3d0>
+     %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<32xi32, #slice2d1>) -> tensor<32x1xi32, #block2>
+@@ -619,7 +619,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK: basic_splat
+-  func @basic_splat(%ptr: !tt.ptr<f32>) {
++  func.func @basic_splat(%ptr: !tt.ptr<f32>) {
+     // CHECK: llvm.mlir.undef
+     // CHECK: llvm.insertvalue
+     // CHECK: llvm.insertvalue
+@@ -633,7 +633,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: basic_store
+-  func @basic_store(%ptrs: tensor<256x!tt.ptr<f32>, #blocked0>, %vals: tensor<256xf32, #blocked0>, %mask: tensor<256xi1, #blocked0>) {
++  func.func @basic_store(%ptrs: tensor<256x!tt.ptr<f32>, #blocked0>, %vals: tensor<256xf32, #blocked0>, %mask: tensor<256xi1, #blocked0>) {
+     // CHECK: llvm.inline_asm
+     // CHECK-SAME: st.global.b32 [ ${{.*}} + 0 ], { ${{.*}} };
+     // CHECK: llvm.inline_asm
+@@ -650,7 +650,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8>
+   // CHECK-LABEL: convert_layout_blocked_blocked
+-  func @convert_layout_blocked_blocked(%arg0: tensor<16x16xf32, #blocked0>) {
++  func.func @convert_layout_blocked_blocked(%arg0: tensor<16x16xf32, #blocked0>) {
+     // CHECK: llvm.mlir.addressof @global_smem
+     // CHECK: llvm.store
+     // CHECK-SAME: !llvm.ptr<vector<1xf32>, 3>
+@@ -697,7 +697,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8>
+   // CHECK-LABEL: convert_layout_blocked_blocked_vec
+-  func @convert_layout_blocked_blocked_vec(%arg0: tensor<16x16xf32, #blocked0>) {
++  func.func @convert_layout_blocked_blocked_vec(%arg0: tensor<16x16xf32, #blocked0>) {
+     // CHECK: llvm.mlir.addressof @global_smem
+     // CHECK: llvm.store
+     // CHECK-SAME: !llvm.ptr<vector<4xf32>, 3>
+@@ -720,7 +720,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8>
+   // CHECK-LABEL: convert_layout_blocked_blocked_multi_rep
+-  func @convert_layout_blocked_blocked_multi_rep(%arg0: tensor<16x16xf32, #blocked0>) {
++  func.func @convert_layout_blocked_blocked_multi_rep(%arg0: tensor<16x16xf32, #blocked0>) {
+     // CHECK: llvm.mlir.addressof @global_smem
+     // CHECK: llvm.store
+     // CHECK-SAME: !llvm.ptr<vector<4xf32>, 3>
+@@ -751,7 +751,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma0}>
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK-LABEL: convert_dot
+-  func @convert_dot(%A: tensor<16x16xf16, #blocked0>, %B: tensor<16x16xf16, #blocked0>) {
++  func.func @convert_dot(%A: tensor<16x16xf16, #blocked0>, %B: tensor<16x16xf16, #blocked0>) {
+     %AA = triton_gpu.convert_layout %A : (tensor<16x16xf16, #blocked0>) -> tensor<16x16xf16, #shared0>
+     %BB = triton_gpu.convert_layout %B : (tensor<16x16xf16, #blocked0>) -> tensor<16x16xf16, #shared0>
+     // CHECK: llvm.inline_asm
+@@ -775,7 +775,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ // TODO: problems in MLIR's parser on slice layout
+ // #blocked0 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [8, 4], warpsPerCTA = [1, 1], order = [1, 0]}>
+ // module attributes {"triton_gpu.num-warps" = 1 : i32} {
+-//   func @make_range_sliced_layout() {
++//   func.func @make_range_sliced_layout() {
+ //     %0 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #triton_gpu.slice<{dim = 0, parent = #blocked0}>>
+ //     return
+ //   }
+@@ -788,7 +788,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8>
+   // CHECK-LABEL: convert_layout_mmav2_block
+-  func @convert_layout_mmav2_blocked(%arg0: tensor<32x16xf32, #mma>) {
++  func.func @convert_layout_mmav2_blocked(%arg0: tensor<32x16xf32, #mma>) {
+     // CHECK: llvm.store
+     // CHECK-SAME: !llvm.ptr<vector<2xf32>, 3>
+     // CHECK: llvm.store
+@@ -808,7 +808,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8>
+   // CHECK-LABEL: convert_layout_mmav1_block
+-  func @convert_layout_mmav1_blocked(%arg0: tensor<32x64xf32, #mma>) {
++  func.func @convert_layout_mmav1_blocked(%arg0: tensor<32x64xf32, #mma>) {
+     // CHECK: llvm.store
+     // CHECK-SAME: !llvm.ptr<vector<2xf32>, 3>
+     // CHECK: llvm.store
+@@ -831,7 +831,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8>
+   // CHECK-LABEL: convert_layout_blocked_shared
+-  func @convert_layout_blocked_shared(%arg0: tensor<128x32xf32, #blocked0>) {
++  func.func @convert_layout_blocked_shared(%arg0: tensor<128x32xf32, #blocked0>) {
+     // CHECK: llvm.store
+     // CHECK-SAME: !llvm.ptr<vector<8xf32>, 3>
+     // CHECK: llvm.store
+@@ -847,7 +847,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ #blocked1 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [4, 8], warpsPerCTA = [1, 1], order = [1, 0]}>
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK-LABEL: convert_blocked1d_to_slice0
+-  func @convert_blocked1d_to_slice0(%src:tensor<32xi32, #blocked0>) {
++  func.func @convert_blocked1d_to_slice0(%src:tensor<32xi32, #blocked0>) {
+     // CHECK-COUNT-4: llvm.load {{.*}} : !llvm.ptr<vector<1xi32>, 3>
+     %cvt = triton_gpu.convert_layout %src : (tensor<32xi32, #blocked0>) -> tensor<32xi32, #triton_gpu.slice<{dim = 0, parent = #blocked1}>>
+     return
+@@ -860,7 +860,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ #blocked1 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [4, 8], warpsPerCTA = [1, 1], order = [1, 0]}>
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK-LABEL: convert_blocked1d_to_slice1
+-  func @convert_blocked1d_to_slice1(%src:tensor<32xi32, #blocked0>) {
++  func.func @convert_blocked1d_to_slice1(%src:tensor<32xi32, #blocked0>) {
+     // CHECK-COUNT-32: llvm.load {{.*}} : !llvm.ptr<vector<1xi32>, 3>
+     %cvt = triton_gpu.convert_layout %src : (tensor<32xi32, #blocked0>) -> tensor<32xi32, #triton_gpu.slice<{dim = 1, parent = #blocked1}>>
+     return
+@@ -873,7 +873,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ #blocked1 = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK-LABEL: convert_blocked_to_blocked_ptr
+-  func @convert_blocked_to_blocked_ptr(%src:tensor<32x!tt.ptr<f32>, #blocked0>) {
++  func.func @convert_blocked_to_blocked_ptr(%src:tensor<32x!tt.ptr<f32>, #blocked0>) {
+     // CHECK: llvm.ptrtoint
+     // CHECK: llvm.store
+     // CHECK: nvvm.barrier0
+@@ -892,7 +892,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ #dot_operand_a = #triton_gpu.dot_op<{opIdx=0, parent=#mma}>
+ #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+-  func @matmul_kernel_dot_operand_layout(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32},
++  func.func @matmul_kernel_dot_operand_layout(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32},
+   %a:tensor<128x32xf16, #shared>, %b:tensor<32x256xf16, #shared>) {
+     %cst = arith.constant dense<0.000000e+00> : tensor<128x256xf32, #mma>
+     // CHECK: ldmatrix.sync.aligned.m8n8.x4.shared.b16
+@@ -918,7 +918,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #dot_operand_a = #triton_gpu.dot_op<{opIdx=0, parent=#mma, isMMAv1Row=true}>
+ #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma, isMMAv1Row=true}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+-  func @matmul884_kernel_dot_operand_layout(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32},
++  func.func @matmul884_kernel_dot_operand_layout(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32},
+   %a:tensor<32x64xf16, #shared0>, %b:tensor<64x64xf16, #shared1>) {
+     %cst = arith.constant dense<0.000000e+00> : tensor<32x64xf32, #mma>
+     // CHECK: ldmatrix.sync.aligned.m8n8.x4.shared.b16
+@@ -941,7 +941,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #dot_operand_a = #triton_gpu.dot_op<{opIdx=0, parent=#blocked}>
+ #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#blocked}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+-  func @matmul_fmadot(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32},
++  func.func @matmul_fmadot(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32},
+   %a:tensor<32x16xf32, #shared>, %b:tensor<16x32xf32, #shared>) {
+     %cst = arith.constant dense<0.000000e+00> : tensor<32x32xf32, #blocked>
+     // CHECK: llvm.intr.fmuladd
+@@ -965,7 +965,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: matmul_tf32dot
+-  func @matmul_tf32dot(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32},
++  func.func @matmul_tf32dot(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32},
+   %a:tensor<32x16xf32, #shared>, %b:tensor<16x32xf32, #shared>) {
+     %cst = arith.constant dense<0.000000e+00> : tensor<32x32xf32, #mma>
+     // CHECK: llvm.inline_asm
+@@ -1000,7 +1000,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: atomic_add_f32
+-  func @atomic_add_f32(%arg0 : tensor<256x!tt.ptr<f32>, #blocked0>, %arg1 : tensor<256xi1, #blocked0>, %arg2 : tensor<256xf32, #blocked0>) {
++  func.func @atomic_add_f32(%arg0 : tensor<256x!tt.ptr<f32>, #blocked0>, %arg1 : tensor<256xi1, #blocked0>, %arg2 : tensor<256xf32, #blocked0>) {
+     // CHECK: llvm.inline_asm
+     // CHECK-SAME: atom.global.gpu.add.f32
+     %0 = "tt.atomic_rmw" (%arg0, %arg2, %arg1) {atomic_rmw_op = 5 : i32} : (tensor<256x!tt.ptr<f32>, #blocked0>, tensor<256xf32, #blocked0>, tensor<256xi1, #blocked0>) -> tensor<256xf32, #blocked0>
+@@ -1012,7 +1012,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ 
+-func @test_get_program_id(%a: tensor<32x!tt.ptr<i32>, #blocked0>) {
++func.func @test_get_program_id(%a: tensor<32x!tt.ptr<i32>, #blocked0>) {
+   %blockidx = tt.get_program_id {axis=0:i32} : i32
+   %blockidy = tt.get_program_id {axis=1:i32} : i32
+   %blockidz = tt.get_program_id {axis=2:i32} : i32
+@@ -1032,7 +1032,7 @@ func @test_get_program_id(%a: tensor<32x!tt.ptr<i32>, #blocked0>) {
+ // -----
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+-  func @test_get_num_program(%a: tensor<32x!tt.ptr<i32>, #blocked0>) {
++  func.func @test_get_num_program(%a: tensor<32x!tt.ptr<i32>, #blocked0>) {
+     // CHECK: nvvm.read.ptx.sreg.nctaid.x
+     // CHECK: nvvm.read.ptx.sreg.nctaid.y
+     // CHECK: nvvm.read.ptx.sreg.nctaid.z
+@@ -1052,7 +1052,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #blocked0 = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}>
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+   // CHECK-LABEL: test_index_cache
+-  func @test_index_cache() {
++  func.func @test_index_cache() {
+     // CHECK: nvvm.read.ptx.sreg.tid.x
+     %0 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32, #blocked0>
+     // CHECK-NOT: nvvm.read.ptx.sreg.tid.x
+@@ -1066,7 +1066,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ #shared0 = #triton_gpu.shared<{vec = 8, perPhase = 2, maxPhase = 4, order = [1, 0]}>
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK-LABEL: test_base_index_cache
+-  func @test_base_index_cache(%arg0: tensor<128x32xf32, #blocked0>) {
++  func.func @test_base_index_cache(%arg0: tensor<128x32xf32, #blocked0>) {
+     // CHECK: nvvm.read.ptx.sreg.tid.x
+     %0 = triton_gpu.convert_layout %arg0 : (tensor<128x32xf32, #blocked0>) -> tensor<128x32xf32, #shared0>
+     // CHECK-NOT: nvvm.read.ptx.sreg.tid.x
+@@ -1080,7 +1080,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+ #shared0 = #triton_gpu.shared<{vec = 8, perPhase = 2, maxPhase = 4, order = [1, 0]}>
+ module attributes {"triton_gpu.num-warps" = 1 : i32} {
+   // CHECK-LABEL: test_index_cache_different_block
+-  func @test_index_cache_different_block(%arg0: tensor<128x32xf32, #blocked0>, %arg1: i1) {
++  func.func @test_index_cache_different_block(%arg0: tensor<128x32xf32, #blocked0>, %arg1: i1) {
+     // CHECK: nvvm.read.ptx.sreg.tid.x
+     %0 = triton_gpu.convert_layout %arg0 : (tensor<128x32xf32, #blocked0>) -> tensor<128x32xf32, #shared0>
+     scf.if %arg1 {
+diff --git a/test/Target/tritongpu_to_llvmir.mlir b/test/Target/tritongpu_to_llvmir.mlir
+index cafff3ca60..114d3a9eb2 100644
+--- a/test/Target/tritongpu_to_llvmir.mlir
++++ b/test/Target/tritongpu_to_llvmir.mlir
+@@ -4,11 +4,11 @@
+ // CHECK-LABEL: ; ModuleID = 'LLVMDialectModule'
+ // CHECK: define void @test_empty_kernel
+ // CHECK: !nvvm.annotations
+-// CHECK: !{void (i32, half addrspace(1)*)* @test_empty_kernel, !"maxntidx", i32 128}
++// CHECK: !{ptr @test_empty_kernel, !"maxntidx", i32 128}
+ 
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ 
+-func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) {
++func.func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) {
+ 
+   return
+ }
+diff --git a/test/Target/tritongpu_to_ptx.mlir b/test/Target/tritongpu_to_ptx.mlir
+index 404e970a29..12742ad9e2 100644
+--- a/test/Target/tritongpu_to_ptx.mlir
++++ b/test/Target/tritongpu_to_ptx.mlir
+@@ -6,7 +6,7 @@
+ 
+ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ 
+-func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) {
++func.func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) {
+ 
+   return
+ }
+diff --git a/test/Triton/combine.mlir b/test/Triton/combine.mlir
+index 050a3f7565..5ef6790e69 100644
+--- a/test/Triton/combine.mlir
++++ b/test/Triton/combine.mlir
+@@ -2,10 +2,10 @@
+ // RUN: triton-opt %s -split-input-file -canonicalize -triton-combine | FileCheck %s
+ 
+ // CHECK-LABEL: @test_combine_dot_add_pattern
+-func @test_combine_dot_add_pattern() -> (tensor<128x128xf32>, tensor<128x128xf32>) {
+-    // CHECK: %[[d:.*]] = arith.constant dense<3.000000e+00> : tensor<128x128xf32>
+-    // CHECK: %[[b:.*]] = arith.constant dense<2.000000e+00> : tensor<128x128xf32>
+-    // CHECK: %[[a:.*]] = arith.constant dense<1.000000e+00> : tensor<128x128xf32>
++func.func @test_combine_dot_add_pattern() -> (tensor<128x128xf32>, tensor<128x128xf32>) {
++    // CHECK-DAG: %[[d:.*]] = arith.constant dense<3.000000e+00> : tensor<128x128xf32>
++    // CHECK-DAG: %[[b:.*]] = arith.constant dense<2.000000e+00> : tensor<128x128xf32>
++    // CHECK-DAG: %[[a:.*]] = arith.constant dense<1.000000e+00> : tensor<128x128xf32>
+     %a = arith.constant dense<1.0> : tensor<128x128xf32>
+     %b = arith.constant dense<2.0> : tensor<128x128xf32>
+     %zero = arith.constant dense<0.0> : tensor<128x128xf32>
+@@ -24,7 +24,7 @@ func @test_combine_dot_add_pattern() -> (tensor<128x128xf32>, tensor<128x128xf32
+ 
+ 
+ // COM: CHECK-LABEL: @test_combine_addptr_pattern
+-func @test_combine_addptr_pattern(%base: !tt.ptr<f32>) -> tensor<8x!tt.ptr<f32>> {
++func.func @test_combine_addptr_pattern(%base: !tt.ptr<f32>) -> tensor<8x!tt.ptr<f32>> {
+     %off0 = arith.constant 10 : i32
+     %off1 = arith.constant 15 : i32
+ 
+@@ -47,46 +47,46 @@ func @test_combine_addptr_pattern(%base: !tt.ptr<f32>) -> tensor<8x!tt.ptr<f32>>
+ 
+ 
+ // CHECK-LABEL: @test_combine_select_masked_load_pattern
+-func @test_combine_select_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %cond: i1) -> (tensor<8xf32>, tensor<8xf32>) {
++func.func @test_combine_select_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %cond: i1) -> (tensor<8xf32>, tensor<8xf32>) {
+     %mask = tt.broadcast %cond : (i1) -> tensor<8xi1>
+     %false_val = arith.constant dense<0.0> : tensor<8xf32>
+ 
+     // CHECK: %[[res1:.*]] = tt.load %{{.*}}, %{{.*}}, %{{.*}} {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32>
+     %x = tt.load %ptr, %mask, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32>
+-    %0 = select %cond, %x, %false_val : tensor<8xf32>
++    %0 = arith.select %cond, %x, %false_val : tensor<8xf32>
+ 
+     // CHECK: %[[res2:.*]] = tt.load %{{.*}}, %{{.*}}, %{{.*}} {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32>
+     %y = tt.load %ptr, %mask, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32>
+-    %1 = select %cond, %y, %false_val : tensor<8xf32>
++    %1 = arith.select %cond, %y, %false_val : tensor<8xf32>
+ 
+     // CHECK: return %[[res1]], %[[res2]] : tensor<8xf32>, tensor<8xf32>
+     return %0, %1 : tensor<8xf32>, tensor<8xf32>
+ }
+ 
+ // CHECK-LABEL: @test_combine_select_masked_load_fail_pattern
+-func @test_combine_select_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %dummy_load: tensor<8xf32>, %dummy_broadcast: tensor<8xi1>, %cond0: i1, %cond1: i1) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) {
++func.func @test_combine_select_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %dummy_load: tensor<8xf32>, %dummy_broadcast: tensor<8xi1>, %cond0: i1, %cond1: i1) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) {
+     %false_val = arith.constant dense<0.0> : tensor<8xf32>
+ 
+     // Case 1: value at the "load" position is not an "op".  Select should not be canonicalized.
+-    // CHECK: %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32>
+-    %0 = select %cond0, %dummy_load, %false_val : tensor<8xf32>
++    // CHECK: %{{.*}} = arith.select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32>
++    %0 = arith.select %cond0, %dummy_load, %false_val : tensor<8xf32>
+ 
+     // Case 2: value at the "broadcast" position is not an "op".  Select should not be canonicalized.
+     %real_load0 = tt.load %ptr, %dummy_broadcast, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32>
+-    // CHECK: %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32>
+-    %1 = select %cond0, %real_load0, %false_val : tensor<8xf32>
++    // CHECK: %{{.*}} = arith.select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32>
++    %1 = arith.select %cond0, %real_load0, %false_val : tensor<8xf32>
+ 
+     // Case 3: condition of "broadcast" is not the same as the condition of "select".  Select should not be canonicalized.
+     %cond0_ = tt.broadcast %cond0 : (i1) -> tensor<8xi1>
+     %real_load1 = tt.load %ptr, %cond0_, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32>
+-    // CHECK: %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32>
+-    %2 = select %cond1, %real_load1, %false_val : tensor<8xf32>
++    // CHECK: %{{.*}} = arith.select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32>
++    %2 = arith.select %cond1, %real_load1, %false_val : tensor<8xf32>
+ 
+     return %0, %1, %2 : tensor<8xf32>, tensor<8xf32>, tensor<8xf32>
+ }
+ 
+ // CHECK-LABEL: @test_combine_broadcast_constant_pattern
+-func @test_combine_broadcast_constant_pattern(%cst : f32) -> tensor<8x2xf32> {
++func.func @test_combine_broadcast_constant_pattern(%cst : f32) -> tensor<8x2xf32> {
+     // CHECK: %[[cst:.*]] = arith.constant dense<1.000000e+00> : tensor<8x2xf32>
+     %const = arith.constant dense<1.0> : tensor<8xf32>
+     %bst_out = tt.broadcast %const : (tensor<8xf32>) -> tensor<8x2xf32>
+@@ -96,7 +96,7 @@ func @test_combine_broadcast_constant_pattern(%cst : f32) -> tensor<8x2xf32> {
+ }
+ 
+ // CHECK-LABEL: @test_canonicalize_masked_load_pattern
+-func @test_canonicalize_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) {
++func.func @test_canonicalize_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) {
+     %true_mask = arith.constant dense<true> : tensor<8xi1>
+     %false_mask = arith.constant dense<false> : tensor<8xi1>
+     %other_val = arith.constant dense<0.0> : tensor<8xf32>
+@@ -117,7 +117,7 @@ func @test_canonicalize_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>) -> (te
+ }
+ 
+ // CHECK-LABEL: @test_canonicalize_masked_load_fail_pattern
+-func @test_canonicalize_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %mask: tensor<8xi1>) -> (tensor<8xf32>, tensor<8xf32>) {
++func.func @test_canonicalize_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %mask: tensor<8xi1>) -> (tensor<8xf32>, tensor<8xf32>) {
+     %other_val = arith.constant dense<0.0> : tensor<8xf32>
+ 
+     // Case: value at the "mask" position is not an "op".  Load should not be canonicalized.
+@@ -130,7 +130,7 @@ func @test_canonicalize_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %
+ }
+ 
+ // CHECK-LABEL: @test_canonicalize_masked_store_pattern
+-func @test_canonicalize_masked_store_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val: tensor<8xf32>) {
++func.func @test_canonicalize_masked_store_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val: tensor<8xf32>) {
+     %true_mask = arith.constant dense<true> : tensor<8xi1>
+     %false_mask = arith.constant dense<false> : tensor<8xi1>
+ 
+@@ -144,7 +144,7 @@ func @test_canonicalize_masked_store_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val:
+ }
+ 
+ // CHECK-LABEL: @test_canonicalize_masked_store_fail_pattern
+-func @test_canonicalize_masked_store_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val: tensor<8xf32>, %mask: tensor<8xi1>) {
++func.func @test_canonicalize_masked_store_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val: tensor<8xf32>, %mask: tensor<8xi1>) {
+     // Case: value at the "mask" position is not an "op".  Store should not be canonicalized.
+     // CHECK: tt.store %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32>
+     tt.store %ptr, %val, %mask : tensor<8xf32>
+diff --git a/test/Triton/vecadd.mlir b/test/Triton/vecadd.mlir
+index 0b69ef3054..f5019b1cdd 100644
+--- a/test/Triton/vecadd.mlir
++++ b/test/Triton/vecadd.mlir
+@@ -1,7 +1,7 @@
+ // RUN: triton-opt %s -verify-diagnostics
+ 
+ module {
+-  func @add_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32__(%arg0: !tt.ptr<f32>, %arg1: !tt.ptr<f32>, %arg2: !tt.ptr<f32>, %arg3: i32, %arg4: i32, %arg5: i32) {
++  func.func @add_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32__(%arg0: !tt.ptr<f32>, %arg1: !tt.ptr<f32>, %arg2: !tt.ptr<f32>, %arg3: i32, %arg4: i32, %arg5: i32) {
+     %0 = tt.get_program_id {axis = 0 : i32} : i32
+     %c256_i32 = arith.constant 256 : i32
+     %1 = arith.muli %0, %c256_i32 : i32
+@@ -43,7 +43,7 @@ module {
+   }
+ }
+ // module {
+-//   func @add_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32__(%arg0: !tt.ptr<f32>, %arg1: !tt.ptr<f32>, %arg2: !tt.ptr<f32>, %arg3: i32, %arg4: i32, %arg5: i32) {
++//   func.func @add_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32__(%arg0: !tt.ptr<f32>, %arg1: !tt.ptr<f32>, %arg2: !tt.ptr<f32>, %arg3: i32, %arg4: i32, %arg5: i32) {
+ //     %c64 = arith.constant 64 : index
+ //     %c32 = arith.constant 32 : index
+ //     %c0 = arith.constant 0 : index
+diff --git a/test/TritonGPU/coalesce.mlir b/test/TritonGPU/coalesce.mlir
+index 60e359f527..51cccccfbd 100644
+--- a/test/TritonGPU/coalesce.mlir
++++ b/test/TritonGPU/coalesce.mlir
+@@ -19,7 +19,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} {
+ // CHECK: [[store_val:%.*]] = triton_gpu.convert_layout {{.*}} -> tensor<64x64xf32, [[col_layout]]>
+ // CHECK: [[store_mask:%.*]] = triton_gpu.convert_layout {{.*}} -> tensor<64x64xi1, [[col_layout]]>
+ // CHECK: tt.store [[store_ptr]], [[store_val]], [[store_mask]]
+-func @transpose(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32},
++func.func @transpose(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32},
+                 %arg1: i32 {tt.divisibility = 16 : i32},
+                 %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32},
+                 %arg3: i32 {tt.divisibility = 16 : i32}) {
+diff --git a/test/TritonGPU/combine.mlir b/test/TritonGPU/combine.mlir
+index 2c009ffa48..7e9cb9d504 100644
+--- a/test/TritonGPU/combine.mlir
++++ b/test/TritonGPU/combine.mlir
+@@ -9,7 +9,7 @@
+ // CHECK: [[col_layout:#.*]] = #triton_gpu.blocked<{sizePerThread = [4, 1], threadsPerWarp = [16, 2], warpsPerCTA = [4, 1], order = [0, 1]}>
+ // CHECK: [[col_layout_novec:#.*]] = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [4, 1], order = [0, 1]}>
+ // CHECK-LABEL: cst
+-func @cst() -> tensor<1024xi32, #layout1> {
++func.func @cst() -> tensor<1024xi32, #layout1> {
+   %cst = arith.constant dense<0> : tensor<1024xi32, #layout0>
+   %1 = triton_gpu.convert_layout %cst : (tensor<1024xi32, #layout0>) -> tensor<1024xi32, #layout1>
+   // CHECK-NOT: triton_gpu.convert_layout
+@@ -18,7 +18,7 @@ func @cst() -> tensor<1024xi32, #layout1> {
+ }
+ 
+ // CHECK-LABEL: range
+-func @range() -> tensor<1024xi32, #layout1> {
++func.func @range() -> tensor<1024xi32, #layout1> {
+   %0 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #layout0>
+   %1 = triton_gpu.convert_layout %0 : (tensor<1024xi32, #layout0>) -> tensor<1024xi32, #layout1>
+   // CHECK-NOT: triton_gpu.convert_layout
+@@ -27,7 +27,7 @@ func @range() -> tensor<1024xi32, #layout1> {
+ }
+ 
+ // CHECK-LABEL: splat
+-func @splat(%arg0: i32) -> tensor<1024xi32, #layout1> {
++func.func @splat(%arg0: i32) -> tensor<1024xi32, #layout1> {
+   %0 = tt.splat %arg0 : (i32) -> tensor<1024xi32, #layout0>
+   %1 = triton_gpu.convert_layout %0 : (tensor<1024xi32, #layout0>) -> tensor<1024xi32, #layout1>
+   // CHECK-NOT: triton_gpu.convert_layout
+@@ -36,7 +36,7 @@ func @splat(%arg0: i32) -> tensor<1024xi32, #layout1> {
+ }
+ 
+ // CHECK-LABEL: remat
+-func @remat(%arg0: i32) -> tensor<1024xi32, #layout1> {
++func.func @remat(%arg0: i32) -> tensor<1024xi32, #layout1> {
+   %0 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #layout0>
+   %1 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #layout0>
+   %2 = arith.muli %0, %1 : tensor<1024xi32, #layout0>
+@@ -56,7 +56,7 @@ func @remat(%arg0: i32) -> tensor<1024xi32, #layout1> {
+ }
+ 
+ // CHECK-LABEL: remat_load_store
+-func @remat_load_store(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
++func.func @remat_load_store(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+   %0 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #layout0>
+   %1 = tt.splat %arg : (!tt.ptr<i32>) -> tensor<64x!tt.ptr<i32>, #layout0>
+   %2 = tt.addptr %1, %0 : tensor<64x!tt.ptr<i32>, #layout0>, tensor<64xi32, #layout0>
+@@ -70,7 +70,7 @@ func @remat_load_store(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+ 
+ // Don't rematerialize vectorized loads
+ // CHECK-LABEL: remat_expensive
+-func @remat_expensive(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
++func.func @remat_expensive(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+   %0 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #layout1>
+   %1 = tt.splat %arg : (!tt.ptr<i32>) -> tensor<64x!tt.ptr<i32>, #layout1>
+   %2 = tt.addptr %1, %0 : tensor<64x!tt.ptr<i32>, #layout1>, tensor<64xi32, #layout1>
+@@ -85,7 +85,7 @@ func @remat_expensive(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+ 
+ // Don't rematerialize loads when original and target layouts are different
+ // CHECK-LABEL: remat_multi_layout
+-func @remat_multi_layout(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
++func.func @remat_multi_layout(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+   %0 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #layout0>
+   %1 = tt.splat %arg : (!tt.ptr<i32>) -> tensor<64x!tt.ptr<i32>, #layout0>
+   %2 = tt.addptr %1, %0 : tensor<64x!tt.ptr<i32>, #layout0>, tensor<64xi32, #layout0>
+@@ -100,7 +100,7 @@ func @remat_multi_layout(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+ 
+ // Always rematerialize single value loads
+ // CHECK-LABEL: remat_single_value
+-func @remat_single_value(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
++func.func @remat_single_value(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+   %0 = tt.splat %arg : (!tt.ptr<i32>) -> tensor<1x!tt.ptr<i32>, #layout1>
+   %1 = tt.load %0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<1xi32, #layout1>
+   // CHECK-NOT: triton_gpu.convert_layout
+@@ -111,7 +111,7 @@ func @remat_single_value(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+ }
+ 
+ // CHECK-LABEL: if
+-func @if(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
++func.func @if(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+   // CHECK-NOT: triton_gpu.convert_layout
+   %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout1>
+   %0 = tt.get_program_id {axis = 0 : i32} : i32
+@@ -128,7 +128,7 @@ func @if(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+ }
+ 
+ // CHECK-LABEL: if_convert_else_not
+-func @if_convert_else_not(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
++func.func @if_convert_else_not(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+   %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout0>
+   %0 = tt.get_program_id {axis = 0 : i32} : i32
+   %1 = tt.splat %0 : (i32) -> tensor<1024xi32, #layout0>
+@@ -149,7 +149,7 @@ func @if_convert_else_not(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16
+ }
+ 
+ // CHECK-LABEL: if_not_else_convert
+-func @if_not_else_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
++func.func @if_not_else_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+   %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout0>
+   %0 = tt.get_program_id {axis = 0 : i32} : i32
+   %1 = tt.splat %0 : (i32) -> tensor<1024xi32, #layout0>
+@@ -170,7 +170,7 @@ func @if_not_else_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16
+ }
+ 
+ // CHECK-LABEL: if_else_both_convert
+-func @if_else_both_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
++func.func @if_else_both_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) {
+   %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout0>
+   %0 = tt.get_program_id {axis = 0 : i32} : i32
+   %1 = tt.splat %0 : (i32) -> tensor<1024xi32, #layout0>
+@@ -200,7 +200,7 @@ func @if_else_both_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16
+ #blocked4 = #triton_gpu.blocked<{sizePerThread = [4, 1], threadsPerWarp = [16, 2], warpsPerCTA = [4, 1], order = [0, 1]}>
+ 
+ // CHECK-LABEL: transpose
+-func @transpose(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32}) {
++func.func @transpose(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32}) {
+   // CHECK-NOT: triton_gpu.convert_layout
+   // CHECK: [[loaded_val:%.*]] = tt.load {{.*}}, {{%cst.*}}, {{%cst.*}} {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64x64xf32, [[row_layout]]>
+   // CHECK: [[cvt_val:%.*]] = triton_gpu.convert_layout [[loaded_val]] : (tensor<64x64xf32, [[row_layout]]>) -> tensor<64x64xf32, [[col_layout]]>
+@@ -241,7 +241,7 @@ func @transpose(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: i32 {tt
+ }
+ 
+ // CHECK-LABEL: loop
+-func @loop(%arg0: !tt.ptr<f32>, %arg1: i32, %arg2: !tt.ptr<f32>, %arg3: i32, %arg4: i32) {
++func.func @loop(%arg0: !tt.ptr<f32>, %arg1: i32, %arg2: !tt.ptr<f32>, %arg3: i32, %arg4: i32) {
+     // CHECK-NOT: triton_gpu.convert_layout
+     // CHECK: [[loop_ret:%.*]]:2 = scf.for {{.*}} -> (tensor<64x64xf32, [[row_layout]]>, tensor<64x64x!tt.ptr<f32>, [[row_layout]]>)
+     // CHECK-NEXT: {{.*}} = tt.load {{.*}} : tensor<64x64xf32, [[row_layout]]>
+@@ -295,7 +295,7 @@ func @loop(%arg0: !tt.ptr<f32>, %arg1: i32, %arg2: !tt.ptr<f32>, %arg3: i32, %ar
+ }
+ 
+ // CHECK-LABEL: vecadd
+-func @vecadd(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32) {
++func.func @vecadd(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32) {
+   // CHECK-NOT: triton_gpu.convert_layout
+   %c256_i32 = arith.constant 256 : i32
+   %0 = tt.get_program_id {axis = 0 : i32} : i32
+@@ -327,7 +327,7 @@ func @vecadd(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f3
+ 
+ // Select has args with different element types
+ // CHECK-LABEL: select
+-func @select(%arg0: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}) {
++func.func @select(%arg0: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}) {
+   // CHECK-NOT: triton_gpu.convert_layout
+   %cst = arith.constant dense<30000> : tensor<1x1xi32, #blocked2>
+   %cst_0 = arith.constant dense<30000> : tensor<1x512xi32, #blocked2>
+@@ -378,7 +378,7 @@ func @select(%arg0: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f6
+ 
+ // Make sure the following IR doesn't hang the compiler.
+ // CHECK-LABEL: long_func
+-func public @long_func(%arg0: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg3: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg4: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg5: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg6: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg7: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg8: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg9: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg10: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg11: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg12: !tt.ptr<i32> {tt.divisibility = 16 : i32}, %arg13: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg14: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg15: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg16: i32 {tt.divisibility = 16 : i32}) {
++func.func public @long_func(%arg0: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg3: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg4: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg5: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg6: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg7: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg8: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg9: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg10: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg11: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg12: !tt.ptr<i32> {tt.divisibility = 16 : i32}, %arg13: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg14: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg15: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg16: i32 {tt.divisibility = 16 : i32}) {
+   %cst = arith.constant dense<1.000000e+00> : tensor<1024xf32, #blocked0>
+   %cst_0 = arith.constant dense<5.000000e-04> : tensor<1024xf32, #blocked0>
+   %cst_1 = arith.constant dense<0.999499976> : tensor<1024xf32, #blocked0>
+@@ -775,7 +775,7 @@ func public @long_func(%arg0: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg1:
+ // A mnist model from torch inductor.
+ // Check if topological sort is working correct and there's no unnecessary convert
+ // CHECK-LABEL: mnist
+-func public @mnist(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}, %arg3: i32) {
++func.func public @mnist(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}, %arg3: i32) {
+   // CHECK-NOT: triton_gpu.convert_layout
+   %cst = arith.constant dense<10> : tensor<16x1xi32, #blocked2>
+   %cst_0 = arith.constant dense<10> : tensor<1x16xi32, #blocked3>
+@@ -862,7 +862,7 @@ func public @mnist(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.
+ #blocked5 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [2, 16], warpsPerCTA = [8, 1], order = [1, 0]}>
+ // cmpf and cmpi have different operands and result types
+ // CHECK-LABEL: cmp
+-func public @cmp(%arg0: !tt.ptr<f16> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: !tt.ptr<f16> {tt.divisibility = 16 : i32}, %arg4: i32 {tt.divisibility = 16 : i32}, %arg5: i32 {tt.divisibility = 16 : i32}) {
++func.func public @cmp(%arg0: !tt.ptr<f16> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: !tt.ptr<f16> {tt.divisibility = 16 : i32}, %arg4: i32 {tt.divisibility = 16 : i32}, %arg5: i32 {tt.divisibility = 16 : i32}) {
+   %c64 = arith.constant 64 : index
+   %c2048 = arith.constant 2048 : index
+   %c0 = arith.constant 0 : index
+diff --git a/test/TritonGPU/loop-pipeline.mlir b/test/TritonGPU/loop-pipeline.mlir
+index 6ee3b15fbc..663f2da7b0 100644
+--- a/test/TritonGPU/loop-pipeline.mlir
++++ b/test/TritonGPU/loop-pipeline.mlir
+@@ -10,7 +10,7 @@
+ #A = #triton_gpu.dot_op<{opIdx = 0, parent = #C}>
+ #B = #triton_gpu.dot_op<{opIdx = 1, parent = #C}>
+ 
+-// CHECK: func @matmul_loop
++// CHECK: func.func @matmul_loop
+ // CHECK-DAG: %[[CONSTANT_0:.*]] = arith.constant 0 : i32
+ // CHECK-DAG: %[[CONSTANT_1:.*]] = arith.constant 1 : i32
+ // CHECK-DAG: %[[CONSTANT_2:.*]] = arith.constant 2 : i32
+@@ -46,8 +46,8 @@
+ // CHECK-DAG: %[[NEXT_PIPELINE_IDX:.*]] = arith.addi %[[PIPELINE_IDX]], %[[CONSTANT_1]]
+ // CHECK-DAG: %[[NEXT_LOOP_IDX:.*]] = arith.addi %[[LOOP_IDX]], %[[CONSTANT_1]]
+ // CHECK:   scf.yield {{.*}}, {{.*}}, {{.*}}, %[[NEXT_A_BUFFER]], %[[NEXT_B_BUFFER]], %[[NEXT_A]], %[[NEXT_B]], {{.*}}, {{.*}}, {{.*}}, %[[NEXT_PIPELINE_IDX]], %[[NEXT_LOOP_IDX]]
+-func @matmul_loop(%lb : index, %ub : index, %step : index, 
+-                  %A : !tt.ptr<f16> {tt.divisibility = 16 : i32}, 
++func.func @matmul_loop(%lb : index, %ub : index, %step : index,
++                  %A : !tt.ptr<f16> {tt.divisibility = 16 : i32},
+                   %B : !tt.ptr<f16> {tt.divisibility = 16 : i32}) {
+   // A ptrs
+   %a_ptr_splat = tt.splat %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL>
+@@ -61,7 +61,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index,
+   %b_tmp1 = tt.expand_dims %b_tmp0 {axis = 0 : i32} : (tensor<128xi32, #BLs0>) -> tensor<1x128xi32, #BL>
+   %b_offs = tt.broadcast %b_tmp1 : (tensor<1x128xi32, #BL>) -> tensor<32x128xi32, #BL>
+   %b_ptr_init = tt.addptr %b_ptr_splat, %b_offs : tensor<32x128x!tt.ptr<f16>, #BL>, tensor<32x128xi32, #BL>
+-  
++
+ 
+   %a_mask = arith.constant dense<true> : tensor<128x32xi1, #AL>
+   %a_other = arith.constant dense<0.00e+00> : tensor<128x32xf16, #AL>
+@@ -88,7 +88,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index,
+ }
+ 
+ 
+-// CHECK: func @matmul_loop_nested
++// CHECK: func.func @matmul_loop_nested
+ // CHECK-DAG: %[[CONSTANT_0:.*]] = arith.constant 0 : i32
+ // CHECK-DAG: %[[CONSTANT_1:.*]] = arith.constant 1 : i32
+ // CHECK-DAG: %[[CONSTANT_2:.*]] = arith.constant 2 : i32
+@@ -118,8 +118,8 @@ func @matmul_loop(%lb : index, %ub : index, %step : index,
+ // CHECK-DAG: %[[NEXT_PIPELINE_IDX:.*]] = arith.addi %[[PIPELINE_IDX]], %[[CONSTANT_1]]
+ // CHECK-DAG: %[[NEXT_LOOP_IDX:.*]] = arith.addi %[[LOOP_IDX]], %[[CONSTANT_1]]
+ // CHECK:     scf.yield {{.*}}, {{.*}}, {{.*}}, %[[NEXT_A_BUFFER]], %[[NEXT_B_BUFFER]], %[[NEXT_A]], %[[NEXT_B]], {{.*}}, {{.*}}, {{.*}}, %[[NEXT_PIPELINE_IDX]], %[[NEXT_LOOP_IDX]]
+-func @matmul_loop_nested(%lb : index, %ub : index, %step : index, 
+-                         %A : !tt.ptr<f16> {tt.divisibility = 16 : i32}, 
++func.func @matmul_loop_nested(%lb : index, %ub : index, %step : index,
++                         %A : !tt.ptr<f16> {tt.divisibility = 16 : i32},
+                          %B : !tt.ptr<f16> {tt.divisibility = 16 : i32}) {
+   scf.for %iv0 = %lb to %ub step %step {
+     // A ptrs
+@@ -134,7 +134,7 @@ func @matmul_loop_nested(%lb : index, %ub : index, %step : index,
+     %b_tmp1 = tt.expand_dims %b_tmp0 {axis = 0 : i32} : (tensor<128xi32, #BLs0>) -> tensor<1x128xi32, #BL>
+     %b_offs = tt.broadcast %b_tmp1 : (tensor<1x128xi32, #BL>) -> tensor<32x128xi32, #BL>
+     %b_ptr_init = tt.addptr %b_ptr_splat, %b_offs : tensor<32x128x!tt.ptr<f16>, #BL>, tensor<32x128xi32, #BL>
+-  
++
+     %a_mask = arith.constant dense<true> : tensor<128x32xi1, #AL>
+     %a_other = arith.constant dense<0.00e+00> : tensor<128x32xf16, #AL>
+     %b_mask = arith.constant dense<true> : tensor<32x128xi1, #BL>
+@@ -161,7 +161,7 @@ func @matmul_loop_nested(%lb : index, %ub : index, %step : index,
+ }
+ 
+ 
+-// CHECK: func @matmul_loop_single_pipeline
++// CHECK: func.func @matmul_loop_single_pipeline
+ // CHECK-DAG: %[[CONSTANT_0:.*]] = arith.constant 0 : i32
+ // CHECK-DAG: %[[CONSTANT_1:.*]] = arith.constant 1 : i32
+ // CHECK-DAG: %[[CONSTANT_2:.*]] = arith.constant 2 : i32
+@@ -183,8 +183,8 @@ func @matmul_loop_nested(%lb : index, %ub : index, %step : index,
+ // CHECK-DAG: %[[NEXT_PIPELINE_IDX:.*]] = arith.addi %[[PIPELINE_IDX]], %[[CONSTANT_1]]
+ // CHECK-DAG: %[[NEXT_LOOP_IDX:.*]] = arith.addi %[[LOOP_IDX]], %[[CONSTANT_1]]
+ // CHECK:   scf.yield {{.*}}, {{.*}}, %[[NEXT_B_BUFFER]], %[[NEXT_B]], {{.*}}, {{.*}}, %[[NEXT_PIPELINE_IDX]], %[[NEXT_LOOP_IDX]]
+-func @matmul_loop_single_pipeline(%lb : index, %ub : index, %step : index, 
+-                                  %A : !tt.ptr<f16> {tt.divisibility = 16 : i32}, 
++func.func @matmul_loop_single_pipeline(%lb : index, %ub : index, %step : index,
++                                  %A : !tt.ptr<f16> {tt.divisibility = 16 : i32},
+                                   %B : !tt.ptr<f16> {tt.divisibility = 16 : i32}) {
+   // A ptrs
+   %a_ptr_splat = tt.splat %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL>
+diff --git a/test/TritonGPU/matmul.mlir b/test/TritonGPU/matmul.mlir
+index 9bd5318e1e..01dc3f0ab1 100644
+--- a/test/TritonGPU/matmul.mlir
++++ b/test/TritonGPU/matmul.mlir
+@@ -4,7 +4,7 @@
+ // CHECK: offset = 49152, size = 49152
+ // CHECK: size = 98304
+ module {
+-func @matmul_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32_i32_i32_i32_i32_i32_i32__12c64_13c64_14c64_15c8(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32, %arg4: i32, %arg5: i32, %arg6: i32 {tt.divisibility = 16 : i32}, %arg7: i32, %arg8: i32 {tt.divisibility = 16 : i32}, %arg9: i32, %arg10: i32 {tt.divisibility = 16 : i32}, %arg11: i32) {
++func.func @matmul_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32_i32_i32_i32_i32_i32_i32__12c64_13c64_14c64_15c8(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg3: i32, %arg4: i32, %arg5: i32, %arg6: i32 {tt.divisibility = 16 : i32}, %arg7: i32, %arg8: i32 {tt.divisibility = 16 : i32}, %arg9: i32, %arg10: i32 {tt.divisibility = 16 : i32}, %arg11: i32) {
+     %cst = arith.constant dense<true> : tensor<64x64xi1>
+     %c64 = arith.constant 64 : index
+     %c0 = arith.constant 0 : index
+@@ -22,7 +22,7 @@ func @matmul_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32_i32_i32_i32_i32_i32_i32__12c6
+     %7 = arith.muli %6, %c8_i32 : i32
+     %8 = arith.subi %2, %7 : i32
+     %9 = arith.cmpi slt, %8, %c8_i32 : i32
+-    %10 = select %9, %8, %c8_i32 : i32
++    %10 = arith.select %9, %8, %c8_i32 : i32
+     %11 = arith.remsi %0, %10 : i32
+     %12 = arith.addi %7, %11 : i32
+     %13 = arith.remsi %0, %5 : i32
+diff --git a/test/TritonGPU/prefetch.mlir b/test/TritonGPU/prefetch.mlir
+index 52b4dddec1..b427547890 100644
+--- a/test/TritonGPU/prefetch.mlir
++++ b/test/TritonGPU/prefetch.mlir
+@@ -11,7 +11,7 @@
+ #B_OP = #triton_gpu.dot_op<{opIdx = 1, parent = #C}>
+ 
+ 
+-// CHECK: func @matmul_loop
++// CHECK: func.func @matmul_loop
+ // CHECK-DAG: %[[A0_PREFETCH_SMEM:.*]] = tensor.extract_slice %[[A0:.*]][0, 0] [128, 16]
+ // CHECK-DAG: %[[A0_PREFETCH:.*]] = triton_gpu.convert_layout %[[A0_PREFETCH_SMEM]]
+ // CHECK-DAG: %[[B0_PREFETCH_SMEM:.*]] = tensor.extract_slice %[[B0:.*]][0, 0] [16, 128]
+@@ -28,7 +28,7 @@
+ // CHECK-DAG:   %[[NEXT_B_PREFETCH_SMEM:.*]] = tensor.extract_slice {{.*}}[0, 0] [16, 128]
+ // CHECK-DAG:   %[[NEXT_B_PREFETCH:.*]] = triton_gpu.convert_layout %[[NEXT_B_PREFETCH_SMEM]]
+ // CHECK:     scf.yield {{.*}}, {{.*}}, {{.*}}, {{.*}}, {{.*}}, %[[NEXT_A_PREFETCH]], %[[NEXT_B_PREFETCH]]
+-func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
++func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) {
+   %a_ptr_init = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL>
+   %b_ptr_init = tt.broadcast %B : (!tt.ptr<f16>) -> tensor<32x128x!tt.ptr<f16>, #BL>
+ 
+diff --git a/test/TritonGPU/update-mma-for-volta.mlir b/test/TritonGPU/update-mma-for-volta.mlir
+index d587fffcca..7571ec6185 100644
+--- a/test/TritonGPU/update-mma-for-volta.mlir
++++ b/test/TritonGPU/update-mma-for-volta.mlir
+@@ -15,7 +15,7 @@
+ // CHECK: [[new_mma:#mma.*]] = #triton_gpu.mma<{versionMajor = 1, versionMinor = 3, warpsPerCTA = [4, 2]}>
+ module attributes {"triton_gpu.num-warps" = 16 : i32} {
+   // CHECK-LABEL: dot_mmav1
+-  func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> {
++  func.func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> {
+     %C = arith.constant dense<0.000000e+00> : tensor<64x64xf32, #blocked0>
+     %AA = triton_gpu.convert_layout %A : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_a>
+     %BB = triton_gpu.convert_layout %B : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_b>
+@@ -50,7 +50,7 @@ module attributes {"triton_gpu.num-warps" = 16 : i32} {
+ 
+ module attributes {"triton_gpu.num-warps" = 16 : i32} {
+   // CHECK-LABEL: dot_mmav1
+-  func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> {
++  func.func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> {
+     %C = arith.constant dense<0.000000e+00> : tensor<64x64xf32, #blocked0>
+     %AA = triton_gpu.convert_layout %A : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_a>
+     %BB = triton_gpu.convert_layout %B : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_b>
+diff --git a/test/lib/Analysis/TestAlias.cpp b/test/lib/Analysis/TestAlias.cpp
+index 88a4118fe9..3fd0cfd0d3 100644
+--- a/test/lib/Analysis/TestAlias.cpp
++++ b/test/lib/Analysis/TestAlias.cpp
+@@ -9,10 +9,10 @@ using namespace mlir;
+ namespace {
+ 
+ struct TestAliasPass
+-    : public PassWrapper<TestAliasPass, OperationPass<FuncOp>> {
++    : public PassWrapper<TestAliasPass, OperationPass<func::FuncOp>> {
++
++  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAliasPass);
+ 
+-  // LLVM15+
+-  // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAliasPass);
+   static void print(StringRef name, SmallVector<std::string, 4> &vals,
+                     raw_ostream &os) {
+     if (vals.empty())
+@@ -39,23 +39,24 @@ struct TestAliasPass
+     auto opName = SymbolTable::getSymbolName(operation).getValue().str();
+     os << opName << "\n";
+ 
+-    SharedMemoryAliasAnalysis analysis(&getContext());
+-    analysis.run(operation);
++    std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver();
++    SharedMemoryAliasAnalysis *analysis =
++        solver->load<SharedMemoryAliasAnalysis>();
++    if (failed(solver->initializeAndRun(operation)))
++      return signalPassFailure();
+ 
+     AsmState state(operation->getParentOfType<ModuleOp>());
+     // Get operation ids of value's aliases
+     auto getAllocOpNames = [&](Value value) {
+-      LatticeElement<AliasInfo> *latticeElement =
+-          analysis.lookupLatticeElement(value);
++      dataflow::Lattice<AliasInfo> *latticeElement =
++          analysis->getLatticeElement(value);
+       SmallVector<std::string, 4> opNames;
+-      if (latticeElement) {
++      if (latticeElement && !latticeElement->isUninitialized()) {
+         auto &info = latticeElement->getValue();
+-        if (!info.getAllocs().empty()) {
+-          for (auto &alias : info.getAllocs()) {
+-            auto opName =
+-                getValueOperandName(alias.getDefiningOp()->getResult(0), state);
+-            opNames.push_back(std::move(opName));
+-          }
++        for (auto &alias : info.getAllocs()) {
++          auto opName =
++              getValueOperandName(alias.getDefiningOp()->getResult(0), state);
++          opNames.push_back(std::move(opName));
+         }
+       }
+       // Ensure deterministic output
+diff --git a/test/lib/Analysis/TestAllocation.cpp b/test/lib/Analysis/TestAllocation.cpp
+index 84108c4d36..35e42242bd 100644
+--- a/test/lib/Analysis/TestAllocation.cpp
++++ b/test/lib/Analysis/TestAllocation.cpp
+@@ -6,10 +6,9 @@ using namespace mlir;
+ namespace {
+ 
+ struct TestAllocationPass
+-    : public PassWrapper<TestAllocationPass, OperationPass<FuncOp>> {
++    : public PassWrapper<TestAllocationPass, OperationPass<func::FuncOp>> {
+ 
+-  // LLVM15+
+-  // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAllocationPass);
++  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAllocationPass);
+ 
+   StringRef getArgument() const final { return "test-print-allocation"; }
+   StringRef getDescription() const final {
+diff --git a/test/lib/Analysis/TestAxisInfo.cpp b/test/lib/Analysis/TestAxisInfo.cpp
+index a5205bb0a0..22347c32f0 100644
+--- a/test/lib/Analysis/TestAxisInfo.cpp
++++ b/test/lib/Analysis/TestAxisInfo.cpp
+@@ -1,25 +1,15 @@
+ #include "mlir/Pass/Pass.h"
+ #include "triton/Analysis/AxisInfo.h"
++#include "triton/Analysis/Utility.h"
+ 
+ using namespace mlir;
+ 
+ namespace {
+ 
+ struct TestAxisInfoPass
+-    : public PassWrapper<TestAxisInfoPass, OperationPass<FuncOp>> {
++    : public PassWrapper<TestAxisInfoPass, OperationPass<func::FuncOp>> {
+ 
+-  // LLVM15+
+-  // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAlignmentPass);
+-
+-  void print(const std::string &name, raw_ostream &os, ArrayRef<int64_t> vals) {
+-    os << name << ": [";
+-    for (size_t d = 0; d < vals.size(); d++) {
+-      if (d != 0)
+-        os << ", ";
+-      os << vals[d];
+-    }
+-    os << "]";
+-  }
++  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAxisInfoPass);
+ 
+   StringRef getArgument() const final { return "test-print-alignment"; }
+   StringRef getDescription() const final {
+@@ -30,38 +20,19 @@ struct TestAxisInfoPass
+     Operation *operation = getOperation();
+     auto &os = llvm::errs();
+     auto opName = SymbolTable::getSymbolName(operation).getValue().str();
+-    os << opName << "\n";
+-    AxisInfoAnalysis analysis(&getContext());
+-    analysis.run(operation);
++    os << "@" << opName << "\n";
++
++    std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver();
++    AxisInfoAnalysis *analysis = solver->load<AxisInfoAnalysis>();
++    if (failed(solver->initializeAndRun(operation)))
++      return signalPassFailure();
+     operation->walk([&](Operation *op) {
+       if (op->getNumResults() < 1)
+         return;
+       for (Value result : op->getResults()) {
+-        // std::ostringstream oss;
+-        // result.print(oss);
+-        // os << " => ";
+-        LatticeElement<AxisInfo> *latticeElement =
+-            analysis.lookupLatticeElement(result);
+-        if (!latticeElement) {
+-          os << "None\n";
+-          return;
+-        }
+-        AxisInfo &info = latticeElement->getValue();
+-        print("Contiguity", os, info.getContiguity());
+-        os << " ; ";
+-        print("Divisibility", os, info.getDivisibility());
+-        os << " ; ";
+-        print("Constancy", os, info.getConstancy());
+-        os << " ; ";
+-        auto constantValue = info.getConstantValue();
+-        os << "ConstantValue: [";
+-        if (constantValue.has_value())
+-          os << constantValue.value();
+-        else
+-          os << "None";
+-        os << "] ( ";
+         result.print(os);
+-        os << " ) ";
++        os << " => ";
++        analysis->getLatticeElement(result)->getValue().print(os);
+         os << "\n";
+       }
+     });
+diff --git a/test/lib/Analysis/TestMembar.cpp b/test/lib/Analysis/TestMembar.cpp
+index df4279fe24..ab9b9f3fb7 100644
+--- a/test/lib/Analysis/TestMembar.cpp
++++ b/test/lib/Analysis/TestMembar.cpp
+@@ -1,4 +1,4 @@
+-#include "mlir/Dialect/GPU/GPUDialect.h"
++#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+ #include "mlir/IR/Dialect.h"
+ #include "mlir/Pass/Pass.h"
+ #include "triton/Analysis/Allocation.h"
+@@ -9,10 +9,9 @@ using namespace mlir;
+ namespace {
+ 
+ struct TestMembarPass
+-    : public PassWrapper<TestMembarPass, OperationPass<FuncOp>> {
++    : public PassWrapper<TestMembarPass, OperationPass<func::FuncOp>> {
+ 
+-  // LLVM15+
+-  // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestMembarPass);
++  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestMembarPass);
+ 
+   StringRef getArgument() const final { return "test-print-membar"; }
+   StringRef getDescription() const final {
diff --git a/pkgs/development/python-modules/torch/default.nix b/pkgs/development/python-modules/torch/default.nix
index 062fcea4334..d9cf4e5760d 100644
--- a/pkgs/development/python-modules/torch/default.nix
+++ b/pkgs/development/python-modules/torch/default.nix
@@ -6,12 +6,18 @@
 
   # Native build inputs
   cmake, util-linux, linkFarm, symlinkJoin, which, pybind11, removeReferencesTo,
+  pythonRelaxDepsHook,
 
   # Build inputs
   numactl,
   Accelerate, CoreServices, libobjc,
 
   # Propagated build inputs
+  filelock,
+  jinja2,
+  networkx,
+  openai-triton,
+  sympy,
   numpy, pyyaml, cffi, click, typing-extensions,
 
   # Unit tests
@@ -49,9 +55,7 @@ let
   inherit (cudaPackages) cudatoolkit cudaFlags cudnn nccl;
 in
 
-# assert that everything needed for cuda is present and that the correct cuda versions are used
-assert !cudaSupport || (let majorIs = lib.versions.major cudatoolkit.version;
-                        in majorIs == "9" || majorIs == "10" || majorIs == "11");
+assert cudaSupport -> (cudaPackages.cudaMajorVersion == "11");
 
 # confirm that cudatoolkits are sync'd across dependencies
 assert !(MPISupport && cudaSupport) || mpi.cudatoolkit == cudatoolkit;
@@ -129,10 +133,10 @@ let
 in buildPythonPackage rec {
   pname = "torch";
   # Don't forget to update torch-bin to the same version.
-  version = "1.13.1";
+  version = "2.0.0";
   format = "setuptools";
 
-  disabled = pythonOlder "3.7.0";
+  disabled = pythonOlder "3.8.0";
 
   outputs = [
     "out" # output standard python package
@@ -145,7 +149,7 @@ in buildPythonPackage rec {
     repo = "pytorch";
     rev = "refs/tags/v${version}";
     fetchSubmodules = true;
-    hash = "sha256-yQz+xHPw9ODRBkV9hv1th38ZmUr/fXa+K+d+cvmX3Z8=";
+    hash = "sha256-cSw7+AYBUcZLz3UyK/+JWWjQxKwVBXcFvBq0XAcL3tE=";
   };
 
   patches = lib.optionals (stdenv.isDarwin && stdenv.isx86_64) [
@@ -155,15 +159,6 @@ in buildPythonPackage rec {
     # base is 10.12. Until we upgrade, we can fall back on the older
     # pthread support.
     ./pthreadpool-disable-gcd.diff
-  ] ++ [
-    # PyTorch fails to build on gcc 12 due to gloo
-    # https://github.com/pytorch/pytorch/issues/77614
-    (fetchpatch {
-      url = "https://github.com/facebookincubator/gloo/commit/4a5e339b764261d20fc409071dc7a8b8989aa195.patch";
-      stripLen = 1;
-      extraPrefix = "third_party/gloo/";
-      hash = "sha256-UxR1r7F6g76BWj3GBIrSy5t+YZDCWy6mMddwx+hon5w=";
-    })
   ];
 
   postPatch = lib.optionalString rocmSupport ''
@@ -261,7 +256,16 @@ in buildPythonPackage rec {
   # Suppress gcc regression: avx512 math function raises uninitialized variable warning
   # https://gcc.gnu.org/bugzilla/show_bug.cgi?id=105593
   # See also: Fails to compile with GCC 12.1.0 https://github.com/pytorch/pytorch/issues/77939
-  ++ lib.optionals stdenv.cc.isGNU [ "-Wno-error=maybe-uninitialized" "-Wno-error=uninitialized" ]));
+  ++ lib.optionals (stdenv.cc.isGNU && lib.versionAtLeast stdenv.cc.version "12.0.0") [
+    "-Wno-error=maybe-uninitialized"
+    "-Wno-error=uninitialized"
+  ]
+  # Since pytorch 2.0:
+  # gcc-12.2.0/include/c++/12.2.0/bits/new_allocator.h:158:33: error: ‘void operator delete(void*, std::size_t)’
+  # ... called on pointer ‘<unknown>’ with nonzero offset [1, 9223372036854775800] [-Werror=free-nonheap-object]
+  ++ lib.optionals (stdenv.cc.isGNU && lib.versions.major stdenv.cc.version == "12" ) [
+    "-Wno-error=free-nonheap-object"
+  ]));
 
   nativeBuildInputs = [
     cmake
@@ -269,6 +273,7 @@ in buildPythonPackage rec {
     which
     ninja
     pybind11
+    pythonRelaxDepsHook
     removeReferencesTo
   ] ++ lib.optionals cudaSupport [ cudatoolkit_joined ]
     ++ lib.optionals rocmSupport [ rocmtoolkit_joined ];
@@ -286,11 +291,27 @@ in buildPythonPackage rec {
     click
     numpy
     pyyaml
+
+    # From install_requires:
+    filelock
     typing-extensions
+    sympy
+    networkx
+    jinja2
+
     # the following are required for tensorboard support
     pillow six future tensorboard protobuf
-  ] ++ lib.optionals MPISupport [ mpi ]
-    ++ lib.optionals rocmSupport [ rocmtoolkit_joined ];
+  ]
+  ++ lib.optionals MPISupport [ mpi ]
+  ++ lib.optionals rocmSupport [ rocmtoolkit_joined ]
+  # rocm build requires openai-triton;
+  # openai-triton currently requires cuda_nvcc,
+  # so not including it in the cpu-only build;
+  # torch.compile relies on openai-triton,
+  # so we include it for the cuda build as well
+  ++ lib.optionals (rocmSupport || cudaSupport) [
+    openai-triton
+  ];
 
   # Tests take a long time and may be flaky, so just sanity-check imports
   doCheck = false;
@@ -318,6 +339,11 @@ in buildPythonPackage rec {
     "runHook postCheck"
   ];
 
+  pythonRemoveDeps = [
+    # In our dist-info the name is just "triton"
+    "pytorch-triton-rocm"
+  ];
+
   postInstall = ''
     find "$out/${python.sitePackages}/torch/include" "$out/${python.sitePackages}/torch/lib" -type f -exec remove-references-to -t ${stdenv.cc} '{}' +
 
diff --git a/pkgs/development/python-modules/torchinfo/default.nix b/pkgs/development/python-modules/torchinfo/default.nix
index c36372235e7..420a5fd8dfc 100644
--- a/pkgs/development/python-modules/torchinfo/default.nix
+++ b/pkgs/development/python-modules/torchinfo/default.nix
@@ -9,7 +9,7 @@
 
 buildPythonPackage rec {
   pname = "torchinfo";
-  version = "1.64";
+  version = "1.7.2";
   format = "setuptools";
 
   disabled = pythonOlder "3.7";
@@ -18,7 +18,7 @@ buildPythonPackage rec {
     owner = "TylerYep";
     repo = pname;
     rev = "refs/tags/v${version}";
-    hash = "sha256-gcl8RxCD017FP4LtB60WVtOh7jg2Otv/vNd9hKneEAU=";
+    hash = "sha256-O+I7BNQ5moV/ZcbbuP/IFoi0LO0WsGHBbSfgPmFu1Ec=";
   };
 
   propagatedBuildInputs = [
@@ -37,8 +37,13 @@ buildPythonPackage rec {
     "test_google"
   ];
 
+  disabledTestPaths = [
+    # Wants "compressai", which we don't package (2023-03-23)
+    "tests/torchinfo_xl_test.py"
+  ];
+
   pythonImportsCheck = [
-    "torchvision"
+    "torchinfo"
   ];
 
   meta = with lib; {
diff --git a/pkgs/tools/audio/tts/default.nix b/pkgs/tools/audio/tts/default.nix
index 8b279a106fa..5e8d4240e14 100644
--- a/pkgs/tools/audio/tts/default.nix
+++ b/pkgs/tools/audio/tts/default.nix
@@ -188,6 +188,8 @@ python.pkgs.buildPythonApplication rec {
     "tests/vocoder_tests/test_multiband_melgan_train.py"
     "tests/vocoder_tests/test_melgan_train.py"
     "tests/vocoder_tests/test_wavernn_train.py"
+    # only a feed forward test, but still takes too long
+    "tests/tts_tests/test_overflow.py"
   ];
 
   passthru = {
diff --git a/pkgs/tools/misc/partition-manager/default.nix b/pkgs/tools/misc/partition-manager/default.nix
index 0d566eb11b2..4587830e6cd 100644
--- a/pkgs/tools/misc/partition-manager/default.nix
+++ b/pkgs/tools/misc/partition-manager/default.nix
@@ -62,6 +62,13 @@ in mkDerivation rec {
 
   meta = with lib; {
     description = "KDE Partition Manager";
+    longDescription = ''
+      KDE Partition Manager is a utility to help you manage the disks, partitions, and file systems on your computer.
+      It allows you to easily create, copy, move, delete, back up, restore, and resize them without losing data.
+      It supports a large number of file systems, including ext2/3/4, btrfs, reiserfs, NTFS, FAT16/32, JFS, XFS and more.
+
+      To install on NixOS, use the option `programs.partition-manager.enable = true`.
+    '';
     license = with licenses; [ cc-by-40 cc0 gpl3Plus lgpl3Plus mit ];
     homepage = "https://www.kde.org/applications/system/kdepartitionmanager/";
     maintainers = with maintainers; [ peterhoeg oxalica ];
diff --git a/pkgs/tools/networking/ntopng/default.nix b/pkgs/tools/networking/ntopng/default.nix
index d5700c42534..c1feb3126d7 100644
--- a/pkgs/tools/networking/ntopng/default.nix
+++ b/pkgs/tools/networking/ntopng/default.nix
@@ -1,38 +1,77 @@
-{ lib, stdenv, fetchFromGitHub, fetchpatch, pkg-config, autoreconfHook
-, zeromq, ndpi, json_c, openssl, libpcap, libcap, curl, libmaxminddb
-, rrdtool, sqlite, libmysqlclient, expat, net-snmp
+{ lib
+, stdenv
+, autoreconfHook
+, curl
+, expat
+, fetchFromGitHub
+, git
+, json_c
+, libcap
+, libmaxminddb
+, libmysqlclient
+, libpcap
+, libsodium
+, ndpi
+, net-snmp
+, openssl
+, pkg-config
+, rdkafka
+, gtest
+, rrdtool
+, hiredis
+, sqlite
+, which
+, zeromq
 }:
 
 stdenv.mkDerivation rec {
   pname = "ntopng";
-  version = "5.2.1";
+  version = "5.6";
 
   src = fetchFromGitHub {
     owner = "ntop";
     repo = "ntopng";
-    rev = version;
-    sha256 = "sha256-FeRERSq8F3HEelUCkA6pgNNcP94xrWy6EbJgk+cEdqc=";
+    rev = "refs/tags/${version}";
+    hash = "sha256-iGqrS0AneKYwGMEpbKy9if8bnaEu6aEV+QaH+JrF9xs=";
   };
 
-  patches = [
-    (fetchpatch {
-      url = "https://github.com/ntop/ntopng/commit/0aa580e1a45f248fffe6d11729ce40571f08e187.patch";
-      sha256 = "sha256-xqEVwfGgkNS+akbJnLZsVvEQdp9GxxUen8VkFomtcPI=";
-    })
+  preConfigure = ''
+    substituteInPlace Makefile.in \
+      --replace "/bin/rm" "rm"
+  '';
+
+  nativeBuildInputs = [
+    autoreconfHook
+    git
+    pkg-config
+    which
   ];
 
-  nativeBuildInputs = [ autoreconfHook pkg-config ];
-
   buildInputs = [
-    zeromq ndpi json_c openssl libpcap curl libmaxminddb rrdtool sqlite
-    libmysqlclient expat net-snmp libcap
+    curl
+    expat
+    json_c
+    libcap
+    libmaxminddb
+    libmysqlclient
+    libpcap
+    gtest
+    hiredis
+    libsodium
+    net-snmp
+    openssl
+    rdkafka
+    rrdtool
+    sqlite
+    zeromq
   ];
 
   autoreconfPhase = "bash autogen.sh";
 
-  preConfigure = ''
-    substituteInPlace Makefile.in --replace "/bin/rm" "rm"
-  '';
+  configureFlags = [
+    "--with-ndpi-includes=${ndpi}/include/ndpi"
+    "--with-ndpi-static-lib=${ndpi}/lib/"
+  ];
 
   preBuild = ''
     sed -e "s|\(#define CONST_BIN_DIR \).*|\1\"$out/bin\"|g" \
@@ -44,9 +83,10 @@ stdenv.mkDerivation rec {
 
   meta = with lib; {
     description = "High-speed web-based traffic analysis and flow collection tool";
-    homepage = "http://www.ntop.org/products/ntop/";
+    homepage = "https://www.ntop.org/products/traffic-analysis/ntop/";
+    changelog = "https://github.com/ntop/ntopng/blob/${version}/CHANGELOG.md";
     license = licenses.gpl3Plus;
     platforms = platforms.linux ++ platforms.darwin;
-    maintainers = [ maintainers.bjornfor ];
+    maintainers = with maintainers; [ bjornfor ];
   };
 }
diff --git a/pkgs/tools/text/goawk/default.nix b/pkgs/tools/text/goawk/default.nix
index a9d71f1faae..a3ecc22b28d 100644
--- a/pkgs/tools/text/goawk/default.nix
+++ b/pkgs/tools/text/goawk/default.nix
@@ -2,13 +2,13 @@
 
 buildGoModule rec {
   pname = "goawk";
-  version = "1.21.0";
+  version = "1.22.0";
 
   src = fetchFromGitHub {
     owner = "benhoyt";
     repo = "goawk";
     rev = "v${version}";
-    sha256 = "sha256-I6KmNPFD8kkYDyek8lR1ZS7biPA/LYGwJqMoA2fG7Wg=";
+    sha256 = "sha256-8UoYGAmYmC0hcxLfkNac3flKyPBT/xsykuy+TvVosBI=";
   };
 
   vendorSha256 = null;
diff --git a/pkgs/top-level/python-packages.nix b/pkgs/top-level/python-packages.nix
index 06cadd33694..7f48a398157 100644
--- a/pkgs/top-level/python-packages.nix
+++ b/pkgs/top-level/python-packages.nix
@@ -6809,6 +6809,8 @@ self: super: with self; {
 
   open-meteo = callPackage ../development/python-modules/open-meteo { };
 
+  openai-triton = callPackage ../development/python-modules/openai-triton { llvmPackages = pkgs.llvmPackages_rocm; };
+
   openai-whisper = callPackage ../development/python-modules/openai-whisper { };
 
   openant = callPackage ../development/python-modules/openant { };
@@ -11864,6 +11866,7 @@ self: super: with self; {
   torchWithCuda = self.torch.override {
     magma = pkgs.magma-cuda;
     cudaSupport = true;
+    rocmSupport = false;
   };
 
   torchWithoutCuda = self.torch.override {
@@ -11873,6 +11876,7 @@ self: super: with self; {
   torchWithRocm = self.torch.override {
     magma = pkgs.magma-hip;
     rocmSupport = true;
+    cudaSupport = false;
   };
 
   torchWithoutRocm = self.torch.override {