Merge master into staging-next

authored by github-actions[bot] and committed by GitHub 97afa6ec df6db8c5

+10334 -176
+5 -5
nixos/modules/profiles/installation-device.nix
··· 52 52 services.getty.helpLine = '' 53 53 The "nixos" and "root" accounts have empty passwords. 54 54 55 - An ssh daemon is running. You then must set a password 56 - for either "root" or "nixos" with `passwd` or add an ssh key 57 - to /home/nixos/.ssh/authorized_keys be able to login. 55 + To log in over ssh you must set a password for either "nixos" or "root" 56 + with `passwd` (prefix with `sudo` for "root"), or add your public key to 57 + /home/nixos/.ssh/authorized_keys or /root/.ssh/authorized_keys. 58 58 59 59 If you need a wireless connection, type 60 60 `sudo systemctl start wpa_supplicant` and configure a ··· 65 65 start the graphical user interface. 66 66 ''; 67 67 68 - # We run sshd by default. Login via root is only possible after adding a 69 - # password via "passwd" or by adding a ssh key to /home/nixos/.ssh/authorized_keys. 68 + # We run sshd by default. Login is only possible after adding a 69 + # password via "passwd" or by adding a ssh key to ~/.ssh/authorized_keys. 70 70 # The latter one is particular useful if keys are manually added to 71 71 # installation device for head-less systems i.e. arm boards by manually 72 72 # mounting the storage in a different system.
+3 -1
nixos/modules/system/boot/resolved.nix
··· 16 16 default = false; 17 17 type = types.bool; 18 18 description = lib.mdDoc '' 19 - Whether to enable the systemd DNS resolver daemon. 19 + Whether to enable the systemd DNS resolver daemon, `systemd-resolved`. 20 + 21 + Search for `services.resolved` to see all options. 20 22 ''; 21 23 }; 22 24
+2 -2
pkgs/applications/audio/tauon/default.nix
··· 25 25 26 26 stdenv.mkDerivation rec { 27 27 pname = "tauon"; 28 - version = "7.6.2"; 28 + version = "7.6.3"; 29 29 30 30 src = fetchFromGitHub { 31 31 owner = "Taiko2k"; 32 32 repo = "TauonMusicBox"; 33 33 rev = "v${version}"; 34 - hash = "sha256-x/tHCDplC45XEaBaf0aQ0w8AS1SorXtYilJoiOcBDtM="; 34 + hash = "sha256-cNR4Ffn9HvgL5KV4FUSnbzEh6VfoKaIbfpb18/qKEns="; 35 35 }; 36 36 37 37 postUnpack = ''
+1 -1
pkgs/applications/editors/vim/plugins/overrides.nix
··· 706 706 }); 707 707 708 708 noice-nvim = super.noice-nvim.overrideAttrs(old: { 709 - dependencies = with self; [ nui-nvim nvim-notify ]; 709 + dependencies = with self; [ nui-nvim ]; 710 710 }); 711 711 712 712 null-ls-nvim = super.null-ls-nvim.overrideAttrs (old: {
+3 -3
pkgs/applications/misc/binocle/default.nix
··· 16 16 17 17 rustPlatform.buildRustPackage rec { 18 18 pname = "binocle"; 19 - version = "0.3.0"; 19 + version = "0.3.1"; 20 20 21 21 src = fetchFromGitHub { 22 22 owner = "sharkdp"; 23 23 repo = pname; 24 24 rev = "v${version}"; 25 - sha256 = "0b0hf2aq34kxxj0la0yar5sp44k6mqcbyailp6j6q0mksf1l74bc"; 25 + sha256 = "sha256-L4l8Gl7Ok/TWqHjNujPx8Qk3UWebs0SgOQNyBNtpnZo="; 26 26 }; 27 27 28 - cargoSha256 = "sha256-CZWAHWZYaL54Rl6Jrp8B6w6HK+2fIKQle2x4mGHv2/o="; 28 + cargoHash = "sha256-9d0MNQ7jEJKpGbjVtl1XBoOBEVNKDgFouSMrcZ7tXNU="; 29 29 30 30 nativeBuildInputs = [ 31 31 makeWrapper
+2 -2
pkgs/applications/misc/lyx/default.nix
··· 18 18 ''; 19 19 20 20 # LaTeX is used from $PATH, as people often want to have it with extra pkgs 21 - nativeBuildInputs = [ pkg-config makeWrapper ]; 21 + nativeBuildInputs = [ pkg-config makeWrapper python3 ]; 22 22 buildInputs = [ 23 - qtbase qtsvg python3 file/*for libmagic*/ bc 23 + qtbase qtsvg file/*for libmagic*/ bc 24 24 hunspell # enchant 25 25 ]; 26 26
+9 -9
pkgs/applications/networking/browsers/chromium/upstream-info.json
··· 19 19 } 20 20 }, 21 21 "beta": { 22 - "version": "112.0.5615.49", 23 - "sha256": "0hgzbbmz40235binfn3vkkxzvwxilaxg04dclqrz980z7hvkgzfx", 24 - "sha256bin64": "146gd9csj08d1ygwwh6gyqqbi7d34mhv3vv7wv4a8z9hn7jhdifg", 22 + "version": "113.0.5672.24", 23 + "sha256": "1z7yv5lqi1n4ycymkf0kz1v8ig06n430a37ik8hri3a6db8r9lv8", 24 + "sha256bin64": "1y9jaw47dgphqr2l8yc36s7k9lf4qrbmfll1d2d1zdjd5ma3slab", 25 25 "deps": { 26 26 "gn": { 27 - "version": "2023-02-17", 27 + "version": "2023-03-18", 28 28 "url": "https://gn.googlesource.com/gn", 29 - "rev": "b25a2f8c2d33f02082f0f258350f5e22c0973108", 30 - "sha256": "075p4jwk1apvwmqmvhwfw5f669ci7nxwjq9mz5aa2g5lz4fkdm4c" 29 + "rev": "41fef642de70ecdcaaa26be96d56a0398f95abd4", 30 + "sha256": "12w4g2dl58283allclpi1c4i6ih9v2xvdb9hpbmfda12v8lizmlq" 31 31 } 32 32 } 33 33 }, 34 34 "dev": { 35 - "version": "113.0.5672.24", 36 - "sha256": "1z7yv5lqi1n4ycymkf0kz1v8ig06n430a37ik8hri3a6db8r9lv8", 37 - "sha256bin64": "0byksvk781gmh5fmjmc77jh19gvkzadf78yr9b4c42las44g4pn4", 35 + "version": "114.0.5696.0", 36 + "sha256": "0hcf971azy3jjsl211qk53b6nk0f4pzf2dwfv3rjh4f6fa3nbwai", 37 + "sha256bin64": "1ssj633vdg3ghd87az28q262ly1fk7rc3k70l1531xvj2030ijrl", 38 38 "deps": { 39 39 "gn": { 40 40 "version": "2023-03-18",
+5193
pkgs/applications/networking/remote/rustdesk/Cargo.lock
··· 1 + # This file is automatically @generated by Cargo. 2 + # It is not intended for manual editing. 3 + version = 3 4 + 5 + [[package]] 6 + name = "addr2line" 7 + version = "0.19.0" 8 + source = "registry+https://github.com/rust-lang/crates.io-index" 9 + checksum = "a76fd60b23679b7d19bd066031410fb7e458ccc5e958eb5c325888ce4baedc97" 10 + dependencies = [ 11 + "gimli", 12 + ] 13 + 14 + [[package]] 15 + name = "adler" 16 + version = "1.0.2" 17 + source = "registry+https://github.com/rust-lang/crates.io-index" 18 + checksum = "f26201604c87b1e01bd3d98f8d5d9a8fcbb815e8cedb41ffccbeb4bf593a35fe" 19 + 20 + [[package]] 21 + name = "adler32" 22 + version = "1.2.0" 23 + source = "registry+https://github.com/rust-lang/crates.io-index" 24 + checksum = "aae1277d39aeec15cb388266ecc24b11c80469deae6067e17a1a7aa9e5c1f234" 25 + 26 + [[package]] 27 + name = "aho-corasick" 28 + version = "0.7.20" 29 + source = "registry+https://github.com/rust-lang/crates.io-index" 30 + checksum = "cc936419f96fa211c1b9166887b38e5e40b19958e5b895be7c1f93adec7071ac" 31 + dependencies = [ 32 + "memchr", 33 + ] 34 + 35 + [[package]] 36 + name = "alsa" 37 + version = "0.6.0" 38 + source = "registry+https://github.com/rust-lang/crates.io-index" 39 + checksum = "5915f52fe2cf65e83924d037b6c5290b7cee097c6b5c8700746e6168a343fd6b" 40 + dependencies = [ 41 + "alsa-sys", 42 + "bitflags", 43 + "libc", 44 + "nix 0.23.2", 45 + ] 46 + 47 + [[package]] 48 + name = "alsa-sys" 49 + version = "0.3.1" 50 + source = "registry+https://github.com/rust-lang/crates.io-index" 51 + checksum = "db8fee663d06c4e303404ef5f40488a53e062f89ba8bfed81f42325aafad1527" 52 + dependencies = [ 53 + "libc", 54 + "pkg-config", 55 + ] 56 + 57 + [[package]] 58 + name = "android_log-sys" 59 + version = "0.2.0" 60 + source = "registry+https://github.com/rust-lang/crates.io-index" 61 + checksum = "85965b6739a430150bdd138e2374a98af0c3ee0d030b3bb7fc3bddff58d0102e" 62 + 63 + [[package]] 64 + name = "android_logger" 65 + version = "0.11.3" 66 + source = "registry+https://github.com/rust-lang/crates.io-index" 67 + checksum = "8619b80c242aa7bd638b5c7ddd952addeecb71f69c75e33f1d47b2804f8f883a" 68 + dependencies = [ 69 + "android_log-sys", 70 + "env_logger 0.10.0", 71 + "log", 72 + "once_cell", 73 + ] 74 + 75 + [[package]] 76 + name = "android_system_properties" 77 + version = "0.1.5" 78 + source = "registry+https://github.com/rust-lang/crates.io-index" 79 + checksum = "819e7219dbd41043ac279b19830f2efc897156490d7fd6ea916720117ee66311" 80 + dependencies = [ 81 + "libc", 82 + ] 83 + 84 + [[package]] 85 + name = "ansi_term" 86 + version = "0.12.1" 87 + source = "registry+https://github.com/rust-lang/crates.io-index" 88 + checksum = "d52a9bb7ec0cf484c551830a7ce27bd20d67eac647e1befb56b0be4ee39a55d2" 89 + dependencies = [ 90 + "winapi 0.3.9", 91 + ] 92 + 93 + [[package]] 94 + name = "anyhow" 95 + version = "1.0.70" 96 + source = "registry+https://github.com/rust-lang/crates.io-index" 97 + checksum = "7de8ce5e0f9f8d88245311066a578d72b7af3e7088f32783804676302df237e4" 98 + 99 + [[package]] 100 + name = "arboard" 101 + version = "2.1.1" 102 + source = "registry+https://github.com/rust-lang/crates.io-index" 103 + checksum = "dc120354d1b5ec6d7aaf4876b602def75595937b5e15d356eb554ab5177e08bb" 104 + dependencies = [ 105 + "clipboard-win", 106 + "core-graphics", 107 + "image", 108 + "log", 109 + "objc", 110 + "objc-foundation", 111 + "objc_id", 112 + "parking_lot 0.12.1", 113 + "thiserror", 114 + "winapi 0.3.9", 115 + "x11rb", 116 + ] 117 + 118 + [[package]] 119 + name = "async-trait" 120 + version = "0.1.68" 121 + source = "registry+https://github.com/rust-lang/crates.io-index" 122 + checksum = "b9ccdd8f2a161be9bd5c023df56f1b2a0bd1d83872ae53b71a84a12c9bf6e842" 123 + dependencies = [ 124 + "proc-macro2", 125 + "quote", 126 + "syn 2.0.13", 127 + ] 128 + 129 + [[package]] 130 + name = "atk" 131 + version = "0.9.0" 132 + source = "registry+https://github.com/rust-lang/crates.io-index" 133 + checksum = "812b4911e210bd51b24596244523c856ca749e6223c50a7fbbba3f89ee37c426" 134 + dependencies = [ 135 + "atk-sys 0.10.0", 136 + "bitflags", 137 + "glib 0.10.3", 138 + "glib-sys 0.10.1", 139 + "gobject-sys 0.10.0", 140 + "libc", 141 + ] 142 + 143 + [[package]] 144 + name = "atk" 145 + version = "0.15.1" 146 + source = "registry+https://github.com/rust-lang/crates.io-index" 147 + checksum = "2c3d816ce6f0e2909a96830d6911c2aff044370b1ef92d7f267b43bae5addedd" 148 + dependencies = [ 149 + "atk-sys 0.15.1", 150 + "bitflags", 151 + "glib 0.15.12", 152 + "libc", 153 + ] 154 + 155 + [[package]] 156 + name = "atk-sys" 157 + version = "0.10.0" 158 + source = "registry+https://github.com/rust-lang/crates.io-index" 159 + checksum = "f530e4af131d94cc4fa15c5c9d0348f0ef28bac64ba660b6b2a1cf2605dedfce" 160 + dependencies = [ 161 + "glib-sys 0.10.1", 162 + "gobject-sys 0.10.0", 163 + "libc", 164 + "system-deps 1.3.2", 165 + ] 166 + 167 + [[package]] 168 + name = "atk-sys" 169 + version = "0.15.1" 170 + source = "registry+https://github.com/rust-lang/crates.io-index" 171 + checksum = "58aeb089fb698e06db8089971c7ee317ab9644bade33383f63631437b03aafb6" 172 + dependencies = [ 173 + "glib-sys 0.15.10", 174 + "gobject-sys 0.15.10", 175 + "libc", 176 + "system-deps 6.0.4", 177 + ] 178 + 179 + [[package]] 180 + name = "atty" 181 + version = "0.2.14" 182 + source = "registry+https://github.com/rust-lang/crates.io-index" 183 + checksum = "d9b39be18770d11421cdb1b9947a45dd3f37e93092cbf377614828a319d5fee8" 184 + dependencies = [ 185 + "hermit-abi 0.1.19", 186 + "libc", 187 + "winapi 0.3.9", 188 + ] 189 + 190 + [[package]] 191 + name = "autocfg" 192 + version = "0.1.8" 193 + source = "registry+https://github.com/rust-lang/crates.io-index" 194 + checksum = "0dde43e75fd43e8a1bf86103336bc699aa8d17ad1be60c76c0bdfd4828e19b78" 195 + dependencies = [ 196 + "autocfg 1.1.0", 197 + ] 198 + 199 + [[package]] 200 + name = "autocfg" 201 + version = "1.1.0" 202 + source = "registry+https://github.com/rust-lang/crates.io-index" 203 + checksum = "d468802bab17cbc0cc575e9b053f41e72aa36bfa6b7f55e3529ffa43161b97fa" 204 + 205 + [[package]] 206 + name = "backtrace" 207 + version = "0.3.67" 208 + source = "registry+https://github.com/rust-lang/crates.io-index" 209 + checksum = "233d376d6d185f2a3093e58f283f60f880315b6c60075b01f36b3b85154564ca" 210 + dependencies = [ 211 + "addr2line", 212 + "cc", 213 + "cfg-if 1.0.0", 214 + "libc", 215 + "miniz_oxide 0.6.2", 216 + "object", 217 + "rustc-demangle", 218 + ] 219 + 220 + [[package]] 221 + name = "base64" 222 + version = "0.13.1" 223 + source = "registry+https://github.com/rust-lang/crates.io-index" 224 + checksum = "9e1b586273c5702936fe7b7d6896644d8be71e6314cfe09d3167c95f712589e8" 225 + 226 + [[package]] 227 + name = "base64" 228 + version = "0.21.0" 229 + source = "registry+https://github.com/rust-lang/crates.io-index" 230 + checksum = "a4a4ddaa51a5bc52a6948f74c06d20aaaddb71924eab79b8c97a8c556e942d6a" 231 + 232 + [[package]] 233 + name = "bindgen" 234 + version = "0.59.2" 235 + source = "registry+https://github.com/rust-lang/crates.io-index" 236 + checksum = "2bd2a9a458e8f4304c52c43ebb0cfbd520289f8379a52e329a38afda99bf8eb8" 237 + dependencies = [ 238 + "bitflags", 239 + "cexpr", 240 + "clang-sys", 241 + "clap 2.34.0", 242 + "env_logger 0.9.3", 243 + "lazy_static", 244 + "lazycell", 245 + "log", 246 + "peeking_take_while", 247 + "proc-macro2", 248 + "quote", 249 + "regex", 250 + "rustc-hash", 251 + "shlex", 252 + "which 4.4.0", 253 + ] 254 + 255 + [[package]] 256 + name = "bindgen" 257 + version = "0.64.0" 258 + source = "registry+https://github.com/rust-lang/crates.io-index" 259 + checksum = "c4243e6031260db77ede97ad86c27e501d646a27ab57b59a574f725d98ab1fb4" 260 + dependencies = [ 261 + "bitflags", 262 + "cexpr", 263 + "clang-sys", 264 + "lazy_static", 265 + "lazycell", 266 + "peeking_take_while", 267 + "proc-macro2", 268 + "quote", 269 + "regex", 270 + "rustc-hash", 271 + "shlex", 272 + "syn 1.0.109", 273 + ] 274 + 275 + [[package]] 276 + name = "bitflags" 277 + version = "1.3.2" 278 + source = "registry+https://github.com/rust-lang/crates.io-index" 279 + checksum = "bef38d45163c2f1dde094a7dfd33ccf595c92905c8f8f4fdc18d06fb1037718a" 280 + 281 + [[package]] 282 + name = "block" 283 + version = "0.1.6" 284 + source = "registry+https://github.com/rust-lang/crates.io-index" 285 + checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a" 286 + 287 + [[package]] 288 + name = "block-buffer" 289 + version = "0.10.4" 290 + source = "registry+https://github.com/rust-lang/crates.io-index" 291 + checksum = "3078c7629b62d3f0439517fa394996acacc5cbc91c5a20d8c658e77abd503a71" 292 + dependencies = [ 293 + "generic-array", 294 + ] 295 + 296 + [[package]] 297 + name = "bumpalo" 298 + version = "3.12.0" 299 + source = "registry+https://github.com/rust-lang/crates.io-index" 300 + checksum = "0d261e256854913907f67ed06efbc3338dfe6179796deefc1ff763fc1aee5535" 301 + 302 + [[package]] 303 + name = "bytemuck" 304 + version = "1.13.1" 305 + source = "registry+https://github.com/rust-lang/crates.io-index" 306 + checksum = "17febce684fd15d89027105661fec94afb475cb995fbc59d2865198446ba2eea" 307 + 308 + [[package]] 309 + name = "byteorder" 310 + version = "1.4.3" 311 + source = "registry+https://github.com/rust-lang/crates.io-index" 312 + checksum = "14c189c53d098945499cdfa7ecc63567cf3886b3332b312a5b4585d8d3a6a610" 313 + 314 + [[package]] 315 + name = "bytes" 316 + version = "1.4.0" 317 + source = "registry+https://github.com/rust-lang/crates.io-index" 318 + checksum = "89b2fd2a0dcf38d7971e2194b6b6eebab45ae01067456a7fd93d5547a61b70be" 319 + 320 + [[package]] 321 + name = "cairo-rs" 322 + version = "0.9.1" 323 + source = "registry+https://github.com/rust-lang/crates.io-index" 324 + checksum = "c5c0f2e047e8ca53d0ff249c54ae047931d7a6ebe05d00af73e0ffeb6e34bdb8" 325 + dependencies = [ 326 + "bitflags", 327 + "cairo-sys-rs 0.10.0", 328 + "glib 0.10.3", 329 + "glib-sys 0.10.1", 330 + "gobject-sys 0.10.0", 331 + "libc", 332 + "thiserror", 333 + ] 334 + 335 + [[package]] 336 + name = "cairo-rs" 337 + version = "0.15.12" 338 + source = "registry+https://github.com/rust-lang/crates.io-index" 339 + checksum = "c76ee391b03d35510d9fa917357c7f1855bd9a6659c95a1b392e33f49b3369bc" 340 + dependencies = [ 341 + "bitflags", 342 + "cairo-sys-rs 0.15.1", 343 + "glib 0.15.12", 344 + "libc", 345 + "thiserror", 346 + ] 347 + 348 + [[package]] 349 + name = "cairo-sys-rs" 350 + version = "0.10.0" 351 + source = "registry+https://github.com/rust-lang/crates.io-index" 352 + checksum = "2ed2639b9ad5f1d6efa76de95558e11339e7318426d84ac4890b86c03e828ca7" 353 + dependencies = [ 354 + "glib-sys 0.10.1", 355 + "libc", 356 + "system-deps 1.3.2", 357 + ] 358 + 359 + [[package]] 360 + name = "cairo-sys-rs" 361 + version = "0.15.1" 362 + source = "registry+https://github.com/rust-lang/crates.io-index" 363 + checksum = "3c55d429bef56ac9172d25fecb85dc8068307d17acd74b377866b7a1ef25d3c8" 364 + dependencies = [ 365 + "glib-sys 0.15.10", 366 + "libc", 367 + "system-deps 6.0.4", 368 + ] 369 + 370 + [[package]] 371 + name = "cc" 372 + version = "1.0.79" 373 + source = "registry+https://github.com/rust-lang/crates.io-index" 374 + checksum = "50d30906286121d95be3d479533b458f87493b30a4b5f79a607db8f5d11aa91f" 375 + dependencies = [ 376 + "jobserver", 377 + ] 378 + 379 + [[package]] 380 + name = "cesu8" 381 + version = "1.1.0" 382 + source = "registry+https://github.com/rust-lang/crates.io-index" 383 + checksum = "6d43a04d8753f35258c91f8ec639f792891f748a1edbd759cf1dcea3382ad83c" 384 + 385 + [[package]] 386 + name = "cexpr" 387 + version = "0.6.0" 388 + source = "registry+https://github.com/rust-lang/crates.io-index" 389 + checksum = "6fac387a98bb7c37292057cffc56d62ecb629900026402633ae9160df93a8766" 390 + dependencies = [ 391 + "nom", 392 + ] 393 + 394 + [[package]] 395 + name = "cfg-expr" 396 + version = "0.14.0" 397 + source = "registry+https://github.com/rust-lang/crates.io-index" 398 + checksum = "a35b255461940a32985c627ce82900867c61db1659764d3675ea81963f72a4c6" 399 + dependencies = [ 400 + "smallvec", 401 + ] 402 + 403 + [[package]] 404 + name = "cfg-if" 405 + version = "0.1.10" 406 + source = "registry+https://github.com/rust-lang/crates.io-index" 407 + checksum = "4785bdd1c96b2a846b2bd7cc02e86b6b3dbf14e7e53446c4f54c92a361040822" 408 + 409 + [[package]] 410 + name = "cfg-if" 411 + version = "1.0.0" 412 + source = "registry+https://github.com/rust-lang/crates.io-index" 413 + checksum = "baf1de4339761588bc0619e3cbc0120ee582ebb74b53b4efbf79117bd2da40fd" 414 + 415 + [[package]] 416 + name = "chrono" 417 + version = "0.4.24" 418 + source = "registry+https://github.com/rust-lang/crates.io-index" 419 + checksum = "4e3c5919066adf22df73762e50cffcde3a758f2a848b113b586d1f86728b673b" 420 + dependencies = [ 421 + "iana-time-zone", 422 + "num-integer", 423 + "num-traits 0.2.15", 424 + "winapi 0.3.9", 425 + ] 426 + 427 + [[package]] 428 + name = "clang-sys" 429 + version = "1.6.1" 430 + source = "registry+https://github.com/rust-lang/crates.io-index" 431 + checksum = "c688fc74432808e3eb684cae8830a86be1d66a2bd58e1f248ed0960a590baf6f" 432 + dependencies = [ 433 + "glob", 434 + "libc", 435 + "libloading", 436 + ] 437 + 438 + [[package]] 439 + name = "clap" 440 + version = "2.34.0" 441 + source = "registry+https://github.com/rust-lang/crates.io-index" 442 + checksum = "a0610544180c38b88101fecf2dd634b174a62eef6946f84dfc6a7127512b381c" 443 + dependencies = [ 444 + "ansi_term", 445 + "atty", 446 + "bitflags", 447 + "strsim 0.8.0", 448 + "textwrap 0.11.0", 449 + "unicode-width", 450 + "vec_map", 451 + ] 452 + 453 + [[package]] 454 + name = "clap" 455 + version = "3.2.23" 456 + source = "registry+https://github.com/rust-lang/crates.io-index" 457 + checksum = "71655c45cb9845d3270c9d6df84ebe72b4dad3c2ba3f7023ad47c144e4e473a5" 458 + dependencies = [ 459 + "atty", 460 + "bitflags", 461 + "clap_lex", 462 + "indexmap", 463 + "strsim 0.10.0", 464 + "termcolor", 465 + "textwrap 0.16.0", 466 + ] 467 + 468 + [[package]] 469 + name = "clap_lex" 470 + version = "0.2.4" 471 + source = "registry+https://github.com/rust-lang/crates.io-index" 472 + checksum = "2850f2f5a82cbf437dd5af4d49848fbdfc27c157c3d010345776f952765261c5" 473 + dependencies = [ 474 + "os_str_bytes", 475 + ] 476 + 477 + [[package]] 478 + name = "clipboard" 479 + version = "0.1.0" 480 + dependencies = [ 481 + "cc", 482 + "hbb_common", 483 + "lazy_static", 484 + "serde 1.0.159", 485 + "serde_derive", 486 + "thiserror", 487 + ] 488 + 489 + [[package]] 490 + name = "clipboard-master" 491 + version = "3.1.3" 492 + source = "registry+https://github.com/rust-lang/crates.io-index" 493 + checksum = "459887701008d8ee21f8de7f45f0f0707417c7eea3311973f6e67222bd686b7a" 494 + dependencies = [ 495 + "objc", 496 + "objc-foundation", 497 + "objc_id", 498 + "winapi 0.3.9", 499 + "windows-win", 500 + "x11-clipboard", 501 + ] 502 + 503 + [[package]] 504 + name = "clipboard-win" 505 + version = "4.5.0" 506 + source = "registry+https://github.com/rust-lang/crates.io-index" 507 + checksum = "7191c27c2357d9b7ef96baac1773290d4ca63b24205b82a3fd8a0637afcf0362" 508 + dependencies = [ 509 + "error-code", 510 + "str-buf", 511 + "winapi 0.3.9", 512 + ] 513 + 514 + [[package]] 515 + name = "cloudabi" 516 + version = "0.0.3" 517 + source = "registry+https://github.com/rust-lang/crates.io-index" 518 + checksum = "ddfc5b9aa5d4507acaf872de71051dfd0e309860e88966e1051e462a077aac4f" 519 + dependencies = [ 520 + "bitflags", 521 + ] 522 + 523 + [[package]] 524 + name = "cmake" 525 + version = "0.1.50" 526 + source = "registry+https://github.com/rust-lang/crates.io-index" 527 + checksum = "a31c789563b815f77f4250caee12365734369f942439b7defd71e18a48197130" 528 + dependencies = [ 529 + "cc", 530 + ] 531 + 532 + [[package]] 533 + name = "cocoa" 534 + version = "0.24.1" 535 + source = "registry+https://github.com/rust-lang/crates.io-index" 536 + checksum = "f425db7937052c684daec3bd6375c8abe2d146dca4b8b143d6db777c39138f3a" 537 + dependencies = [ 538 + "bitflags", 539 + "block", 540 + "cocoa-foundation", 541 + "core-foundation", 542 + "core-graphics", 543 + "foreign-types", 544 + "libc", 545 + "objc", 546 + ] 547 + 548 + [[package]] 549 + name = "cocoa-foundation" 550 + version = "0.1.1" 551 + source = "registry+https://github.com/rust-lang/crates.io-index" 552 + checksum = "931d3837c286f56e3c58423ce4eba12d08db2374461a785c86f672b08b5650d6" 553 + dependencies = [ 554 + "bitflags", 555 + "block", 556 + "core-foundation", 557 + "core-graphics-types", 558 + "foreign-types", 559 + "libc", 560 + "objc", 561 + ] 562 + 563 + [[package]] 564 + name = "codespan-reporting" 565 + version = "0.11.1" 566 + source = "registry+https://github.com/rust-lang/crates.io-index" 567 + checksum = "3538270d33cc669650c4b093848450d380def10c331d38c768e34cac80576e6e" 568 + dependencies = [ 569 + "termcolor", 570 + "unicode-width", 571 + ] 572 + 573 + [[package]] 574 + name = "color_quant" 575 + version = "1.1.0" 576 + source = "registry+https://github.com/rust-lang/crates.io-index" 577 + checksum = "3d7b894f5411737b7867f4827955924d7c254fc9f4d91a6aad6b097804b1018b" 578 + 579 + [[package]] 580 + name = "combine" 581 + version = "4.6.6" 582 + source = "registry+https://github.com/rust-lang/crates.io-index" 583 + checksum = "35ed6e9d84f0b51a7f52daf1c7d71dd136fd7a3f41a8462b8cdb8c78d920fad4" 584 + dependencies = [ 585 + "bytes", 586 + "memchr", 587 + ] 588 + 589 + [[package]] 590 + name = "confy" 591 + version = "0.4.0" 592 + source = "git+https://github.com/open-trade/confy#630cc28a396cb7d01eefdd9f3824486fe4d8554b" 593 + dependencies = [ 594 + "directories-next", 595 + "serde 1.0.159", 596 + "thiserror", 597 + "toml 0.5.11", 598 + ] 599 + 600 + [[package]] 601 + name = "core-foundation" 602 + version = "0.9.3" 603 + source = "registry+https://github.com/rust-lang/crates.io-index" 604 + checksum = "194a7a9e6de53fa55116934067c844d9d749312f75c6f6d0980e8c252f8c2146" 605 + dependencies = [ 606 + "core-foundation-sys", 607 + "libc", 608 + ] 609 + 610 + [[package]] 611 + name = "core-foundation-sys" 612 + version = "0.8.4" 613 + source = "registry+https://github.com/rust-lang/crates.io-index" 614 + checksum = "e496a50fda8aacccc86d7529e2c1e0892dbd0f898a6b5645b5561b89c3210efa" 615 + 616 + [[package]] 617 + name = "core-graphics" 618 + version = "0.22.3" 619 + source = "registry+https://github.com/rust-lang/crates.io-index" 620 + checksum = "2581bbab3b8ffc6fcbd550bf46c355135d16e9ff2a6ea032ad6b9bf1d7efe4fb" 621 + dependencies = [ 622 + "bitflags", 623 + "core-foundation", 624 + "core-graphics-types", 625 + "foreign-types", 626 + "libc", 627 + ] 628 + 629 + [[package]] 630 + name = "core-graphics-types" 631 + version = "0.1.1" 632 + source = "registry+https://github.com/rust-lang/crates.io-index" 633 + checksum = "3a68b68b3446082644c91ac778bf50cd4104bfb002b5a6a7c44cca5a2c70788b" 634 + dependencies = [ 635 + "bitflags", 636 + "core-foundation", 637 + "foreign-types", 638 + "libc", 639 + ] 640 + 641 + [[package]] 642 + name = "coreaudio-rs" 643 + version = "0.10.0" 644 + source = "registry+https://github.com/rust-lang/crates.io-index" 645 + checksum = "11894b20ebfe1ff903cbdc52259693389eea03b94918a2def2c30c3bf227ad88" 646 + dependencies = [ 647 + "bitflags", 648 + "coreaudio-sys", 649 + ] 650 + 651 + [[package]] 652 + name = "coreaudio-sys" 653 + version = "0.2.12" 654 + source = "registry+https://github.com/rust-lang/crates.io-index" 655 + checksum = "f034b2258e6c4ade2f73bf87b21047567fb913ee9550837c2316d139b0262b24" 656 + dependencies = [ 657 + "bindgen 0.64.0", 658 + ] 659 + 660 + [[package]] 661 + name = "cpal" 662 + version = "0.13.5" 663 + source = "registry+https://github.com/rust-lang/crates.io-index" 664 + checksum = "74117836a5124f3629e4b474eed03e479abaf98988b4bb317e29f08cfe0e4116" 665 + dependencies = [ 666 + "alsa", 667 + "core-foundation-sys", 668 + "coreaudio-rs", 669 + "jni", 670 + "js-sys", 671 + "lazy_static", 672 + "libc", 673 + "mach", 674 + "ndk", 675 + "ndk-glue", 676 + "nix 0.23.2", 677 + "oboe", 678 + "parking_lot 0.11.2", 679 + "stdweb", 680 + "thiserror", 681 + "web-sys", 682 + "winapi 0.3.9", 683 + ] 684 + 685 + [[package]] 686 + name = "cpufeatures" 687 + version = "0.2.6" 688 + source = "registry+https://github.com/rust-lang/crates.io-index" 689 + checksum = "280a9f2d8b3a38871a3c8a46fb80db65e5e5ed97da80c4d08bf27fb63e35e181" 690 + dependencies = [ 691 + "libc", 692 + ] 693 + 694 + [[package]] 695 + name = "crc32fast" 696 + version = "1.3.2" 697 + source = "registry+https://github.com/rust-lang/crates.io-index" 698 + checksum = "b540bd8bc810d3885c6ea91e2018302f68baba2129ab3e88f32389ee9370880d" 699 + dependencies = [ 700 + "cfg-if 1.0.0", 701 + ] 702 + 703 + [[package]] 704 + name = "crossbeam-channel" 705 + version = "0.5.7" 706 + source = "registry+https://github.com/rust-lang/crates.io-index" 707 + checksum = "cf2b3e8478797446514c91ef04bafcb59faba183e621ad488df88983cc14128c" 708 + dependencies = [ 709 + "cfg-if 1.0.0", 710 + "crossbeam-utils", 711 + ] 712 + 713 + [[package]] 714 + name = "crossbeam-deque" 715 + version = "0.8.3" 716 + source = "registry+https://github.com/rust-lang/crates.io-index" 717 + checksum = "ce6fd6f855243022dcecf8702fef0c297d4338e226845fe067f6341ad9fa0cef" 718 + dependencies = [ 719 + "cfg-if 1.0.0", 720 + "crossbeam-epoch", 721 + "crossbeam-utils", 722 + ] 723 + 724 + [[package]] 725 + name = "crossbeam-epoch" 726 + version = "0.9.14" 727 + source = "registry+https://github.com/rust-lang/crates.io-index" 728 + checksum = "46bd5f3f85273295a9d14aedfb86f6aadbff6d8f5295c4a9edb08e819dcf5695" 729 + dependencies = [ 730 + "autocfg 1.1.0", 731 + "cfg-if 1.0.0", 732 + "crossbeam-utils", 733 + "memoffset 0.8.0", 734 + "scopeguard", 735 + ] 736 + 737 + [[package]] 738 + name = "crossbeam-queue" 739 + version = "0.3.8" 740 + source = "registry+https://github.com/rust-lang/crates.io-index" 741 + checksum = "d1cfb3ea8a53f37c40dea2c7bedcbd88bdfae54f5e2175d6ecaff1c988353add" 742 + dependencies = [ 743 + "cfg-if 1.0.0", 744 + "crossbeam-utils", 745 + ] 746 + 747 + [[package]] 748 + name = "crossbeam-utils" 749 + version = "0.8.15" 750 + source = "registry+https://github.com/rust-lang/crates.io-index" 751 + checksum = "3c063cd8cc95f5c377ed0d4b49a4b21f632396ff690e8470c29b3359b346984b" 752 + dependencies = [ 753 + "cfg-if 1.0.0", 754 + ] 755 + 756 + [[package]] 757 + name = "crypto-common" 758 + version = "0.1.6" 759 + source = "registry+https://github.com/rust-lang/crates.io-index" 760 + checksum = "1bfb12502f3fc46cca1bb51ac28df9d618d813cdc3d2f25b9fe775a34af26bb3" 761 + dependencies = [ 762 + "generic-array", 763 + "typenum", 764 + ] 765 + 766 + [[package]] 767 + name = "ctrlc" 768 + version = "3.2.5" 769 + source = "registry+https://github.com/rust-lang/crates.io-index" 770 + checksum = "bbcf33c2a618cbe41ee43ae6e9f2e48368cd9f9db2896f10167d8d762679f639" 771 + dependencies = [ 772 + "nix 0.26.2", 773 + "windows-sys 0.45.0", 774 + ] 775 + 776 + [[package]] 777 + name = "cxx" 778 + version = "1.0.94" 779 + source = "registry+https://github.com/rust-lang/crates.io-index" 780 + checksum = "f61f1b6389c3fe1c316bf8a4dccc90a38208354b330925bce1f74a6c4756eb93" 781 + dependencies = [ 782 + "cc", 783 + "cxxbridge-flags", 784 + "cxxbridge-macro", 785 + "link-cplusplus", 786 + ] 787 + 788 + [[package]] 789 + name = "cxx-build" 790 + version = "1.0.94" 791 + source = "registry+https://github.com/rust-lang/crates.io-index" 792 + checksum = "12cee708e8962df2aeb38f594aae5d827c022b6460ac71a7a3e2c3c2aae5a07b" 793 + dependencies = [ 794 + "cc", 795 + "codespan-reporting", 796 + "once_cell", 797 + "proc-macro2", 798 + "quote", 799 + "scratch", 800 + "syn 2.0.13", 801 + ] 802 + 803 + [[package]] 804 + name = "cxxbridge-flags" 805 + version = "1.0.94" 806 + source = "registry+https://github.com/rust-lang/crates.io-index" 807 + checksum = "7944172ae7e4068c533afbb984114a56c46e9ccddda550499caa222902c7f7bb" 808 + 809 + [[package]] 810 + name = "cxxbridge-macro" 811 + version = "1.0.94" 812 + source = "registry+https://github.com/rust-lang/crates.io-index" 813 + checksum = "2345488264226bf682893e25de0769f3360aac9957980ec49361b083ddaa5bc5" 814 + dependencies = [ 815 + "proc-macro2", 816 + "quote", 817 + "syn 2.0.13", 818 + ] 819 + 820 + [[package]] 821 + name = "darling" 822 + version = "0.13.4" 823 + source = "registry+https://github.com/rust-lang/crates.io-index" 824 + checksum = "a01d95850c592940db9b8194bc39f4bc0e89dee5c4265e4b1807c34a9aba453c" 825 + dependencies = [ 826 + "darling_core", 827 + "darling_macro", 828 + ] 829 + 830 + [[package]] 831 + name = "darling_core" 832 + version = "0.13.4" 833 + source = "registry+https://github.com/rust-lang/crates.io-index" 834 + checksum = "859d65a907b6852c9361e3185c862aae7fafd2887876799fa55f5f99dc40d610" 835 + dependencies = [ 836 + "fnv", 837 + "ident_case", 838 + "proc-macro2", 839 + "quote", 840 + "strsim 0.10.0", 841 + "syn 1.0.109", 842 + ] 843 + 844 + [[package]] 845 + name = "darling_macro" 846 + version = "0.13.4" 847 + source = "registry+https://github.com/rust-lang/crates.io-index" 848 + checksum = "9c972679f83bdf9c42bd905396b6c3588a843a17f0f16dfcfa3e2c5d57441835" 849 + dependencies = [ 850 + "darling_core", 851 + "quote", 852 + "syn 1.0.109", 853 + ] 854 + 855 + [[package]] 856 + name = "dasp" 857 + version = "0.11.0" 858 + source = "registry+https://github.com/rust-lang/crates.io-index" 859 + checksum = "7381b67da416b639690ac77c73b86a7b5e64a29e31d1f75fb3b1102301ef355a" 860 + dependencies = [ 861 + "dasp_envelope", 862 + "dasp_frame", 863 + "dasp_interpolate", 864 + "dasp_peak", 865 + "dasp_ring_buffer", 866 + "dasp_rms", 867 + "dasp_sample", 868 + "dasp_signal", 869 + "dasp_slice", 870 + "dasp_window", 871 + ] 872 + 873 + [[package]] 874 + name = "dasp_envelope" 875 + version = "0.11.0" 876 + source = "registry+https://github.com/rust-lang/crates.io-index" 877 + checksum = "8ec617ce7016f101a87fe85ed44180839744265fae73bb4aa43e7ece1b7668b6" 878 + dependencies = [ 879 + "dasp_frame", 880 + "dasp_peak", 881 + "dasp_ring_buffer", 882 + "dasp_rms", 883 + "dasp_sample", 884 + ] 885 + 886 + [[package]] 887 + name = "dasp_frame" 888 + version = "0.11.0" 889 + source = "registry+https://github.com/rust-lang/crates.io-index" 890 + checksum = "b2a3937f5fe2135702897535c8d4a5553f8b116f76c1529088797f2eee7c5cd6" 891 + dependencies = [ 892 + "dasp_sample", 893 + ] 894 + 895 + [[package]] 896 + name = "dasp_interpolate" 897 + version = "0.11.0" 898 + source = "registry+https://github.com/rust-lang/crates.io-index" 899 + checksum = "7fc975a6563bb7ca7ec0a6c784ead49983a21c24835b0bc96eea11ee407c7486" 900 + dependencies = [ 901 + "dasp_frame", 902 + "dasp_ring_buffer", 903 + "dasp_sample", 904 + ] 905 + 906 + [[package]] 907 + name = "dasp_peak" 908 + version = "0.11.0" 909 + source = "registry+https://github.com/rust-lang/crates.io-index" 910 + checksum = "5cf88559d79c21f3d8523d91250c397f9a15b5fc72fbb3f87fdb0a37b79915bf" 911 + dependencies = [ 912 + "dasp_frame", 913 + "dasp_sample", 914 + ] 915 + 916 + [[package]] 917 + name = "dasp_ring_buffer" 918 + version = "0.11.0" 919 + source = "registry+https://github.com/rust-lang/crates.io-index" 920 + checksum = "07d79e19b89618a543c4adec9c5a347fe378a19041699b3278e616e387511ea1" 921 + 922 + [[package]] 923 + name = "dasp_rms" 924 + version = "0.11.0" 925 + source = "registry+https://github.com/rust-lang/crates.io-index" 926 + checksum = "a6c5dcb30b7e5014486e2822537ea2beae50b19722ffe2ed7549ab03774575aa" 927 + dependencies = [ 928 + "dasp_frame", 929 + "dasp_ring_buffer", 930 + "dasp_sample", 931 + ] 932 + 933 + [[package]] 934 + name = "dasp_sample" 935 + version = "0.11.0" 936 + source = "registry+https://github.com/rust-lang/crates.io-index" 937 + checksum = "0c87e182de0887fd5361989c677c4e8f5000cd9491d6d563161a8f3a5519fc7f" 938 + 939 + [[package]] 940 + name = "dasp_signal" 941 + version = "0.11.0" 942 + source = "registry+https://github.com/rust-lang/crates.io-index" 943 + checksum = "aa1ab7d01689c6ed4eae3d38fe1cea08cba761573fbd2d592528d55b421077e7" 944 + dependencies = [ 945 + "dasp_envelope", 946 + "dasp_frame", 947 + "dasp_interpolate", 948 + "dasp_peak", 949 + "dasp_ring_buffer", 950 + "dasp_rms", 951 + "dasp_sample", 952 + "dasp_window", 953 + ] 954 + 955 + [[package]] 956 + name = "dasp_slice" 957 + version = "0.11.0" 958 + source = "registry+https://github.com/rust-lang/crates.io-index" 959 + checksum = "4e1c7335d58e7baedafa516cb361360ff38d6f4d3f9d9d5ee2a2fc8e27178fa1" 960 + dependencies = [ 961 + "dasp_frame", 962 + "dasp_sample", 963 + ] 964 + 965 + [[package]] 966 + name = "dasp_window" 967 + version = "0.11.1" 968 + source = "registry+https://github.com/rust-lang/crates.io-index" 969 + checksum = "99ded7b88821d2ce4e8b842c9f1c86ac911891ab89443cc1de750cae764c5076" 970 + dependencies = [ 971 + "dasp_sample", 972 + ] 973 + 974 + [[package]] 975 + name = "dbus" 976 + version = "0.9.7" 977 + source = "registry+https://github.com/rust-lang/crates.io-index" 978 + checksum = "1bb21987b9fb1613058ba3843121dd18b163b254d8a6e797e144cbac14d96d1b" 979 + dependencies = [ 980 + "libc", 981 + "libdbus-sys", 982 + "winapi 0.3.9", 983 + ] 984 + 985 + [[package]] 986 + name = "deflate" 987 + version = "0.8.6" 988 + source = "registry+https://github.com/rust-lang/crates.io-index" 989 + checksum = "73770f8e1fe7d64df17ca66ad28994a0a623ea497fa69486e14984e715c5d174" 990 + dependencies = [ 991 + "adler32", 992 + "byteorder", 993 + ] 994 + 995 + [[package]] 996 + name = "digest" 997 + version = "0.10.6" 998 + source = "registry+https://github.com/rust-lang/crates.io-index" 999 + checksum = "8168378f4e5023e7218c89c891c0fd8ecdb5e5e4f18cb78f38cf245dd021e76f" 1000 + dependencies = [ 1001 + "block-buffer", 1002 + "crypto-common", 1003 + ] 1004 + 1005 + [[package]] 1006 + name = "directories-next" 1007 + version = "2.0.0" 1008 + source = "registry+https://github.com/rust-lang/crates.io-index" 1009 + checksum = "339ee130d97a610ea5a5872d2bbb130fdf68884ff09d3028b81bec8a1ac23bbc" 1010 + dependencies = [ 1011 + "cfg-if 1.0.0", 1012 + "dirs-sys-next", 1013 + ] 1014 + 1015 + [[package]] 1016 + name = "dirs-next" 1017 + version = "2.0.0" 1018 + source = "registry+https://github.com/rust-lang/crates.io-index" 1019 + checksum = "b98cf8ebf19c3d1b223e151f99a4f9f0690dca41414773390fc824184ac833e1" 1020 + dependencies = [ 1021 + "cfg-if 1.0.0", 1022 + "dirs-sys-next", 1023 + ] 1024 + 1025 + [[package]] 1026 + name = "dirs-sys-next" 1027 + version = "0.1.2" 1028 + source = "registry+https://github.com/rust-lang/crates.io-index" 1029 + checksum = "4ebda144c4fe02d1f7ea1a7d9641b6fc6b580adcfa024ae48797ecdeb6825b4d" 1030 + dependencies = [ 1031 + "libc", 1032 + "redox_users", 1033 + "winapi 0.3.9", 1034 + ] 1035 + 1036 + [[package]] 1037 + name = "dispatch" 1038 + version = "0.2.0" 1039 + source = "registry+https://github.com/rust-lang/crates.io-index" 1040 + checksum = "bd0c93bb4b0c6d9b77f4435b0ae98c24d17f1c45b2ff844c6151a07256ca923b" 1041 + 1042 + [[package]] 1043 + name = "docopt" 1044 + version = "1.1.1" 1045 + source = "registry+https://github.com/rust-lang/crates.io-index" 1046 + checksum = "7f3f119846c823f9eafcf953a8f6ffb6ed69bf6240883261a7f13b634579a51f" 1047 + dependencies = [ 1048 + "lazy_static", 1049 + "regex", 1050 + "serde 1.0.159", 1051 + "strsim 0.10.0", 1052 + ] 1053 + 1054 + [[package]] 1055 + name = "dtoa" 1056 + version = "0.4.8" 1057 + source = "registry+https://github.com/rust-lang/crates.io-index" 1058 + checksum = "56899898ce76aaf4a0f24d914c97ea6ed976d42fec6ad33fcbb0a1103e07b2b0" 1059 + 1060 + [[package]] 1061 + name = "ed25519" 1062 + version = "1.5.3" 1063 + source = "registry+https://github.com/rust-lang/crates.io-index" 1064 + checksum = "91cff35c70bba8a626e3185d8cd48cc11b5437e1a5bcd15b9b5fa3c64b6dfee7" 1065 + dependencies = [ 1066 + "signature", 1067 + ] 1068 + 1069 + [[package]] 1070 + name = "either" 1071 + version = "1.8.1" 1072 + source = "registry+https://github.com/rust-lang/crates.io-index" 1073 + checksum = "7fcaabb2fef8c910e7f4c7ce9f67a1283a1715879a7c230ca9d6d1ae31f16d91" 1074 + 1075 + [[package]] 1076 + name = "enigo" 1077 + version = "0.0.14" 1078 + dependencies = [ 1079 + "core-graphics", 1080 + "libc", 1081 + "log", 1082 + "objc", 1083 + "pkg-config", 1084 + "serde 1.0.159", 1085 + "serde_derive", 1086 + "unicode-segmentation", 1087 + "winapi 0.3.9", 1088 + ] 1089 + 1090 + [[package]] 1091 + name = "enum-map" 1092 + version = "2.5.0" 1093 + source = "registry+https://github.com/rust-lang/crates.io-index" 1094 + checksum = "988f0d17a0fa38291e5f41f71ea8d46a5d5497b9054d5a759fae2cbb819f2356" 1095 + dependencies = [ 1096 + "enum-map-derive", 1097 + ] 1098 + 1099 + [[package]] 1100 + name = "enum-map-derive" 1101 + version = "0.11.0" 1102 + source = "registry+https://github.com/rust-lang/crates.io-index" 1103 + checksum = "2a4da76b3b6116d758c7ba93f7ec6a35d2e2cf24feda76c6e38a375f4d5c59f2" 1104 + dependencies = [ 1105 + "proc-macro2", 1106 + "quote", 1107 + "syn 1.0.109", 1108 + ] 1109 + 1110 + [[package]] 1111 + name = "env_logger" 1112 + version = "0.9.3" 1113 + source = "registry+https://github.com/rust-lang/crates.io-index" 1114 + checksum = "a12e6657c4c97ebab115a42dcee77225f7f482cdd841cf7088c657a42e9e00e7" 1115 + dependencies = [ 1116 + "atty", 1117 + "humantime", 1118 + "log", 1119 + "regex", 1120 + "termcolor", 1121 + ] 1122 + 1123 + [[package]] 1124 + name = "env_logger" 1125 + version = "0.10.0" 1126 + source = "registry+https://github.com/rust-lang/crates.io-index" 1127 + checksum = "85cdab6a89accf66733ad5a1693a4dcced6aeff64602b634530dd73c1f3ee9f0" 1128 + dependencies = [ 1129 + "log", 1130 + "regex", 1131 + ] 1132 + 1133 + [[package]] 1134 + name = "err-derive" 1135 + version = "0.3.1" 1136 + source = "registry+https://github.com/rust-lang/crates.io-index" 1137 + checksum = "c34a887c8df3ed90498c1c437ce21f211c8e27672921a8ffa293cb8d6d4caa9e" 1138 + dependencies = [ 1139 + "proc-macro-error", 1140 + "proc-macro2", 1141 + "quote", 1142 + "rustversion", 1143 + "syn 1.0.109", 1144 + "synstructure", 1145 + ] 1146 + 1147 + [[package]] 1148 + name = "errno" 1149 + version = "0.3.1" 1150 + source = "registry+https://github.com/rust-lang/crates.io-index" 1151 + checksum = "4bcfec3a70f97c962c307b2d2c56e358cf1d00b558d74262b5f929ee8cc7e73a" 1152 + dependencies = [ 1153 + "errno-dragonfly", 1154 + "libc", 1155 + "windows-sys 0.48.0", 1156 + ] 1157 + 1158 + [[package]] 1159 + name = "errno-dragonfly" 1160 + version = "0.1.2" 1161 + source = "registry+https://github.com/rust-lang/crates.io-index" 1162 + checksum = "aa68f1b12764fab894d2755d2518754e71b4fd80ecfb822714a1206c2aab39bf" 1163 + dependencies = [ 1164 + "cc", 1165 + "libc", 1166 + ] 1167 + 1168 + [[package]] 1169 + name = "error-code" 1170 + version = "2.3.1" 1171 + source = "registry+https://github.com/rust-lang/crates.io-index" 1172 + checksum = "64f18991e7bf11e7ffee451b5318b5c1a73c52d0d0ada6e5a3017c8c1ced6a21" 1173 + dependencies = [ 1174 + "libc", 1175 + "str-buf", 1176 + ] 1177 + 1178 + [[package]] 1179 + name = "failure" 1180 + version = "0.1.8" 1181 + source = "registry+https://github.com/rust-lang/crates.io-index" 1182 + checksum = "d32e9bd16cc02eae7db7ef620b392808b89f6a5e16bb3497d159c6b92a0f4f86" 1183 + dependencies = [ 1184 + "backtrace", 1185 + ] 1186 + 1187 + [[package]] 1188 + name = "fastrand" 1189 + version = "1.9.0" 1190 + source = "registry+https://github.com/rust-lang/crates.io-index" 1191 + checksum = "e51093e27b0797c359783294ca4f0a911c270184cb10f85783b118614a1501be" 1192 + dependencies = [ 1193 + "instant", 1194 + ] 1195 + 1196 + [[package]] 1197 + name = "field-offset" 1198 + version = "0.3.5" 1199 + source = "registry+https://github.com/rust-lang/crates.io-index" 1200 + checksum = "a3cf3a800ff6e860c863ca6d4b16fd999db8b752819c1606884047b73e468535" 1201 + dependencies = [ 1202 + "memoffset 0.8.0", 1203 + "rustc_version", 1204 + ] 1205 + 1206 + [[package]] 1207 + name = "filetime" 1208 + version = "0.2.21" 1209 + source = "registry+https://github.com/rust-lang/crates.io-index" 1210 + checksum = "5cbc844cecaee9d4443931972e1289c8ff485cb4cc2767cb03ca139ed6885153" 1211 + dependencies = [ 1212 + "cfg-if 1.0.0", 1213 + "libc", 1214 + "redox_syscall 0.2.16", 1215 + "windows-sys 0.48.0", 1216 + ] 1217 + 1218 + [[package]] 1219 + name = "flate2" 1220 + version = "1.0.25" 1221 + source = "registry+https://github.com/rust-lang/crates.io-index" 1222 + checksum = "a8a2db397cb1c8772f31494cb8917e48cd1e64f0fa7efac59fbd741a0a8ce841" 1223 + dependencies = [ 1224 + "crc32fast", 1225 + "miniz_oxide 0.6.2", 1226 + ] 1227 + 1228 + [[package]] 1229 + name = "flexi_logger" 1230 + version = "0.22.6" 1231 + source = "registry+https://github.com/rust-lang/crates.io-index" 1232 + checksum = "0c76a80dd14a27fc3d8bc696502132cb52b3f227256fd8601166c3a35e45f409" 1233 + dependencies = [ 1234 + "ansi_term", 1235 + "atty", 1236 + "chrono", 1237 + "crossbeam-channel", 1238 + "crossbeam-queue", 1239 + "glob", 1240 + "lazy_static", 1241 + "log", 1242 + "regex", 1243 + "rustversion", 1244 + "thiserror", 1245 + "time", 1246 + ] 1247 + 1248 + [[package]] 1249 + name = "fnv" 1250 + version = "1.0.7" 1251 + source = "registry+https://github.com/rust-lang/crates.io-index" 1252 + checksum = "3f9eec918d3f24069decb9af1554cad7c880e2da24a9afd88aca000531ab82c1" 1253 + 1254 + [[package]] 1255 + name = "foreign-types" 1256 + version = "0.3.2" 1257 + source = "registry+https://github.com/rust-lang/crates.io-index" 1258 + checksum = "f6f339eb8adc052cd2ca78910fda869aefa38d22d5cb648e6485e4d3fc06f3b1" 1259 + dependencies = [ 1260 + "foreign-types-shared", 1261 + ] 1262 + 1263 + [[package]] 1264 + name = "foreign-types-shared" 1265 + version = "0.1.1" 1266 + source = "registry+https://github.com/rust-lang/crates.io-index" 1267 + checksum = "00b0228411908ca8685dba7fc2cdd70ec9990a6e753e89b6ac91a84c40fbaf4b" 1268 + 1269 + [[package]] 1270 + name = "fuchsia-cprng" 1271 + version = "0.1.1" 1272 + source = "registry+https://github.com/rust-lang/crates.io-index" 1273 + checksum = "a06f77d526c1a601b7c4cdd98f54b5eaabffc14d5f2f0296febdc7f357c6d3ba" 1274 + 1275 + [[package]] 1276 + name = "fuchsia-zircon" 1277 + version = "0.3.3" 1278 + source = "registry+https://github.com/rust-lang/crates.io-index" 1279 + checksum = "2e9763c69ebaae630ba35f74888db465e49e259ba1bc0eda7d06f4a067615d82" 1280 + dependencies = [ 1281 + "bitflags", 1282 + "fuchsia-zircon-sys", 1283 + ] 1284 + 1285 + [[package]] 1286 + name = "fuchsia-zircon-sys" 1287 + version = "0.3.3" 1288 + source = "registry+https://github.com/rust-lang/crates.io-index" 1289 + checksum = "3dcaa9ae7725d12cdb85b3ad99a434db70b468c09ded17e012d86b5c1010f7a7" 1290 + 1291 + [[package]] 1292 + name = "futures" 1293 + version = "0.3.28" 1294 + source = "registry+https://github.com/rust-lang/crates.io-index" 1295 + checksum = "23342abe12aba583913b2e62f22225ff9c950774065e4bfb61a19cd9770fec40" 1296 + dependencies = [ 1297 + "futures-channel", 1298 + "futures-core", 1299 + "futures-executor", 1300 + "futures-io", 1301 + "futures-sink", 1302 + "futures-task", 1303 + "futures-util", 1304 + ] 1305 + 1306 + [[package]] 1307 + name = "futures-channel" 1308 + version = "0.3.28" 1309 + source = "registry+https://github.com/rust-lang/crates.io-index" 1310 + checksum = "955518d47e09b25bbebc7a18df10b81f0c766eaf4c4f1cccef2fca5f2a4fb5f2" 1311 + dependencies = [ 1312 + "futures-core", 1313 + "futures-sink", 1314 + ] 1315 + 1316 + [[package]] 1317 + name = "futures-core" 1318 + version = "0.3.28" 1319 + source = "registry+https://github.com/rust-lang/crates.io-index" 1320 + checksum = "4bca583b7e26f571124fe5b7561d49cb2868d79116cfa0eefce955557c6fee8c" 1321 + 1322 + [[package]] 1323 + name = "futures-executor" 1324 + version = "0.3.28" 1325 + source = "registry+https://github.com/rust-lang/crates.io-index" 1326 + checksum = "ccecee823288125bd88b4d7f565c9e58e41858e47ab72e8ea2d64e93624386e0" 1327 + dependencies = [ 1328 + "futures-core", 1329 + "futures-task", 1330 + "futures-util", 1331 + ] 1332 + 1333 + [[package]] 1334 + name = "futures-io" 1335 + version = "0.3.28" 1336 + source = "registry+https://github.com/rust-lang/crates.io-index" 1337 + checksum = "4fff74096e71ed47f8e023204cfd0aa1289cd54ae5430a9523be060cdb849964" 1338 + 1339 + [[package]] 1340 + name = "futures-macro" 1341 + version = "0.3.28" 1342 + source = "registry+https://github.com/rust-lang/crates.io-index" 1343 + checksum = "89ca545a94061b6365f2c7355b4b32bd20df3ff95f02da9329b34ccc3bd6ee72" 1344 + dependencies = [ 1345 + "proc-macro2", 1346 + "quote", 1347 + "syn 2.0.13", 1348 + ] 1349 + 1350 + [[package]] 1351 + name = "futures-sink" 1352 + version = "0.3.28" 1353 + source = "registry+https://github.com/rust-lang/crates.io-index" 1354 + checksum = "f43be4fe21a13b9781a69afa4985b0f6ee0e1afab2c6f454a8cf30e2b2237b6e" 1355 + 1356 + [[package]] 1357 + name = "futures-task" 1358 + version = "0.3.28" 1359 + source = "registry+https://github.com/rust-lang/crates.io-index" 1360 + checksum = "76d3d132be6c0e6aa1534069c705a74a5997a356c0dc2f86a47765e5617c5b65" 1361 + 1362 + [[package]] 1363 + name = "futures-util" 1364 + version = "0.3.28" 1365 + source = "registry+https://github.com/rust-lang/crates.io-index" 1366 + checksum = "26b01e40b772d54cf6c6d721c1d1abd0647a0106a12ecaa1c186273392a69533" 1367 + dependencies = [ 1368 + "futures-channel", 1369 + "futures-core", 1370 + "futures-io", 1371 + "futures-macro", 1372 + "futures-sink", 1373 + "futures-task", 1374 + "memchr", 1375 + "pin-project-lite", 1376 + "pin-utils", 1377 + "slab", 1378 + ] 1379 + 1380 + [[package]] 1381 + name = "fxhash" 1382 + version = "0.2.1" 1383 + source = "registry+https://github.com/rust-lang/crates.io-index" 1384 + checksum = "c31b6d751ae2c7f11320402d34e41349dd1016f8d5d45e48c4312bc8625af50c" 1385 + dependencies = [ 1386 + "byteorder", 1387 + ] 1388 + 1389 + [[package]] 1390 + name = "gdk" 1391 + version = "0.13.2" 1392 + source = "registry+https://github.com/rust-lang/crates.io-index" 1393 + checksum = "db00839b2a68a7a10af3fa28dfb3febaba3a20c3a9ac2425a33b7df1f84a6b7d" 1394 + dependencies = [ 1395 + "bitflags", 1396 + "cairo-rs 0.9.1", 1397 + "cairo-sys-rs 0.10.0", 1398 + "gdk-pixbuf 0.9.0", 1399 + "gdk-sys 0.10.0", 1400 + "gio 0.9.1", 1401 + "gio-sys 0.10.1", 1402 + "glib 0.10.3", 1403 + "glib-sys 0.10.1", 1404 + "gobject-sys 0.10.0", 1405 + "libc", 1406 + "pango 0.9.1", 1407 + ] 1408 + 1409 + [[package]] 1410 + name = "gdk" 1411 + version = "0.15.4" 1412 + source = "registry+https://github.com/rust-lang/crates.io-index" 1413 + checksum = "a6e05c1f572ab0e1f15be94217f0dc29088c248b14f792a5ff0af0d84bcda9e8" 1414 + dependencies = [ 1415 + "bitflags", 1416 + "cairo-rs 0.15.12", 1417 + "gdk-pixbuf 0.15.11", 1418 + "gdk-sys 0.15.1", 1419 + "gio 0.15.12", 1420 + "glib 0.15.12", 1421 + "libc", 1422 + "pango 0.15.10", 1423 + ] 1424 + 1425 + [[package]] 1426 + name = "gdk-pixbuf" 1427 + version = "0.9.0" 1428 + source = "registry+https://github.com/rust-lang/crates.io-index" 1429 + checksum = "8f6dae3cb99dd49b758b88f0132f8d401108e63ae8edd45f432d42cdff99998a" 1430 + dependencies = [ 1431 + "gdk-pixbuf-sys 0.10.0", 1432 + "gio 0.9.1", 1433 + "gio-sys 0.10.1", 1434 + "glib 0.10.3", 1435 + "glib-sys 0.10.1", 1436 + "gobject-sys 0.10.0", 1437 + "libc", 1438 + ] 1439 + 1440 + [[package]] 1441 + name = "gdk-pixbuf" 1442 + version = "0.15.11" 1443 + source = "registry+https://github.com/rust-lang/crates.io-index" 1444 + checksum = "ad38dd9cc8b099cceecdf41375bb6d481b1b5a7cd5cd603e10a69a9383f8619a" 1445 + dependencies = [ 1446 + "bitflags", 1447 + "gdk-pixbuf-sys 0.15.10", 1448 + "gio 0.15.12", 1449 + "glib 0.15.12", 1450 + "libc", 1451 + ] 1452 + 1453 + [[package]] 1454 + name = "gdk-pixbuf-sys" 1455 + version = "0.10.0" 1456 + source = "registry+https://github.com/rust-lang/crates.io-index" 1457 + checksum = "3bfe468a7f43e97b8d193a762b6c5cf67a7d36cacbc0b9291dbcae24bfea1e8f" 1458 + dependencies = [ 1459 + "gio-sys 0.10.1", 1460 + "glib-sys 0.10.1", 1461 + "gobject-sys 0.10.0", 1462 + "libc", 1463 + "system-deps 1.3.2", 1464 + ] 1465 + 1466 + [[package]] 1467 + name = "gdk-pixbuf-sys" 1468 + version = "0.15.10" 1469 + source = "registry+https://github.com/rust-lang/crates.io-index" 1470 + checksum = "140b2f5378256527150350a8346dbdb08fadc13453a7a2d73aecd5fab3c402a7" 1471 + dependencies = [ 1472 + "gio-sys 0.15.10", 1473 + "glib-sys 0.15.10", 1474 + "gobject-sys 0.15.10", 1475 + "libc", 1476 + "system-deps 6.0.4", 1477 + ] 1478 + 1479 + [[package]] 1480 + name = "gdk-sys" 1481 + version = "0.10.0" 1482 + source = "registry+https://github.com/rust-lang/crates.io-index" 1483 + checksum = "0a9653cfc500fd268015b1ac055ddbc3df7a5c9ea3f4ccef147b3957bd140d69" 1484 + dependencies = [ 1485 + "cairo-sys-rs 0.10.0", 1486 + "gdk-pixbuf-sys 0.10.0", 1487 + "gio-sys 0.10.1", 1488 + "glib-sys 0.10.1", 1489 + "gobject-sys 0.10.0", 1490 + "libc", 1491 + "pango-sys 0.10.0", 1492 + "pkg-config", 1493 + "system-deps 1.3.2", 1494 + ] 1495 + 1496 + [[package]] 1497 + name = "gdk-sys" 1498 + version = "0.15.1" 1499 + source = "registry+https://github.com/rust-lang/crates.io-index" 1500 + checksum = "32e7a08c1e8f06f4177fb7e51a777b8c1689f743a7bc11ea91d44d2226073a88" 1501 + dependencies = [ 1502 + "cairo-sys-rs 0.15.1", 1503 + "gdk-pixbuf-sys 0.15.10", 1504 + "gio-sys 0.15.10", 1505 + "glib-sys 0.15.10", 1506 + "gobject-sys 0.15.10", 1507 + "libc", 1508 + "pango-sys 0.15.10", 1509 + "pkg-config", 1510 + "system-deps 6.0.4", 1511 + ] 1512 + 1513 + [[package]] 1514 + name = "generic-array" 1515 + version = "0.14.7" 1516 + source = "registry+https://github.com/rust-lang/crates.io-index" 1517 + checksum = "85649ca51fd72272d7821adaf274ad91c288277713d9c18820d8499a7ff69e9a" 1518 + dependencies = [ 1519 + "typenum", 1520 + "version_check", 1521 + ] 1522 + 1523 + [[package]] 1524 + name = "gethostname" 1525 + version = "0.2.3" 1526 + source = "registry+https://github.com/rust-lang/crates.io-index" 1527 + checksum = "c1ebd34e35c46e00bb73e81363248d627782724609fe1b6396f553f68fe3862e" 1528 + dependencies = [ 1529 + "libc", 1530 + "winapi 0.3.9", 1531 + ] 1532 + 1533 + [[package]] 1534 + name = "getrandom" 1535 + version = "0.2.9" 1536 + source = "registry+https://github.com/rust-lang/crates.io-index" 1537 + checksum = "c85e1d9ab2eadba7e5040d4e09cbd6d072b76a557ad64e797c2cb9d4da21d7e4" 1538 + dependencies = [ 1539 + "cfg-if 1.0.0", 1540 + "libc", 1541 + "wasi", 1542 + ] 1543 + 1544 + [[package]] 1545 + name = "gimli" 1546 + version = "0.27.2" 1547 + source = "registry+https://github.com/rust-lang/crates.io-index" 1548 + checksum = "ad0a93d233ebf96623465aad4046a8d3aa4da22d4f4beba5388838c8a434bbb4" 1549 + 1550 + [[package]] 1551 + name = "gio" 1552 + version = "0.9.1" 1553 + source = "registry+https://github.com/rust-lang/crates.io-index" 1554 + checksum = "1fb60242bfff700772dae5d9e3a1f7aa2e4ebccf18b89662a16acb2822568561" 1555 + dependencies = [ 1556 + "bitflags", 1557 + "futures", 1558 + "futures-channel", 1559 + "futures-core", 1560 + "futures-io", 1561 + "futures-util", 1562 + "gio-sys 0.10.1", 1563 + "glib 0.10.3", 1564 + "glib-sys 0.10.1", 1565 + "gobject-sys 0.10.0", 1566 + "libc", 1567 + "once_cell", 1568 + "thiserror", 1569 + ] 1570 + 1571 + [[package]] 1572 + name = "gio" 1573 + version = "0.15.12" 1574 + source = "registry+https://github.com/rust-lang/crates.io-index" 1575 + checksum = "68fdbc90312d462781a395f7a16d96a2b379bb6ef8cd6310a2df272771c4283b" 1576 + dependencies = [ 1577 + "bitflags", 1578 + "futures-channel", 1579 + "futures-core", 1580 + "futures-io", 1581 + "gio-sys 0.15.10", 1582 + "glib 0.15.12", 1583 + "libc", 1584 + "once_cell", 1585 + "thiserror", 1586 + ] 1587 + 1588 + [[package]] 1589 + name = "gio-sys" 1590 + version = "0.10.1" 1591 + source = "registry+https://github.com/rust-lang/crates.io-index" 1592 + checksum = "5e24fb752f8f5d2cf6bbc2c606fd2bc989c81c5e2fe321ab974d54f8b6344eac" 1593 + dependencies = [ 1594 + "glib-sys 0.10.1", 1595 + "gobject-sys 0.10.0", 1596 + "libc", 1597 + "system-deps 1.3.2", 1598 + "winapi 0.3.9", 1599 + ] 1600 + 1601 + [[package]] 1602 + name = "gio-sys" 1603 + version = "0.15.10" 1604 + source = "registry+https://github.com/rust-lang/crates.io-index" 1605 + checksum = "32157a475271e2c4a023382e9cab31c4584ee30a97da41d3c4e9fdd605abcf8d" 1606 + dependencies = [ 1607 + "glib-sys 0.15.10", 1608 + "gobject-sys 0.15.10", 1609 + "libc", 1610 + "system-deps 6.0.4", 1611 + "winapi 0.3.9", 1612 + ] 1613 + 1614 + [[package]] 1615 + name = "glib" 1616 + version = "0.10.3" 1617 + source = "registry+https://github.com/rust-lang/crates.io-index" 1618 + checksum = "0c685013b7515e668f1b57a165b009d4d28cb139a8a989bbd699c10dad29d0c5" 1619 + dependencies = [ 1620 + "bitflags", 1621 + "futures-channel", 1622 + "futures-core", 1623 + "futures-executor", 1624 + "futures-task", 1625 + "futures-util", 1626 + "glib-macros 0.10.1", 1627 + "glib-sys 0.10.1", 1628 + "gobject-sys 0.10.0", 1629 + "libc", 1630 + "once_cell", 1631 + ] 1632 + 1633 + [[package]] 1634 + name = "glib" 1635 + version = "0.15.12" 1636 + source = "registry+https://github.com/rust-lang/crates.io-index" 1637 + checksum = "edb0306fbad0ab5428b0ca674a23893db909a98582969c9b537be4ced78c505d" 1638 + dependencies = [ 1639 + "bitflags", 1640 + "futures-channel", 1641 + "futures-core", 1642 + "futures-executor", 1643 + "futures-task", 1644 + "glib-macros 0.15.13", 1645 + "glib-sys 0.15.10", 1646 + "gobject-sys 0.15.10", 1647 + "libc", 1648 + "once_cell", 1649 + "smallvec", 1650 + "thiserror", 1651 + ] 1652 + 1653 + [[package]] 1654 + name = "glib-macros" 1655 + version = "0.10.1" 1656 + source = "registry+https://github.com/rust-lang/crates.io-index" 1657 + checksum = "41486a26d1366a8032b160b59065a59fb528530a46a49f627e7048fb8c064039" 1658 + dependencies = [ 1659 + "anyhow", 1660 + "heck 0.3.3", 1661 + "itertools", 1662 + "proc-macro-crate 0.1.5", 1663 + "proc-macro-error", 1664 + "proc-macro2", 1665 + "quote", 1666 + "syn 1.0.109", 1667 + ] 1668 + 1669 + [[package]] 1670 + name = "glib-macros" 1671 + version = "0.15.13" 1672 + source = "registry+https://github.com/rust-lang/crates.io-index" 1673 + checksum = "10c6ae9f6fa26f4fb2ac16b528d138d971ead56141de489f8111e259b9df3c4a" 1674 + dependencies = [ 1675 + "anyhow", 1676 + "heck 0.4.1", 1677 + "proc-macro-crate 1.3.1", 1678 + "proc-macro-error", 1679 + "proc-macro2", 1680 + "quote", 1681 + "syn 1.0.109", 1682 + ] 1683 + 1684 + [[package]] 1685 + name = "glib-sys" 1686 + version = "0.10.1" 1687 + source = "registry+https://github.com/rust-lang/crates.io-index" 1688 + checksum = "c7e9b997a66e9a23d073f2b1abb4dbfc3925e0b8952f67efd8d9b6e168e4cdc1" 1689 + dependencies = [ 1690 + "libc", 1691 + "system-deps 1.3.2", 1692 + ] 1693 + 1694 + [[package]] 1695 + name = "glib-sys" 1696 + version = "0.15.10" 1697 + source = "registry+https://github.com/rust-lang/crates.io-index" 1698 + checksum = "ef4b192f8e65e9cf76cbf4ea71fa8e3be4a0e18ffe3d68b8da6836974cc5bad4" 1699 + dependencies = [ 1700 + "libc", 1701 + "system-deps 6.0.4", 1702 + ] 1703 + 1704 + [[package]] 1705 + name = "glob" 1706 + version = "0.3.1" 1707 + source = "registry+https://github.com/rust-lang/crates.io-index" 1708 + checksum = "d2fabcfbdc87f4758337ca535fb41a6d701b65693ce38287d856d1674551ec9b" 1709 + 1710 + [[package]] 1711 + name = "gobject-sys" 1712 + version = "0.10.0" 1713 + source = "registry+https://github.com/rust-lang/crates.io-index" 1714 + checksum = "952133b60c318a62bf82ee75b93acc7e84028a093e06b9e27981c2b6fe68218c" 1715 + dependencies = [ 1716 + "glib-sys 0.10.1", 1717 + "libc", 1718 + "system-deps 1.3.2", 1719 + ] 1720 + 1721 + [[package]] 1722 + name = "gobject-sys" 1723 + version = "0.15.10" 1724 + source = "registry+https://github.com/rust-lang/crates.io-index" 1725 + checksum = "0d57ce44246becd17153bd035ab4d32cfee096a657fc01f2231c9278378d1e0a" 1726 + dependencies = [ 1727 + "glib-sys 0.15.10", 1728 + "libc", 1729 + "system-deps 6.0.4", 1730 + ] 1731 + 1732 + [[package]] 1733 + name = "gstreamer" 1734 + version = "0.16.7" 1735 + source = "registry+https://github.com/rust-lang/crates.io-index" 1736 + checksum = "9ff5d0f7ff308ae37e6eb47b6ded17785bdea06e438a708cd09e0288c1862f33" 1737 + dependencies = [ 1738 + "bitflags", 1739 + "cfg-if 1.0.0", 1740 + "futures-channel", 1741 + "futures-core", 1742 + "futures-util", 1743 + "glib 0.10.3", 1744 + "glib-sys 0.10.1", 1745 + "gobject-sys 0.10.0", 1746 + "gstreamer-sys", 1747 + "libc", 1748 + "muldiv", 1749 + "num-rational", 1750 + "once_cell", 1751 + "paste", 1752 + "pretty-hex", 1753 + "thiserror", 1754 + ] 1755 + 1756 + [[package]] 1757 + name = "gstreamer-app" 1758 + version = "0.16.5" 1759 + source = "registry+https://github.com/rust-lang/crates.io-index" 1760 + checksum = "cc80888271338c3ede875d8cafc452eb207476ff5539dcbe0018a8f5b827af0e" 1761 + dependencies = [ 1762 + "bitflags", 1763 + "futures-core", 1764 + "futures-sink", 1765 + "glib 0.10.3", 1766 + "glib-sys 0.10.1", 1767 + "gobject-sys 0.10.0", 1768 + "gstreamer", 1769 + "gstreamer-app-sys", 1770 + "gstreamer-base", 1771 + "gstreamer-sys", 1772 + "libc", 1773 + "once_cell", 1774 + ] 1775 + 1776 + [[package]] 1777 + name = "gstreamer-app-sys" 1778 + version = "0.9.1" 1779 + source = "registry+https://github.com/rust-lang/crates.io-index" 1780 + checksum = "813f64275c9e7b33b828b9efcf9dfa64b95996766d4de996e84363ac65b87e3d" 1781 + dependencies = [ 1782 + "glib-sys 0.10.1", 1783 + "gstreamer-base-sys", 1784 + "gstreamer-sys", 1785 + "libc", 1786 + "system-deps 1.3.2", 1787 + ] 1788 + 1789 + [[package]] 1790 + name = "gstreamer-base" 1791 + version = "0.16.5" 1792 + source = "registry+https://github.com/rust-lang/crates.io-index" 1793 + checksum = "bafd01c56f59cb10f4b5a10f97bb4bdf8c2b2784ae5b04da7e2d400cf6e6afcf" 1794 + dependencies = [ 1795 + "bitflags", 1796 + "glib 0.10.3", 1797 + "glib-sys 0.10.1", 1798 + "gobject-sys 0.10.0", 1799 + "gstreamer", 1800 + "gstreamer-base-sys", 1801 + "gstreamer-sys", 1802 + "libc", 1803 + ] 1804 + 1805 + [[package]] 1806 + name = "gstreamer-base-sys" 1807 + version = "0.9.1" 1808 + source = "registry+https://github.com/rust-lang/crates.io-index" 1809 + checksum = "a4b7b6dc2d6e160a1ae28612f602bd500b3fa474ce90bf6bb2f08072682beef5" 1810 + dependencies = [ 1811 + "glib-sys 0.10.1", 1812 + "gobject-sys 0.10.0", 1813 + "gstreamer-sys", 1814 + "libc", 1815 + "system-deps 1.3.2", 1816 + ] 1817 + 1818 + [[package]] 1819 + name = "gstreamer-sys" 1820 + version = "0.9.1" 1821 + source = "registry+https://github.com/rust-lang/crates.io-index" 1822 + checksum = "fc1f154082d01af5718c5f8a8eb4f565a4ea5586ad8833a8fc2c2aa6844b601d" 1823 + dependencies = [ 1824 + "glib-sys 0.10.1", 1825 + "gobject-sys 0.10.0", 1826 + "libc", 1827 + "system-deps 1.3.2", 1828 + ] 1829 + 1830 + [[package]] 1831 + name = "gstreamer-video" 1832 + version = "0.16.7" 1833 + source = "registry+https://github.com/rust-lang/crates.io-index" 1834 + checksum = "f7bbb1485d87469849ec45c08e03c2f280d3ea20ff3c439d03185be54e3ce98e" 1835 + dependencies = [ 1836 + "bitflags", 1837 + "futures-channel", 1838 + "futures-util", 1839 + "glib 0.10.3", 1840 + "glib-sys 0.10.1", 1841 + "gobject-sys 0.10.0", 1842 + "gstreamer", 1843 + "gstreamer-base", 1844 + "gstreamer-base-sys", 1845 + "gstreamer-sys", 1846 + "gstreamer-video-sys", 1847 + "libc", 1848 + "once_cell", 1849 + ] 1850 + 1851 + [[package]] 1852 + name = "gstreamer-video-sys" 1853 + version = "0.9.1" 1854 + source = "registry+https://github.com/rust-lang/crates.io-index" 1855 + checksum = "92347e46438007d6a2386302125f62cb9df6769cdacb931af5c0f12c1ee21de4" 1856 + dependencies = [ 1857 + "glib-sys 0.10.1", 1858 + "gobject-sys 0.10.0", 1859 + "gstreamer-base-sys", 1860 + "gstreamer-sys", 1861 + "libc", 1862 + "system-deps 1.3.2", 1863 + ] 1864 + 1865 + [[package]] 1866 + name = "gtk" 1867 + version = "0.9.2" 1868 + source = "registry+https://github.com/rust-lang/crates.io-index" 1869 + checksum = "2f022f2054072b3af07666341984562c8e626a79daa8be27b955d12d06a5ad6a" 1870 + dependencies = [ 1871 + "atk 0.9.0", 1872 + "bitflags", 1873 + "cairo-rs 0.9.1", 1874 + "cairo-sys-rs 0.10.0", 1875 + "cc", 1876 + "gdk 0.13.2", 1877 + "gdk-pixbuf 0.9.0", 1878 + "gdk-pixbuf-sys 0.10.0", 1879 + "gdk-sys 0.10.0", 1880 + "gio 0.9.1", 1881 + "gio-sys 0.10.1", 1882 + "glib 0.10.3", 1883 + "glib-sys 0.10.1", 1884 + "gobject-sys 0.10.0", 1885 + "gtk-sys 0.10.0", 1886 + "libc", 1887 + "once_cell", 1888 + "pango 0.9.1", 1889 + "pango-sys 0.10.0", 1890 + "pkg-config", 1891 + ] 1892 + 1893 + [[package]] 1894 + name = "gtk" 1895 + version = "0.15.5" 1896 + source = "registry+https://github.com/rust-lang/crates.io-index" 1897 + checksum = "92e3004a2d5d6d8b5057d2b57b3712c9529b62e82c77f25c1fecde1fd5c23bd0" 1898 + dependencies = [ 1899 + "atk 0.15.1", 1900 + "bitflags", 1901 + "cairo-rs 0.15.12", 1902 + "field-offset", 1903 + "futures-channel", 1904 + "gdk 0.15.4", 1905 + "gdk-pixbuf 0.15.11", 1906 + "gio 0.15.12", 1907 + "glib 0.15.12", 1908 + "gtk-sys 0.15.3", 1909 + "gtk3-macros", 1910 + "libc", 1911 + "once_cell", 1912 + "pango 0.15.10", 1913 + "pkg-config", 1914 + ] 1915 + 1916 + [[package]] 1917 + name = "gtk-sys" 1918 + version = "0.10.0" 1919 + source = "registry+https://github.com/rust-lang/crates.io-index" 1920 + checksum = "89acda6f084863307d948ba64a4b1ef674e8527dddab147ee4cdcc194c880457" 1921 + dependencies = [ 1922 + "atk-sys 0.10.0", 1923 + "cairo-sys-rs 0.10.0", 1924 + "gdk-pixbuf-sys 0.10.0", 1925 + "gdk-sys 0.10.0", 1926 + "gio-sys 0.10.1", 1927 + "glib-sys 0.10.1", 1928 + "gobject-sys 0.10.0", 1929 + "libc", 1930 + "pango-sys 0.10.0", 1931 + "system-deps 1.3.2", 1932 + ] 1933 + 1934 + [[package]] 1935 + name = "gtk-sys" 1936 + version = "0.15.3" 1937 + source = "registry+https://github.com/rust-lang/crates.io-index" 1938 + checksum = "d5bc2f0587cba247f60246a0ca11fe25fb733eabc3de12d1965fc07efab87c84" 1939 + dependencies = [ 1940 + "atk-sys 0.15.1", 1941 + "cairo-sys-rs 0.15.1", 1942 + "gdk-pixbuf-sys 0.15.10", 1943 + "gdk-sys 0.15.1", 1944 + "gio-sys 0.15.10", 1945 + "glib-sys 0.15.10", 1946 + "gobject-sys 0.15.10", 1947 + "libc", 1948 + "pango-sys 0.15.10", 1949 + "system-deps 6.0.4", 1950 + ] 1951 + 1952 + [[package]] 1953 + name = "gtk3-macros" 1954 + version = "0.15.6" 1955 + source = "registry+https://github.com/rust-lang/crates.io-index" 1956 + checksum = "684c0456c086e8e7e9af73ec5b84e35938df394712054550e81558d21c44ab0d" 1957 + dependencies = [ 1958 + "anyhow", 1959 + "proc-macro-crate 1.3.1", 1960 + "proc-macro-error", 1961 + "proc-macro2", 1962 + "quote", 1963 + "syn 1.0.109", 1964 + ] 1965 + 1966 + [[package]] 1967 + name = "hashbrown" 1968 + version = "0.12.3" 1969 + source = "registry+https://github.com/rust-lang/crates.io-index" 1970 + checksum = "8a9ee70c43aaf417c914396645a0fa852624801b24ebb7ae78fe8272889ac888" 1971 + 1972 + [[package]] 1973 + name = "hbb_common" 1974 + version = "0.1.0" 1975 + dependencies = [ 1976 + "anyhow", 1977 + "bytes", 1978 + "confy", 1979 + "directories-next", 1980 + "dirs-next", 1981 + "env_logger 0.9.3", 1982 + "filetime", 1983 + "futures", 1984 + "futures-util", 1985 + "lazy_static", 1986 + "log", 1987 + "mac_address", 1988 + "protobuf", 1989 + "protobuf-codegen-pure", 1990 + "quinn", 1991 + "rand 0.8.5", 1992 + "regex", 1993 + "serde 1.0.159", 1994 + "serde_derive", 1995 + "serde_json 1.0.95", 1996 + "socket2 0.3.19", 1997 + "sodiumoxide", 1998 + "tokio", 1999 + "tokio-socks", 2000 + "tokio-util 0.6.10", 2001 + "toml 0.5.11", 2002 + "winapi 0.3.9", 2003 + "zstd", 2004 + ] 2005 + 2006 + [[package]] 2007 + name = "heck" 2008 + version = "0.3.3" 2009 + source = "registry+https://github.com/rust-lang/crates.io-index" 2010 + checksum = "6d621efb26863f0e9924c6ac577e8275e5e6b77455db64ffa6c65c904e9e132c" 2011 + dependencies = [ 2012 + "unicode-segmentation", 2013 + ] 2014 + 2015 + [[package]] 2016 + name = "heck" 2017 + version = "0.4.1" 2018 + source = "registry+https://github.com/rust-lang/crates.io-index" 2019 + checksum = "95505c38b4572b2d910cecb0281560f54b440a19336cbbcb27bf6ce6adc6f5a8" 2020 + 2021 + [[package]] 2022 + name = "hermit-abi" 2023 + version = "0.1.19" 2024 + source = "registry+https://github.com/rust-lang/crates.io-index" 2025 + checksum = "62b467343b94ba476dcb2500d242dadbb39557df889310ac77c5d99100aaac33" 2026 + dependencies = [ 2027 + "libc", 2028 + ] 2029 + 2030 + [[package]] 2031 + name = "hermit-abi" 2032 + version = "0.2.6" 2033 + source = "registry+https://github.com/rust-lang/crates.io-index" 2034 + checksum = "ee512640fe35acbfb4bb779db6f0d80704c2cacfa2e39b601ef3e3f47d1ae4c7" 2035 + dependencies = [ 2036 + "libc", 2037 + ] 2038 + 2039 + [[package]] 2040 + name = "hermit-abi" 2041 + version = "0.3.1" 2042 + source = "registry+https://github.com/rust-lang/crates.io-index" 2043 + checksum = "fed44880c466736ef9a5c5b5facefb5ed0785676d0c02d612db14e54f0d84286" 2044 + 2045 + [[package]] 2046 + name = "hound" 2047 + version = "3.5.0" 2048 + source = "registry+https://github.com/rust-lang/crates.io-index" 2049 + checksum = "4d13cdbd5dbb29f9c88095bbdc2590c9cba0d0a1269b983fef6b2cdd7e9f4db1" 2050 + 2051 + [[package]] 2052 + name = "humantime" 2053 + version = "2.1.0" 2054 + source = "registry+https://github.com/rust-lang/crates.io-index" 2055 + checksum = "9a3a5bfb195931eeb336b2a7b4d761daec841b97f947d34394601737a7bba5e4" 2056 + 2057 + [[package]] 2058 + name = "iana-time-zone" 2059 + version = "0.1.56" 2060 + source = "registry+https://github.com/rust-lang/crates.io-index" 2061 + checksum = "0722cd7114b7de04316e7ea5456a0bbb20e4adb46fd27a3697adb812cff0f37c" 2062 + dependencies = [ 2063 + "android_system_properties", 2064 + "core-foundation-sys", 2065 + "iana-time-zone-haiku", 2066 + "js-sys", 2067 + "wasm-bindgen", 2068 + "windows", 2069 + ] 2070 + 2071 + [[package]] 2072 + name = "iana-time-zone-haiku" 2073 + version = "0.1.1" 2074 + source = "registry+https://github.com/rust-lang/crates.io-index" 2075 + checksum = "0703ae284fc167426161c2e3f1da3ea71d94b21bedbcc9494e92b28e334e3dca" 2076 + dependencies = [ 2077 + "cxx", 2078 + "cxx-build", 2079 + ] 2080 + 2081 + [[package]] 2082 + name = "ident_case" 2083 + version = "1.0.1" 2084 + source = "registry+https://github.com/rust-lang/crates.io-index" 2085 + checksum = "b9e0384b61958566e926dc50660321d12159025e767c18e043daf26b70104c39" 2086 + 2087 + [[package]] 2088 + name = "image" 2089 + version = "0.23.14" 2090 + source = "registry+https://github.com/rust-lang/crates.io-index" 2091 + checksum = "24ffcb7e7244a9bf19d35bf2883b9c080c4ced3c07a9895572178cdb8f13f6a1" 2092 + dependencies = [ 2093 + "bytemuck", 2094 + "byteorder", 2095 + "color_quant", 2096 + "num-iter", 2097 + "num-rational", 2098 + "num-traits 0.2.15", 2099 + "png", 2100 + "tiff", 2101 + ] 2102 + 2103 + [[package]] 2104 + name = "include_dir" 2105 + version = "0.7.3" 2106 + source = "registry+https://github.com/rust-lang/crates.io-index" 2107 + checksum = "18762faeff7122e89e0857b02f7ce6fcc0d101d5e9ad2ad7846cc01d61b7f19e" 2108 + dependencies = [ 2109 + "include_dir_macros", 2110 + ] 2111 + 2112 + [[package]] 2113 + name = "include_dir_macros" 2114 + version = "0.7.3" 2115 + source = "registry+https://github.com/rust-lang/crates.io-index" 2116 + checksum = "b139284b5cf57ecfa712bcc66950bb635b31aff41c188e8a4cfc758eca374a3f" 2117 + dependencies = [ 2118 + "proc-macro2", 2119 + "quote", 2120 + ] 2121 + 2122 + [[package]] 2123 + name = "indexmap" 2124 + version = "1.9.3" 2125 + source = "registry+https://github.com/rust-lang/crates.io-index" 2126 + checksum = "bd070e393353796e801d209ad339e89596eb4c8d430d18ede6a1cced8fafbd99" 2127 + dependencies = [ 2128 + "autocfg 1.1.0", 2129 + "hashbrown", 2130 + ] 2131 + 2132 + [[package]] 2133 + name = "instant" 2134 + version = "0.1.12" 2135 + source = "registry+https://github.com/rust-lang/crates.io-index" 2136 + checksum = "7a5bbe824c507c5da5956355e86a746d82e0e1464f65d862cc5e71da70e94b2c" 2137 + dependencies = [ 2138 + "cfg-if 1.0.0", 2139 + ] 2140 + 2141 + [[package]] 2142 + name = "io-lifetimes" 2143 + version = "1.0.10" 2144 + source = "registry+https://github.com/rust-lang/crates.io-index" 2145 + checksum = "9c66c74d2ae7e79a5a8f7ac924adbe38ee42a859c6539ad869eb51f0b52dc220" 2146 + dependencies = [ 2147 + "hermit-abi 0.3.1", 2148 + "libc", 2149 + "windows-sys 0.48.0", 2150 + ] 2151 + 2152 + [[package]] 2153 + name = "iovec" 2154 + version = "0.1.4" 2155 + source = "registry+https://github.com/rust-lang/crates.io-index" 2156 + checksum = "b2b3ea6ff95e175473f8ffe6a7eb7c00d054240321b84c57051175fe3c1e075e" 2157 + dependencies = [ 2158 + "libc", 2159 + ] 2160 + 2161 + [[package]] 2162 + name = "itertools" 2163 + version = "0.9.0" 2164 + source = "registry+https://github.com/rust-lang/crates.io-index" 2165 + checksum = "284f18f85651fe11e8a991b2adb42cb078325c996ed026d994719efcfca1d54b" 2166 + dependencies = [ 2167 + "either", 2168 + ] 2169 + 2170 + [[package]] 2171 + name = "itoa" 2172 + version = "0.3.4" 2173 + source = "registry+https://github.com/rust-lang/crates.io-index" 2174 + checksum = "8324a32baf01e2ae060e9de58ed0bc2320c9a2833491ee36cd3b4c414de4db8c" 2175 + 2176 + [[package]] 2177 + name = "itoa" 2178 + version = "1.0.6" 2179 + source = "registry+https://github.com/rust-lang/crates.io-index" 2180 + checksum = "453ad9f582a441959e5f0d088b02ce04cfe8d51a8eaf077f12ac6d3e94164ca6" 2181 + 2182 + [[package]] 2183 + name = "jni" 2184 + version = "0.19.0" 2185 + source = "registry+https://github.com/rust-lang/crates.io-index" 2186 + checksum = "c6df18c2e3db7e453d3c6ac5b3e9d5182664d28788126d39b91f2d1e22b017ec" 2187 + dependencies = [ 2188 + "cesu8", 2189 + "combine", 2190 + "jni-sys", 2191 + "log", 2192 + "thiserror", 2193 + "walkdir", 2194 + ] 2195 + 2196 + [[package]] 2197 + name = "jni-sys" 2198 + version = "0.3.0" 2199 + source = "registry+https://github.com/rust-lang/crates.io-index" 2200 + checksum = "8eaf4bc02d17cbdd7ff4c7438cafcdf7fb9a4613313ad11b4f8fefe7d3fa0130" 2201 + 2202 + [[package]] 2203 + name = "jobserver" 2204 + version = "0.1.26" 2205 + source = "registry+https://github.com/rust-lang/crates.io-index" 2206 + checksum = "936cfd212a0155903bcbc060e316fb6cc7cbf2e1907329391ebadc1fe0ce77c2" 2207 + dependencies = [ 2208 + "libc", 2209 + ] 2210 + 2211 + [[package]] 2212 + name = "jpeg-decoder" 2213 + version = "0.1.22" 2214 + source = "registry+https://github.com/rust-lang/crates.io-index" 2215 + checksum = "229d53d58899083193af11e15917b5640cd40b29ff475a1fe4ef725deb02d0f2" 2216 + 2217 + [[package]] 2218 + name = "js-sys" 2219 + version = "0.3.61" 2220 + source = "registry+https://github.com/rust-lang/crates.io-index" 2221 + checksum = "445dde2150c55e483f3d8416706b97ec8e8237c307e5b7b4b8dd15e6af2a0730" 2222 + dependencies = [ 2223 + "wasm-bindgen", 2224 + ] 2225 + 2226 + [[package]] 2227 + name = "kernel32-sys" 2228 + version = "0.2.2" 2229 + source = "registry+https://github.com/rust-lang/crates.io-index" 2230 + checksum = "7507624b29483431c0ba2d82aece8ca6cdba9382bff4ddd0f7490560c056098d" 2231 + dependencies = [ 2232 + "winapi 0.2.8", 2233 + "winapi-build", 2234 + ] 2235 + 2236 + [[package]] 2237 + name = "lazy_static" 2238 + version = "1.4.0" 2239 + source = "registry+https://github.com/rust-lang/crates.io-index" 2240 + checksum = "e2abad23fbc42b3700f2f279844dc832adb2b2eb069b2df918f455c4e18cc646" 2241 + 2242 + [[package]] 2243 + name = "lazycell" 2244 + version = "1.3.0" 2245 + source = "registry+https://github.com/rust-lang/crates.io-index" 2246 + checksum = "830d08ce1d1d941e6b30645f1a0eb5643013d835ce3779a5fc208261dbe10f55" 2247 + 2248 + [[package]] 2249 + name = "libappindicator" 2250 + version = "0.6.1" 2251 + source = "git+https://github.com/liyue201/libappindicator-rs#3763cfd629dd90050af1feafa643cbfca0bf487e" 2252 + dependencies = [ 2253 + "glib 0.10.3", 2254 + "gtk 0.9.2", 2255 + "gtk-sys 0.10.0", 2256 + "libappindicator-sys 0.6.1", 2257 + "log", 2258 + ] 2259 + 2260 + [[package]] 2261 + name = "libappindicator" 2262 + version = "0.7.1" 2263 + source = "registry+https://github.com/rust-lang/crates.io-index" 2264 + checksum = "db2d3cb96d092b4824cb306c9e544c856a4cb6210c1081945187f7f1924b47e8" 2265 + dependencies = [ 2266 + "glib 0.15.12", 2267 + "gtk 0.15.5", 2268 + "gtk-sys 0.15.3", 2269 + "libappindicator-sys 0.7.3", 2270 + "log", 2271 + ] 2272 + 2273 + [[package]] 2274 + name = "libappindicator-sys" 2275 + version = "0.6.1" 2276 + source = "git+https://github.com/liyue201/libappindicator-rs#3763cfd629dd90050af1feafa643cbfca0bf487e" 2277 + dependencies = [ 2278 + "gtk-sys 0.10.0", 2279 + "pkg-config", 2280 + ] 2281 + 2282 + [[package]] 2283 + name = "libappindicator-sys" 2284 + version = "0.7.3" 2285 + source = "registry+https://github.com/rust-lang/crates.io-index" 2286 + checksum = "f1b3b6681973cea8cc3bce7391e6d7d5502720b80a581c9a95c9cbaf592826aa" 2287 + dependencies = [ 2288 + "gtk-sys 0.15.3", 2289 + "libloading", 2290 + "once_cell", 2291 + ] 2292 + 2293 + [[package]] 2294 + name = "libc" 2295 + version = "0.2.141" 2296 + source = "registry+https://github.com/rust-lang/crates.io-index" 2297 + checksum = "3304a64d199bb964be99741b7a14d26972741915b3649639149b2479bb46f4b5" 2298 + 2299 + [[package]] 2300 + name = "libdbus-sys" 2301 + version = "0.2.4" 2302 + source = "registry+https://github.com/rust-lang/crates.io-index" 2303 + checksum = "9f8d7ae751e1cb825c840ae5e682f59b098cdfd213c350ac268b61449a5f58a0" 2304 + dependencies = [ 2305 + "pkg-config", 2306 + ] 2307 + 2308 + [[package]] 2309 + name = "libloading" 2310 + version = "0.7.4" 2311 + source = "registry+https://github.com/rust-lang/crates.io-index" 2312 + checksum = "b67380fd3b2fbe7527a606e18729d21c6f3951633d0500574c4dc22d2d638b9f" 2313 + dependencies = [ 2314 + "cfg-if 1.0.0", 2315 + "winapi 0.3.9", 2316 + ] 2317 + 2318 + [[package]] 2319 + name = "libpulse-binding" 2320 + version = "2.27.1" 2321 + source = "registry+https://github.com/rust-lang/crates.io-index" 2322 + checksum = "1745b20bfc194ac12ef828f144f0ec2d4a7fe993281fa3567a0bd4969aee6890" 2323 + dependencies = [ 2324 + "bitflags", 2325 + "libc", 2326 + "libpulse-sys", 2327 + "num-derive", 2328 + "num-traits 0.2.15", 2329 + "winapi 0.3.9", 2330 + ] 2331 + 2332 + [[package]] 2333 + name = "libpulse-simple-binding" 2334 + version = "2.27.1" 2335 + source = "registry+https://github.com/rust-lang/crates.io-index" 2336 + checksum = "5ced94199e6e44133431374e4043f34e1f0697ebfb7b7d6c244a65bfaedf0e31" 2337 + dependencies = [ 2338 + "libpulse-binding", 2339 + "libpulse-simple-sys", 2340 + "libpulse-sys", 2341 + ] 2342 + 2343 + [[package]] 2344 + name = "libpulse-simple-sys" 2345 + version = "1.20.1" 2346 + source = "registry+https://github.com/rust-lang/crates.io-index" 2347 + checksum = "84e423d9c619c908ce9b4916080e65ab586ca55b8c4939379f15e6e72fb43842" 2348 + dependencies = [ 2349 + "libpulse-sys", 2350 + "pkg-config", 2351 + ] 2352 + 2353 + [[package]] 2354 + name = "libpulse-sys" 2355 + version = "1.20.1" 2356 + source = "registry+https://github.com/rust-lang/crates.io-index" 2357 + checksum = "2191e6880818d1df4cf72eac8e91dce7a5a52ba0da4b2a5cdafabc22b937eadb" 2358 + dependencies = [ 2359 + "libc", 2360 + "num-derive", 2361 + "num-traits 0.2.15", 2362 + "pkg-config", 2363 + "winapi 0.3.9", 2364 + ] 2365 + 2366 + [[package]] 2367 + name = "libsamplerate-sys" 2368 + version = "0.1.12" 2369 + source = "registry+https://github.com/rust-lang/crates.io-index" 2370 + checksum = "28853b399f78f8281cd88d333b54a63170c4275f6faea66726a2bea5cca72e0d" 2371 + dependencies = [ 2372 + "cmake", 2373 + ] 2374 + 2375 + [[package]] 2376 + name = "libsodium-sys" 2377 + version = "0.2.7" 2378 + source = "registry+https://github.com/rust-lang/crates.io-index" 2379 + checksum = "6b779387cd56adfbc02ea4a668e704f729be8d6a6abd2c27ca5ee537849a92fd" 2380 + dependencies = [ 2381 + "cc", 2382 + "libc", 2383 + "pkg-config", 2384 + "walkdir", 2385 + ] 2386 + 2387 + [[package]] 2388 + name = "link-cplusplus" 2389 + version = "1.0.8" 2390 + source = "registry+https://github.com/rust-lang/crates.io-index" 2391 + checksum = "ecd207c9c713c34f95a097a5b029ac2ce6010530c7b49d7fea24d977dede04f5" 2392 + dependencies = [ 2393 + "cc", 2394 + ] 2395 + 2396 + [[package]] 2397 + name = "linux-raw-sys" 2398 + version = "0.3.1" 2399 + source = "registry+https://github.com/rust-lang/crates.io-index" 2400 + checksum = "d59d8c75012853d2e872fb56bc8a2e53718e2cafe1a4c823143141c6d90c322f" 2401 + 2402 + [[package]] 2403 + name = "lock_api" 2404 + version = "0.4.9" 2405 + source = "registry+https://github.com/rust-lang/crates.io-index" 2406 + checksum = "435011366fe56583b16cf956f9df0095b405b82d76425bc8981c0e22e60ec4df" 2407 + dependencies = [ 2408 + "autocfg 1.1.0", 2409 + "scopeguard", 2410 + ] 2411 + 2412 + [[package]] 2413 + name = "log" 2414 + version = "0.4.17" 2415 + source = "registry+https://github.com/rust-lang/crates.io-index" 2416 + checksum = "abb12e687cfb44aa40f41fc3978ef76448f9b6038cad6aef4259d3c095a2382e" 2417 + dependencies = [ 2418 + "cfg-if 1.0.0", 2419 + ] 2420 + 2421 + [[package]] 2422 + name = "mac_address" 2423 + version = "1.1.4" 2424 + source = "registry+https://github.com/rust-lang/crates.io-index" 2425 + checksum = "b238e3235c8382b7653c6408ed1b08dd379bdb9fdf990fb0bbae3db2cc0ae963" 2426 + dependencies = [ 2427 + "nix 0.23.2", 2428 + "winapi 0.3.9", 2429 + ] 2430 + 2431 + [[package]] 2432 + name = "mach" 2433 + version = "0.3.2" 2434 + source = "registry+https://github.com/rust-lang/crates.io-index" 2435 + checksum = "b823e83b2affd8f40a9ee8c29dbc56404c1e34cd2710921f2801e2cf29527afa" 2436 + dependencies = [ 2437 + "libc", 2438 + ] 2439 + 2440 + [[package]] 2441 + name = "machine-uid" 2442 + version = "0.2.0" 2443 + source = "registry+https://github.com/rust-lang/crates.io-index" 2444 + checksum = "1f1595709b0a7386bcd56ba34d250d626e5503917d05d32cdccddcd68603e212" 2445 + dependencies = [ 2446 + "winreg 0.6.2", 2447 + ] 2448 + 2449 + [[package]] 2450 + name = "magnum-opus" 2451 + version = "0.4.0" 2452 + source = "git+https://github.com/TheRadioGuy/magnum-opus#171e1d021004626f7444d1e39b98f50bc3cb2604" 2453 + dependencies = [ 2454 + "libc", 2455 + "opusic-sys", 2456 + ] 2457 + 2458 + [[package]] 2459 + name = "malloc_buf" 2460 + version = "0.0.6" 2461 + source = "registry+https://github.com/rust-lang/crates.io-index" 2462 + checksum = "62bb907fe88d54d8d9ce32a3cceab4218ed2f6b7d35617cafe9adf84e43919cb" 2463 + dependencies = [ 2464 + "libc", 2465 + ] 2466 + 2467 + [[package]] 2468 + name = "memchr" 2469 + version = "2.5.0" 2470 + source = "registry+https://github.com/rust-lang/crates.io-index" 2471 + checksum = "2dffe52ecf27772e601905b7522cb4ef790d2cc203488bbd0e2fe85fcb74566d" 2472 + 2473 + [[package]] 2474 + name = "memoffset" 2475 + version = "0.6.5" 2476 + source = "registry+https://github.com/rust-lang/crates.io-index" 2477 + checksum = "5aa361d4faea93603064a027415f07bd8e1d5c88c9fbf68bf56a285428fd79ce" 2478 + dependencies = [ 2479 + "autocfg 1.1.0", 2480 + ] 2481 + 2482 + [[package]] 2483 + name = "memoffset" 2484 + version = "0.8.0" 2485 + source = "registry+https://github.com/rust-lang/crates.io-index" 2486 + checksum = "d61c719bcfbcf5d62b3a09efa6088de8c54bc0bfcd3ea7ae39fcc186108b8de1" 2487 + dependencies = [ 2488 + "autocfg 1.1.0", 2489 + ] 2490 + 2491 + [[package]] 2492 + name = "minimal-lexical" 2493 + version = "0.2.1" 2494 + source = "registry+https://github.com/rust-lang/crates.io-index" 2495 + checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a" 2496 + 2497 + [[package]] 2498 + name = "miniz_oxide" 2499 + version = "0.3.7" 2500 + source = "registry+https://github.com/rust-lang/crates.io-index" 2501 + checksum = "791daaae1ed6889560f8c4359194f56648355540573244a5448a83ba1ecc7435" 2502 + dependencies = [ 2503 + "adler32", 2504 + ] 2505 + 2506 + [[package]] 2507 + name = "miniz_oxide" 2508 + version = "0.4.4" 2509 + source = "registry+https://github.com/rust-lang/crates.io-index" 2510 + checksum = "a92518e98c078586bc6c934028adcca4c92a53d6a958196de835170a01d84e4b" 2511 + dependencies = [ 2512 + "adler", 2513 + "autocfg 1.1.0", 2514 + ] 2515 + 2516 + [[package]] 2517 + name = "miniz_oxide" 2518 + version = "0.6.2" 2519 + source = "registry+https://github.com/rust-lang/crates.io-index" 2520 + checksum = "b275950c28b37e794e8c55d88aeb5e139d0ce23fdbbeda68f8d7174abdf9e8fa" 2521 + dependencies = [ 2522 + "adler", 2523 + ] 2524 + 2525 + [[package]] 2526 + name = "mio" 2527 + version = "0.6.23" 2528 + source = "registry+https://github.com/rust-lang/crates.io-index" 2529 + checksum = "4afd66f5b91bf2a3bc13fad0e21caedac168ca4c707504e75585648ae80e4cc4" 2530 + dependencies = [ 2531 + "cfg-if 0.1.10", 2532 + "fuchsia-zircon", 2533 + "fuchsia-zircon-sys", 2534 + "iovec", 2535 + "kernel32-sys", 2536 + "libc", 2537 + "log", 2538 + "miow 0.2.2", 2539 + "net2", 2540 + "slab", 2541 + "winapi 0.2.8", 2542 + ] 2543 + 2544 + [[package]] 2545 + name = "mio" 2546 + version = "0.8.6" 2547 + source = "registry+https://github.com/rust-lang/crates.io-index" 2548 + checksum = "5b9d9a46eff5b4ff64b45a9e316a6d1e0bc719ef429cbec4dc630684212bfdf9" 2549 + dependencies = [ 2550 + "libc", 2551 + "log", 2552 + "wasi", 2553 + "windows-sys 0.45.0", 2554 + ] 2555 + 2556 + [[package]] 2557 + name = "mio-named-pipes" 2558 + version = "0.1.7" 2559 + source = "registry+https://github.com/rust-lang/crates.io-index" 2560 + checksum = "0840c1c50fd55e521b247f949c241c9997709f23bd7f023b9762cd561e935656" 2561 + dependencies = [ 2562 + "log", 2563 + "mio 0.6.23", 2564 + "miow 0.3.7", 2565 + "winapi 0.3.9", 2566 + ] 2567 + 2568 + [[package]] 2569 + name = "miow" 2570 + version = "0.2.2" 2571 + source = "registry+https://github.com/rust-lang/crates.io-index" 2572 + checksum = "ebd808424166322d4a38da87083bfddd3ac4c131334ed55856112eb06d46944d" 2573 + dependencies = [ 2574 + "kernel32-sys", 2575 + "net2", 2576 + "winapi 0.2.8", 2577 + "ws2_32-sys", 2578 + ] 2579 + 2580 + [[package]] 2581 + name = "miow" 2582 + version = "0.3.7" 2583 + source = "registry+https://github.com/rust-lang/crates.io-index" 2584 + checksum = "b9f1c5b025cda876f66ef43a113f91ebc9f4ccef34843000e0adf6ebbab84e21" 2585 + dependencies = [ 2586 + "winapi 0.3.9", 2587 + ] 2588 + 2589 + [[package]] 2590 + name = "miow" 2591 + version = "0.4.0" 2592 + source = "registry+https://github.com/rust-lang/crates.io-index" 2593 + checksum = "a7377f7792b3afb6a3cba68daa54ca23c032137010460d667fda53a8d66be00e" 2594 + dependencies = [ 2595 + "windows-sys 0.28.0", 2596 + ] 2597 + 2598 + [[package]] 2599 + name = "muldiv" 2600 + version = "0.2.1" 2601 + source = "registry+https://github.com/rust-lang/crates.io-index" 2602 + checksum = "0419348c027fa7be448d2ae7ea0e4e04c2334c31dc4e74ab29f00a2a7ca69204" 2603 + 2604 + [[package]] 2605 + name = "ndk" 2606 + version = "0.6.0" 2607 + source = "registry+https://github.com/rust-lang/crates.io-index" 2608 + checksum = "2032c77e030ddee34a6787a64166008da93f6a352b629261d0fee232b8742dd4" 2609 + dependencies = [ 2610 + "bitflags", 2611 + "jni-sys", 2612 + "ndk-sys", 2613 + "num_enum", 2614 + "thiserror", 2615 + ] 2616 + 2617 + [[package]] 2618 + name = "ndk-context" 2619 + version = "0.1.1" 2620 + source = "registry+https://github.com/rust-lang/crates.io-index" 2621 + checksum = "27b02d87554356db9e9a873add8782d4ea6e3e58ea071a9adb9a2e8ddb884a8b" 2622 + 2623 + [[package]] 2624 + name = "ndk-glue" 2625 + version = "0.6.2" 2626 + source = "registry+https://github.com/rust-lang/crates.io-index" 2627 + checksum = "0d0c4a7b83860226e6b4183edac21851f05d5a51756e97a1144b7f5a6b63e65f" 2628 + dependencies = [ 2629 + "lazy_static", 2630 + "libc", 2631 + "log", 2632 + "ndk", 2633 + "ndk-context", 2634 + "ndk-macro", 2635 + "ndk-sys", 2636 + ] 2637 + 2638 + [[package]] 2639 + name = "ndk-macro" 2640 + version = "0.3.0" 2641 + source = "registry+https://github.com/rust-lang/crates.io-index" 2642 + checksum = "0df7ac00c4672f9d5aece54ee3347520b7e20f158656c7db2e6de01902eb7a6c" 2643 + dependencies = [ 2644 + "darling", 2645 + "proc-macro-crate 1.3.1", 2646 + "proc-macro2", 2647 + "quote", 2648 + "syn 1.0.109", 2649 + ] 2650 + 2651 + [[package]] 2652 + name = "ndk-sys" 2653 + version = "0.3.0" 2654 + source = "registry+https://github.com/rust-lang/crates.io-index" 2655 + checksum = "6e5a6ae77c8ee183dcbbba6150e2e6b9f3f4196a7666c02a715a95692ec1fa97" 2656 + dependencies = [ 2657 + "jni-sys", 2658 + ] 2659 + 2660 + [[package]] 2661 + name = "net2" 2662 + version = "0.2.38" 2663 + source = "registry+https://github.com/rust-lang/crates.io-index" 2664 + checksum = "74d0df99cfcd2530b2e694f6e17e7f37b8e26bb23983ac530c0c97408837c631" 2665 + dependencies = [ 2666 + "cfg-if 0.1.10", 2667 + "libc", 2668 + "winapi 0.3.9", 2669 + ] 2670 + 2671 + [[package]] 2672 + name = "nix" 2673 + version = "0.22.3" 2674 + source = "registry+https://github.com/rust-lang/crates.io-index" 2675 + checksum = "e4916f159ed8e5de0082076562152a76b7a1f64a01fd9d1e0fea002c37624faf" 2676 + dependencies = [ 2677 + "bitflags", 2678 + "cc", 2679 + "cfg-if 1.0.0", 2680 + "libc", 2681 + "memoffset 0.6.5", 2682 + ] 2683 + 2684 + [[package]] 2685 + name = "nix" 2686 + version = "0.23.2" 2687 + source = "registry+https://github.com/rust-lang/crates.io-index" 2688 + checksum = "8f3790c00a0150112de0f4cd161e3d7fc4b2d8a5542ffc35f099a2562aecb35c" 2689 + dependencies = [ 2690 + "bitflags", 2691 + "cc", 2692 + "cfg-if 1.0.0", 2693 + "libc", 2694 + "memoffset 0.6.5", 2695 + ] 2696 + 2697 + [[package]] 2698 + name = "nix" 2699 + version = "0.26.2" 2700 + source = "registry+https://github.com/rust-lang/crates.io-index" 2701 + checksum = "bfdda3d196821d6af13126e40375cdf7da646a96114af134d5f417a9a1dc8e1a" 2702 + dependencies = [ 2703 + "bitflags", 2704 + "cfg-if 1.0.0", 2705 + "libc", 2706 + "static_assertions", 2707 + ] 2708 + 2709 + [[package]] 2710 + name = "nom" 2711 + version = "7.1.3" 2712 + source = "registry+https://github.com/rust-lang/crates.io-index" 2713 + checksum = "d273983c5a657a70a3e8f2a01329822f3b8c8172b73826411a55751e404a0a4a" 2714 + dependencies = [ 2715 + "memchr", 2716 + "minimal-lexical", 2717 + ] 2718 + 2719 + [[package]] 2720 + name = "ntapi" 2721 + version = "0.3.7" 2722 + source = "registry+https://github.com/rust-lang/crates.io-index" 2723 + checksum = "c28774a7fd2fbb4f0babd8237ce554b73af68021b5f695a3cebd6c59bac0980f" 2724 + dependencies = [ 2725 + "winapi 0.3.9", 2726 + ] 2727 + 2728 + [[package]] 2729 + name = "num-complex" 2730 + version = "0.4.3" 2731 + source = "registry+https://github.com/rust-lang/crates.io-index" 2732 + checksum = "02e0d21255c828d6f128a1e41534206671e8c3ea0c62f32291e808dc82cff17d" 2733 + dependencies = [ 2734 + "num-traits 0.2.15", 2735 + ] 2736 + 2737 + [[package]] 2738 + name = "num-derive" 2739 + version = "0.3.3" 2740 + source = "registry+https://github.com/rust-lang/crates.io-index" 2741 + checksum = "876a53fff98e03a936a674b29568b0e605f06b29372c2489ff4de23f1949743d" 2742 + dependencies = [ 2743 + "proc-macro2", 2744 + "quote", 2745 + "syn 1.0.109", 2746 + ] 2747 + 2748 + [[package]] 2749 + name = "num-integer" 2750 + version = "0.1.45" 2751 + source = "registry+https://github.com/rust-lang/crates.io-index" 2752 + checksum = "225d3389fb3509a24c93f5c29eb6bde2586b98d9f016636dff58d7c6f7569cd9" 2753 + dependencies = [ 2754 + "autocfg 1.1.0", 2755 + "num-traits 0.2.15", 2756 + ] 2757 + 2758 + [[package]] 2759 + name = "num-iter" 2760 + version = "0.1.43" 2761 + source = "registry+https://github.com/rust-lang/crates.io-index" 2762 + checksum = "7d03e6c028c5dc5cac6e2dec0efda81fc887605bb3d884578bb6d6bf7514e252" 2763 + dependencies = [ 2764 + "autocfg 1.1.0", 2765 + "num-integer", 2766 + "num-traits 0.2.15", 2767 + ] 2768 + 2769 + [[package]] 2770 + name = "num-rational" 2771 + version = "0.3.2" 2772 + source = "registry+https://github.com/rust-lang/crates.io-index" 2773 + checksum = "12ac428b1cb17fce6f731001d307d351ec70a6d202fc2e60f7d4c5e42d8f4f07" 2774 + dependencies = [ 2775 + "autocfg 1.1.0", 2776 + "num-integer", 2777 + "num-traits 0.2.15", 2778 + ] 2779 + 2780 + [[package]] 2781 + name = "num-traits" 2782 + version = "0.1.43" 2783 + source = "registry+https://github.com/rust-lang/crates.io-index" 2784 + checksum = "92e5113e9fd4cc14ded8e499429f396a20f98c772a47cc8622a736e1ec843c31" 2785 + dependencies = [ 2786 + "num-traits 0.2.15", 2787 + ] 2788 + 2789 + [[package]] 2790 + name = "num-traits" 2791 + version = "0.2.15" 2792 + source = "registry+https://github.com/rust-lang/crates.io-index" 2793 + checksum = "578ede34cf02f8924ab9447f50c28075b4d3e5b269972345e7e0372b38c6cdcd" 2794 + dependencies = [ 2795 + "autocfg 1.1.0", 2796 + ] 2797 + 2798 + [[package]] 2799 + name = "num_cpus" 2800 + version = "1.15.0" 2801 + source = "registry+https://github.com/rust-lang/crates.io-index" 2802 + checksum = "0fac9e2da13b5eb447a6ce3d392f23a29d8694bff781bf03a16cd9ac8697593b" 2803 + dependencies = [ 2804 + "hermit-abi 0.2.6", 2805 + "libc", 2806 + ] 2807 + 2808 + [[package]] 2809 + name = "num_enum" 2810 + version = "0.5.11" 2811 + source = "registry+https://github.com/rust-lang/crates.io-index" 2812 + checksum = "1f646caf906c20226733ed5b1374287eb97e3c2a5c227ce668c1f2ce20ae57c9" 2813 + dependencies = [ 2814 + "num_enum_derive", 2815 + ] 2816 + 2817 + [[package]] 2818 + name = "num_enum_derive" 2819 + version = "0.5.11" 2820 + source = "registry+https://github.com/rust-lang/crates.io-index" 2821 + checksum = "dcbff9bc912032c62bf65ef1d5aea88983b420f4f839db1e9b0c281a25c9c799" 2822 + dependencies = [ 2823 + "proc-macro-crate 1.3.1", 2824 + "proc-macro2", 2825 + "quote", 2826 + "syn 1.0.109", 2827 + ] 2828 + 2829 + [[package]] 2830 + name = "num_threads" 2831 + version = "0.1.6" 2832 + source = "registry+https://github.com/rust-lang/crates.io-index" 2833 + checksum = "2819ce041d2ee131036f4fc9d6ae7ae125a3a40e97ba64d04fe799ad9dabbb44" 2834 + dependencies = [ 2835 + "libc", 2836 + ] 2837 + 2838 + [[package]] 2839 + name = "objc" 2840 + version = "0.2.7" 2841 + source = "registry+https://github.com/rust-lang/crates.io-index" 2842 + checksum = "915b1b472bc21c53464d6c8461c9d3af805ba1ef837e1cac254428f4a77177b1" 2843 + dependencies = [ 2844 + "malloc_buf", 2845 + ] 2846 + 2847 + [[package]] 2848 + name = "objc-foundation" 2849 + version = "0.1.1" 2850 + source = "registry+https://github.com/rust-lang/crates.io-index" 2851 + checksum = "1add1b659e36c9607c7aab864a76c7a4c2760cd0cd2e120f3fb8b952c7e22bf9" 2852 + dependencies = [ 2853 + "block", 2854 + "objc", 2855 + "objc_id", 2856 + ] 2857 + 2858 + [[package]] 2859 + name = "objc_id" 2860 + version = "0.1.1" 2861 + source = "registry+https://github.com/rust-lang/crates.io-index" 2862 + checksum = "c92d4ddb4bd7b50d730c215ff871754d0da6b2178849f8a2a2ab69712d0c073b" 2863 + dependencies = [ 2864 + "objc", 2865 + ] 2866 + 2867 + [[package]] 2868 + name = "object" 2869 + version = "0.30.3" 2870 + source = "registry+https://github.com/rust-lang/crates.io-index" 2871 + checksum = "ea86265d3d3dcb6a27fc51bd29a4bf387fae9d2986b823079d4986af253eb439" 2872 + dependencies = [ 2873 + "memchr", 2874 + ] 2875 + 2876 + [[package]] 2877 + name = "oboe" 2878 + version = "0.4.6" 2879 + source = "registry+https://github.com/rust-lang/crates.io-index" 2880 + checksum = "27f63c358b4fa0fbcfefd7c8be5cfc39c08ce2389f5325687e7762a48d30a5c1" 2881 + dependencies = [ 2882 + "jni", 2883 + "ndk", 2884 + "ndk-context", 2885 + "num-derive", 2886 + "num-traits 0.2.15", 2887 + "oboe-sys", 2888 + ] 2889 + 2890 + [[package]] 2891 + name = "oboe-sys" 2892 + version = "0.4.5" 2893 + source = "registry+https://github.com/rust-lang/crates.io-index" 2894 + checksum = "3370abb7372ed744232c12954d920d1a40f1c4686de9e79e800021ef492294bd" 2895 + dependencies = [ 2896 + "cc", 2897 + ] 2898 + 2899 + [[package]] 2900 + name = "once_cell" 2901 + version = "1.17.1" 2902 + source = "registry+https://github.com/rust-lang/crates.io-index" 2903 + checksum = "b7e5500299e16ebb147ae15a00a942af264cf3688f47923b8fc2cd5858f23ad3" 2904 + 2905 + [[package]] 2906 + name = "openssl-probe" 2907 + version = "0.1.5" 2908 + source = "registry+https://github.com/rust-lang/crates.io-index" 2909 + checksum = "ff011a302c396a5197692431fc1948019154afc178baf7d8e37367442a4601cf" 2910 + 2911 + [[package]] 2912 + name = "opusic-sys" 2913 + version = "0.3.6" 2914 + source = "registry+https://github.com/rust-lang/crates.io-index" 2915 + checksum = "5eace752ce07a037241dba8f02c654799f051e431b27028056bcb480e83b54f5" 2916 + dependencies = [ 2917 + "cmake", 2918 + "libc", 2919 + ] 2920 + 2921 + [[package]] 2922 + name = "os_str_bytes" 2923 + version = "6.5.0" 2924 + source = "registry+https://github.com/rust-lang/crates.io-index" 2925 + checksum = "ceedf44fb00f2d1984b0bc98102627ce622e083e49a5bacdb3e514fa4238e267" 2926 + 2927 + [[package]] 2928 + name = "padlock" 2929 + version = "0.2.0" 2930 + source = "registry+https://github.com/rust-lang/crates.io-index" 2931 + checksum = "c10569378a1dacd9f30dbe7ae49e054d2c45dc2f8ee49899903e09c3924e8b6f" 2932 + 2933 + [[package]] 2934 + name = "pango" 2935 + version = "0.9.1" 2936 + source = "registry+https://github.com/rust-lang/crates.io-index" 2937 + checksum = "9937068580bebd8ced19975938573803273ccbcbd598c58d4906efd4ac87c438" 2938 + dependencies = [ 2939 + "bitflags", 2940 + "glib 0.10.3", 2941 + "glib-sys 0.10.1", 2942 + "gobject-sys 0.10.0", 2943 + "libc", 2944 + "once_cell", 2945 + "pango-sys 0.10.0", 2946 + ] 2947 + 2948 + [[package]] 2949 + name = "pango" 2950 + version = "0.15.10" 2951 + source = "registry+https://github.com/rust-lang/crates.io-index" 2952 + checksum = "22e4045548659aee5313bde6c582b0d83a627b7904dd20dc2d9ef0895d414e4f" 2953 + dependencies = [ 2954 + "bitflags", 2955 + "glib 0.15.12", 2956 + "libc", 2957 + "once_cell", 2958 + "pango-sys 0.15.10", 2959 + ] 2960 + 2961 + [[package]] 2962 + name = "pango-sys" 2963 + version = "0.10.0" 2964 + source = "registry+https://github.com/rust-lang/crates.io-index" 2965 + checksum = "24d2650c8b62d116c020abd0cea26a4ed96526afda89b1c4ea567131fdefc890" 2966 + dependencies = [ 2967 + "glib-sys 0.10.1", 2968 + "gobject-sys 0.10.0", 2969 + "libc", 2970 + "system-deps 1.3.2", 2971 + ] 2972 + 2973 + [[package]] 2974 + name = "pango-sys" 2975 + version = "0.15.10" 2976 + source = "registry+https://github.com/rust-lang/crates.io-index" 2977 + checksum = "d2a00081cde4661982ed91d80ef437c20eacaf6aa1a5962c0279ae194662c3aa" 2978 + dependencies = [ 2979 + "glib-sys 0.15.10", 2980 + "gobject-sys 0.15.10", 2981 + "libc", 2982 + "system-deps 6.0.4", 2983 + ] 2984 + 2985 + [[package]] 2986 + name = "parity-tokio-ipc" 2987 + version = "0.7.3-1" 2988 + source = "git+https://github.com/open-trade/parity-tokio-ipc#20b2895910161605210657f3e751edd55321f698" 2989 + dependencies = [ 2990 + "futures", 2991 + "libc", 2992 + "log", 2993 + "mio-named-pipes", 2994 + "miow 0.4.0", 2995 + "rand 0.8.5", 2996 + "tokio", 2997 + "winapi 0.3.9", 2998 + ] 2999 + 3000 + [[package]] 3001 + name = "parking_lot" 3002 + version = "0.11.2" 3003 + source = "registry+https://github.com/rust-lang/crates.io-index" 3004 + checksum = "7d17b78036a60663b797adeaee46f5c9dfebb86948d1255007a1d6be0271ff99" 3005 + dependencies = [ 3006 + "instant", 3007 + "lock_api", 3008 + "parking_lot_core 0.8.6", 3009 + ] 3010 + 3011 + [[package]] 3012 + name = "parking_lot" 3013 + version = "0.12.1" 3014 + source = "registry+https://github.com/rust-lang/crates.io-index" 3015 + checksum = "3742b2c103b9f06bc9fff0a37ff4912935851bee6d36f3c02bcc755bcfec228f" 3016 + dependencies = [ 3017 + "lock_api", 3018 + "parking_lot_core 0.9.7", 3019 + ] 3020 + 3021 + [[package]] 3022 + name = "parking_lot_core" 3023 + version = "0.8.6" 3024 + source = "registry+https://github.com/rust-lang/crates.io-index" 3025 + checksum = "60a2cfe6f0ad2bfc16aefa463b497d5c7a5ecd44a23efa72aa342d90177356dc" 3026 + dependencies = [ 3027 + "cfg-if 1.0.0", 3028 + "instant", 3029 + "libc", 3030 + "redox_syscall 0.2.16", 3031 + "smallvec", 3032 + "winapi 0.3.9", 3033 + ] 3034 + 3035 + [[package]] 3036 + name = "parking_lot_core" 3037 + version = "0.9.7" 3038 + source = "registry+https://github.com/rust-lang/crates.io-index" 3039 + checksum = "9069cbb9f99e3a5083476ccb29ceb1de18b9118cafa53e90c9551235de2b9521" 3040 + dependencies = [ 3041 + "cfg-if 1.0.0", 3042 + "libc", 3043 + "redox_syscall 0.2.16", 3044 + "smallvec", 3045 + "windows-sys 0.45.0", 3046 + ] 3047 + 3048 + [[package]] 3049 + name = "paste" 3050 + version = "1.0.12" 3051 + source = "registry+https://github.com/rust-lang/crates.io-index" 3052 + checksum = "9f746c4065a8fa3fe23974dd82f15431cc8d40779821001404d10d2e79ca7d79" 3053 + 3054 + [[package]] 3055 + name = "peeking_take_while" 3056 + version = "0.1.2" 3057 + source = "registry+https://github.com/rust-lang/crates.io-index" 3058 + checksum = "19b17cddbe7ec3f8bc800887bab5e717348c95ea2ca0b1bf0837fb964dc67099" 3059 + 3060 + [[package]] 3061 + name = "phf" 3062 + version = "0.7.24" 3063 + source = "registry+https://github.com/rust-lang/crates.io-index" 3064 + checksum = "b3da44b85f8e8dfaec21adae67f95d93244b2ecf6ad2a692320598dcc8e6dd18" 3065 + dependencies = [ 3066 + "phf_shared", 3067 + ] 3068 + 3069 + [[package]] 3070 + name = "phf_codegen" 3071 + version = "0.7.24" 3072 + source = "registry+https://github.com/rust-lang/crates.io-index" 3073 + checksum = "b03e85129e324ad4166b06b2c7491ae27fe3ec353af72e72cd1654c7225d517e" 3074 + dependencies = [ 3075 + "phf_generator", 3076 + "phf_shared", 3077 + ] 3078 + 3079 + [[package]] 3080 + name = "phf_generator" 3081 + version = "0.7.24" 3082 + source = "registry+https://github.com/rust-lang/crates.io-index" 3083 + checksum = "09364cc93c159b8b06b1f4dd8a4398984503483891b0c26b867cf431fb132662" 3084 + dependencies = [ 3085 + "phf_shared", 3086 + "rand 0.6.5", 3087 + ] 3088 + 3089 + [[package]] 3090 + name = "phf_shared" 3091 + version = "0.7.24" 3092 + source = "registry+https://github.com/rust-lang/crates.io-index" 3093 + checksum = "234f71a15de2288bcb7e3b6515828d22af7ec8598ee6d24c3b526fa0a80b67a0" 3094 + dependencies = [ 3095 + "siphasher", 3096 + ] 3097 + 3098 + [[package]] 3099 + name = "pin-project" 3100 + version = "1.0.12" 3101 + source = "registry+https://github.com/rust-lang/crates.io-index" 3102 + checksum = "ad29a609b6bcd67fee905812e544992d216af9d755757c05ed2d0e15a74c6ecc" 3103 + dependencies = [ 3104 + "pin-project-internal", 3105 + ] 3106 + 3107 + [[package]] 3108 + name = "pin-project-internal" 3109 + version = "1.0.12" 3110 + source = "registry+https://github.com/rust-lang/crates.io-index" 3111 + checksum = "069bdb1e05adc7a8990dce9cc75370895fbe4e3d58b9b73bf1aee56359344a55" 3112 + dependencies = [ 3113 + "proc-macro2", 3114 + "quote", 3115 + "syn 1.0.109", 3116 + ] 3117 + 3118 + [[package]] 3119 + name = "pin-project-lite" 3120 + version = "0.2.9" 3121 + source = "registry+https://github.com/rust-lang/crates.io-index" 3122 + checksum = "e0a7ae3ac2f1173085d398531c705756c94a4c56843785df85a60c1a0afac116" 3123 + 3124 + [[package]] 3125 + name = "pin-utils" 3126 + version = "0.1.0" 3127 + source = "registry+https://github.com/rust-lang/crates.io-index" 3128 + checksum = "8b870d8c151b6f2fb93e84a13146138f05d02ed11c7e7c54f8826aaaf7c9f184" 3129 + 3130 + [[package]] 3131 + name = "pkg-config" 3132 + version = "0.3.26" 3133 + source = "registry+https://github.com/rust-lang/crates.io-index" 3134 + checksum = "6ac9a59f73473f1b8d852421e59e64809f025994837ef743615c6d0c5b305160" 3135 + 3136 + [[package]] 3137 + name = "png" 3138 + version = "0.16.8" 3139 + source = "registry+https://github.com/rust-lang/crates.io-index" 3140 + checksum = "3c3287920cb847dee3de33d301c463fba14dda99db24214ddf93f83d3021f4c6" 3141 + dependencies = [ 3142 + "bitflags", 3143 + "crc32fast", 3144 + "deflate", 3145 + "miniz_oxide 0.3.7", 3146 + ] 3147 + 3148 + [[package]] 3149 + name = "ppv-lite86" 3150 + version = "0.2.17" 3151 + source = "registry+https://github.com/rust-lang/crates.io-index" 3152 + checksum = "5b40af805b3121feab8a3c29f04d8ad262fa8e0561883e7653e024ae4479e6de" 3153 + 3154 + [[package]] 3155 + name = "pretty-hex" 3156 + version = "0.2.1" 3157 + source = "registry+https://github.com/rust-lang/crates.io-index" 3158 + checksum = "bc5c99d529f0d30937f6f4b8a86d988047327bb88d04d2c4afc356de74722131" 3159 + 3160 + [[package]] 3161 + name = "primal-check" 3162 + version = "0.3.3" 3163 + source = "registry+https://github.com/rust-lang/crates.io-index" 3164 + checksum = "9df7f93fd637f083201473dab4fee2db4c429d32e55e3299980ab3957ab916a0" 3165 + dependencies = [ 3166 + "num-integer", 3167 + ] 3168 + 3169 + [[package]] 3170 + name = "proc-macro-crate" 3171 + version = "0.1.5" 3172 + source = "registry+https://github.com/rust-lang/crates.io-index" 3173 + checksum = "1d6ea3c4595b96363c13943497db34af4460fb474a95c43f4446ad341b8c9785" 3174 + dependencies = [ 3175 + "toml 0.5.11", 3176 + ] 3177 + 3178 + [[package]] 3179 + name = "proc-macro-crate" 3180 + version = "1.3.1" 3181 + source = "registry+https://github.com/rust-lang/crates.io-index" 3182 + checksum = "7f4c021e1093a56626774e81216a4ce732a735e5bad4868a03f3ed65ca0c3919" 3183 + dependencies = [ 3184 + "once_cell", 3185 + "toml_edit", 3186 + ] 3187 + 3188 + [[package]] 3189 + name = "proc-macro-error" 3190 + version = "1.0.4" 3191 + source = "registry+https://github.com/rust-lang/crates.io-index" 3192 + checksum = "da25490ff9892aab3fcf7c36f08cfb902dd3e71ca0f9f9517bea02a73a5ce38c" 3193 + dependencies = [ 3194 + "proc-macro-error-attr", 3195 + "proc-macro2", 3196 + "quote", 3197 + "syn 1.0.109", 3198 + "version_check", 3199 + ] 3200 + 3201 + [[package]] 3202 + name = "proc-macro-error-attr" 3203 + version = "1.0.4" 3204 + source = "registry+https://github.com/rust-lang/crates.io-index" 3205 + checksum = "a1be40180e52ecc98ad80b184934baf3d0d29f979574e439af5a55274b35f869" 3206 + dependencies = [ 3207 + "proc-macro2", 3208 + "quote", 3209 + "version_check", 3210 + ] 3211 + 3212 + [[package]] 3213 + name = "proc-macro2" 3214 + version = "1.0.56" 3215 + source = "registry+https://github.com/rust-lang/crates.io-index" 3216 + checksum = "2b63bdb0cd06f1f4dedf69b254734f9b45af66e4a031e42a7480257d9898b435" 3217 + dependencies = [ 3218 + "unicode-ident", 3219 + ] 3220 + 3221 + [[package]] 3222 + name = "protobuf" 3223 + version = "3.0.0-alpha.2" 3224 + source = "registry+https://github.com/rust-lang/crates.io-index" 3225 + checksum = "9d5ef59c35c7472ce5e1b6c5924b87585143d1fc2cf39eae0009bba6c4df62f1" 3226 + 3227 + [[package]] 3228 + name = "protobuf-codegen" 3229 + version = "3.0.0-alpha.2" 3230 + source = "registry+https://github.com/rust-lang/crates.io-index" 3231 + checksum = "89100ee819f69b77a4cab389fec9dd155a305af4c615e6413ec1ef9341f333ef" 3232 + dependencies = [ 3233 + "anyhow", 3234 + "protobuf", 3235 + "protobuf-parse", 3236 + "thiserror", 3237 + ] 3238 + 3239 + [[package]] 3240 + name = "protobuf-codegen-pure" 3241 + version = "3.0.0-alpha.2" 3242 + source = "registry+https://github.com/rust-lang/crates.io-index" 3243 + checksum = "79453e74d08190551e821533ee42c447f9e21ca26f83520e120e6e8af27f6879" 3244 + dependencies = [ 3245 + "anyhow", 3246 + "protobuf", 3247 + "protobuf-codegen", 3248 + "protobuf-parse", 3249 + "thiserror", 3250 + ] 3251 + 3252 + [[package]] 3253 + name = "protobuf-parse" 3254 + version = "3.0.0-alpha.2" 3255 + source = "registry+https://github.com/rust-lang/crates.io-index" 3256 + checksum = "c265ffc69976efc3056955b881641add3186ad0be893ef10622482d80d1d2b68" 3257 + dependencies = [ 3258 + "anyhow", 3259 + "protobuf", 3260 + "protoc", 3261 + "tempfile", 3262 + "thiserror", 3263 + ] 3264 + 3265 + [[package]] 3266 + name = "protoc" 3267 + version = "3.0.0-alpha.2" 3268 + source = "registry+https://github.com/rust-lang/crates.io-index" 3269 + checksum = "1f1f8b318a54d18fbe542513331e058f4f8ce6502e542e057c50c7e5e803fdab" 3270 + dependencies = [ 3271 + "anyhow", 3272 + "log", 3273 + "thiserror", 3274 + "which 4.4.0", 3275 + ] 3276 + 3277 + [[package]] 3278 + name = "quest" 3279 + version = "0.3.0" 3280 + source = "registry+https://github.com/rust-lang/crates.io-index" 3281 + checksum = "556af5f5c953a2ee13f45753e581a38f9778e6551bc3ccc56d90b14628fe59d8" 3282 + dependencies = [ 3283 + "cfg-if 0.1.10", 3284 + "rpassword 2.1.0", 3285 + "tempfile", 3286 + "termios", 3287 + "winapi 0.3.9", 3288 + ] 3289 + 3290 + [[package]] 3291 + name = "quick-xml" 3292 + version = "0.22.0" 3293 + source = "registry+https://github.com/rust-lang/crates.io-index" 3294 + checksum = "8533f14c8382aaad0d592c812ac3b826162128b65662331e1127b45c3d18536b" 3295 + dependencies = [ 3296 + "memchr", 3297 + ] 3298 + 3299 + [[package]] 3300 + name = "quinn" 3301 + version = "0.8.5" 3302 + source = "registry+https://github.com/rust-lang/crates.io-index" 3303 + checksum = "5b435e71d9bfa0d8889927231970c51fb89c58fa63bffcab117c9c7a41e5ef8f" 3304 + dependencies = [ 3305 + "bytes", 3306 + "futures-channel", 3307 + "futures-util", 3308 + "fxhash", 3309 + "quinn-proto", 3310 + "quinn-udp", 3311 + "rustls", 3312 + "thiserror", 3313 + "tokio", 3314 + "tracing", 3315 + "webpki", 3316 + ] 3317 + 3318 + [[package]] 3319 + name = "quinn-proto" 3320 + version = "0.8.4" 3321 + source = "registry+https://github.com/rust-lang/crates.io-index" 3322 + checksum = "3fce546b9688f767a57530652488420d419a8b1f44a478b451c3d1ab6d992a55" 3323 + dependencies = [ 3324 + "bytes", 3325 + "fxhash", 3326 + "rand 0.8.5", 3327 + "ring", 3328 + "rustls", 3329 + "rustls-native-certs", 3330 + "rustls-pemfile 0.2.1", 3331 + "slab", 3332 + "thiserror", 3333 + "tinyvec", 3334 + "tracing", 3335 + "webpki", 3336 + ] 3337 + 3338 + [[package]] 3339 + name = "quinn-udp" 3340 + version = "0.1.4" 3341 + source = "registry+https://github.com/rust-lang/crates.io-index" 3342 + checksum = "b07946277141531aea269befd949ed16b2c85a780ba1043244eda0969e538e54" 3343 + dependencies = [ 3344 + "futures-util", 3345 + "libc", 3346 + "quinn-proto", 3347 + "socket2 0.4.9", 3348 + "tokio", 3349 + "tracing", 3350 + ] 3351 + 3352 + [[package]] 3353 + name = "quote" 3354 + version = "1.0.26" 3355 + source = "registry+https://github.com/rust-lang/crates.io-index" 3356 + checksum = "4424af4bf778aae2051a77b60283332f386554255d722233d09fbfc7e30da2fc" 3357 + dependencies = [ 3358 + "proc-macro2", 3359 + ] 3360 + 3361 + [[package]] 3362 + name = "rand" 3363 + version = "0.6.5" 3364 + source = "registry+https://github.com/rust-lang/crates.io-index" 3365 + checksum = "6d71dacdc3c88c1fde3885a3be3fbab9f35724e6ce99467f7d9c5026132184ca" 3366 + dependencies = [ 3367 + "autocfg 0.1.8", 3368 + "libc", 3369 + "rand_chacha 0.1.1", 3370 + "rand_core 0.4.2", 3371 + "rand_hc", 3372 + "rand_isaac", 3373 + "rand_jitter", 3374 + "rand_os", 3375 + "rand_pcg", 3376 + "rand_xorshift", 3377 + "winapi 0.3.9", 3378 + ] 3379 + 3380 + [[package]] 3381 + name = "rand" 3382 + version = "0.8.5" 3383 + source = "registry+https://github.com/rust-lang/crates.io-index" 3384 + checksum = "34af8d1a0e25924bc5b7c43c079c942339d8f0a8b57c39049bef581b46327404" 3385 + dependencies = [ 3386 + "libc", 3387 + "rand_chacha 0.3.1", 3388 + "rand_core 0.6.4", 3389 + ] 3390 + 3391 + [[package]] 3392 + name = "rand_chacha" 3393 + version = "0.1.1" 3394 + source = "registry+https://github.com/rust-lang/crates.io-index" 3395 + checksum = "556d3a1ca6600bfcbab7c7c91ccb085ac7fbbcd70e008a98742e7847f4f7bcef" 3396 + dependencies = [ 3397 + "autocfg 0.1.8", 3398 + "rand_core 0.3.1", 3399 + ] 3400 + 3401 + [[package]] 3402 + name = "rand_chacha" 3403 + version = "0.3.1" 3404 + source = "registry+https://github.com/rust-lang/crates.io-index" 3405 + checksum = "e6c10a63a0fa32252be49d21e7709d4d4baf8d231c2dbce1eaa8141b9b127d88" 3406 + dependencies = [ 3407 + "ppv-lite86", 3408 + "rand_core 0.6.4", 3409 + ] 3410 + 3411 + [[package]] 3412 + name = "rand_core" 3413 + version = "0.3.1" 3414 + source = "registry+https://github.com/rust-lang/crates.io-index" 3415 + checksum = "7a6fdeb83b075e8266dcc8762c22776f6877a63111121f5f8c7411e5be7eed4b" 3416 + dependencies = [ 3417 + "rand_core 0.4.2", 3418 + ] 3419 + 3420 + [[package]] 3421 + name = "rand_core" 3422 + version = "0.4.2" 3423 + source = "registry+https://github.com/rust-lang/crates.io-index" 3424 + checksum = "9c33a3c44ca05fa6f1807d8e6743f3824e8509beca625669633be0acbdf509dc" 3425 + 3426 + [[package]] 3427 + name = "rand_core" 3428 + version = "0.6.4" 3429 + source = "registry+https://github.com/rust-lang/crates.io-index" 3430 + checksum = "ec0be4795e2f6a28069bec0b5ff3e2ac9bafc99e6a9a7dc3547996c5c816922c" 3431 + dependencies = [ 3432 + "getrandom", 3433 + ] 3434 + 3435 + [[package]] 3436 + name = "rand_hc" 3437 + version = "0.1.0" 3438 + source = "registry+https://github.com/rust-lang/crates.io-index" 3439 + checksum = "7b40677c7be09ae76218dc623efbf7b18e34bced3f38883af07bb75630a21bc4" 3440 + dependencies = [ 3441 + "rand_core 0.3.1", 3442 + ] 3443 + 3444 + [[package]] 3445 + name = "rand_isaac" 3446 + version = "0.1.1" 3447 + source = "registry+https://github.com/rust-lang/crates.io-index" 3448 + checksum = "ded997c9d5f13925be2a6fd7e66bf1872597f759fd9dd93513dd7e92e5a5ee08" 3449 + dependencies = [ 3450 + "rand_core 0.3.1", 3451 + ] 3452 + 3453 + [[package]] 3454 + name = "rand_jitter" 3455 + version = "0.1.4" 3456 + source = "registry+https://github.com/rust-lang/crates.io-index" 3457 + checksum = "1166d5c91dc97b88d1decc3285bb0a99ed84b05cfd0bc2341bdf2d43fc41e39b" 3458 + dependencies = [ 3459 + "libc", 3460 + "rand_core 0.4.2", 3461 + "winapi 0.3.9", 3462 + ] 3463 + 3464 + [[package]] 3465 + name = "rand_os" 3466 + version = "0.1.3" 3467 + source = "registry+https://github.com/rust-lang/crates.io-index" 3468 + checksum = "7b75f676a1e053fc562eafbb47838d67c84801e38fc1ba459e8f180deabd5071" 3469 + dependencies = [ 3470 + "cloudabi", 3471 + "fuchsia-cprng", 3472 + "libc", 3473 + "rand_core 0.4.2", 3474 + "rdrand", 3475 + "winapi 0.3.9", 3476 + ] 3477 + 3478 + [[package]] 3479 + name = "rand_pcg" 3480 + version = "0.1.2" 3481 + source = "registry+https://github.com/rust-lang/crates.io-index" 3482 + checksum = "abf9b09b01790cfe0364f52bf32995ea3c39f4d2dd011eac241d2914146d0b44" 3483 + dependencies = [ 3484 + "autocfg 0.1.8", 3485 + "rand_core 0.4.2", 3486 + ] 3487 + 3488 + [[package]] 3489 + name = "rand_xorshift" 3490 + version = "0.1.1" 3491 + source = "registry+https://github.com/rust-lang/crates.io-index" 3492 + checksum = "cbf7e9e623549b0e21f6e97cf8ecf247c1a8fd2e8a992ae265314300b2455d5c" 3493 + dependencies = [ 3494 + "rand_core 0.3.1", 3495 + ] 3496 + 3497 + [[package]] 3498 + name = "rayon" 3499 + version = "1.7.0" 3500 + source = "registry+https://github.com/rust-lang/crates.io-index" 3501 + checksum = "1d2df5196e37bcc87abebc0053e20787d73847bb33134a69841207dd0a47f03b" 3502 + dependencies = [ 3503 + "either", 3504 + "rayon-core", 3505 + ] 3506 + 3507 + [[package]] 3508 + name = "rayon-core" 3509 + version = "1.11.0" 3510 + source = "registry+https://github.com/rust-lang/crates.io-index" 3511 + checksum = "4b8f95bd6966f5c87776639160a66bd8ab9895d9d4ab01ddba9fc60661aebe8d" 3512 + dependencies = [ 3513 + "crossbeam-channel", 3514 + "crossbeam-deque", 3515 + "crossbeam-utils", 3516 + "num_cpus", 3517 + ] 3518 + 3519 + [[package]] 3520 + name = "rdev" 3521 + version = "0.5.0-2" 3522 + source = "git+https://github.com/open-trade/rdev#0ad53987fa6f0e37a7bc000358f71c3802de4e7c" 3523 + dependencies = [ 3524 + "cocoa", 3525 + "core-foundation", 3526 + "core-foundation-sys", 3527 + "core-graphics", 3528 + "enum-map", 3529 + "lazy_static", 3530 + "libc", 3531 + "widestring 1.0.2", 3532 + "winapi 0.3.9", 3533 + "x11", 3534 + ] 3535 + 3536 + [[package]] 3537 + name = "rdrand" 3538 + version = "0.4.0" 3539 + source = "registry+https://github.com/rust-lang/crates.io-index" 3540 + checksum = "678054eb77286b51581ba43620cc911abf02758c91f93f479767aed0f90458b2" 3541 + dependencies = [ 3542 + "rand_core 0.3.1", 3543 + ] 3544 + 3545 + [[package]] 3546 + name = "realfft" 3547 + version = "3.2.0" 3548 + source = "registry+https://github.com/rust-lang/crates.io-index" 3549 + checksum = "93d6b8e8f0c6d2234aa58048d7290c60bf92cd36fd2888cd8331c66ad4f2e1d2" 3550 + dependencies = [ 3551 + "rustfft", 3552 + ] 3553 + 3554 + [[package]] 3555 + name = "redox_syscall" 3556 + version = "0.2.16" 3557 + source = "registry+https://github.com/rust-lang/crates.io-index" 3558 + checksum = "fb5a58c1855b4b6819d59012155603f0b22ad30cad752600aadfcb695265519a" 3559 + dependencies = [ 3560 + "bitflags", 3561 + ] 3562 + 3563 + [[package]] 3564 + name = "redox_syscall" 3565 + version = "0.3.5" 3566 + source = "registry+https://github.com/rust-lang/crates.io-index" 3567 + checksum = "567664f262709473930a4bf9e51bf2ebf3348f2e748ccc50dea20646858f8f29" 3568 + dependencies = [ 3569 + "bitflags", 3570 + ] 3571 + 3572 + [[package]] 3573 + name = "redox_users" 3574 + version = "0.4.3" 3575 + source = "registry+https://github.com/rust-lang/crates.io-index" 3576 + checksum = "b033d837a7cf162d7993aded9304e30a83213c648b6e389db233191f891e5c2b" 3577 + dependencies = [ 3578 + "getrandom", 3579 + "redox_syscall 0.2.16", 3580 + "thiserror", 3581 + ] 3582 + 3583 + [[package]] 3584 + name = "regex" 3585 + version = "1.7.3" 3586 + source = "registry+https://github.com/rust-lang/crates.io-index" 3587 + checksum = "8b1f693b24f6ac912f4893ef08244d70b6067480d2f1a46e950c9691e6749d1d" 3588 + dependencies = [ 3589 + "aho-corasick", 3590 + "memchr", 3591 + "regex-syntax", 3592 + ] 3593 + 3594 + [[package]] 3595 + name = "regex-syntax" 3596 + version = "0.6.29" 3597 + source = "registry+https://github.com/rust-lang/crates.io-index" 3598 + checksum = "f162c6dd7b008981e4d40210aca20b4bd0f9b60ca9271061b07f78537722f2e1" 3599 + 3600 + [[package]] 3601 + name = "repng" 3602 + version = "0.2.2" 3603 + source = "registry+https://github.com/rust-lang/crates.io-index" 3604 + checksum = "0dd57cd2cb5cc699b3eb4824d654e5a32f3bc013766da4966f71fe94805abbda" 3605 + dependencies = [ 3606 + "byteorder", 3607 + "flate2", 3608 + ] 3609 + 3610 + [[package]] 3611 + name = "ring" 3612 + version = "0.16.20" 3613 + source = "registry+https://github.com/rust-lang/crates.io-index" 3614 + checksum = "3053cf52e236a3ed746dfc745aa9cacf1b791d846bdaf412f60a8d7d6e17c8fc" 3615 + dependencies = [ 3616 + "cc", 3617 + "libc", 3618 + "once_cell", 3619 + "spin", 3620 + "untrusted", 3621 + "web-sys", 3622 + "winapi 0.3.9", 3623 + ] 3624 + 3625 + [[package]] 3626 + name = "rpassword" 3627 + version = "2.1.0" 3628 + source = "registry+https://github.com/rust-lang/crates.io-index" 3629 + checksum = "d37473170aedbe66ffa3ad3726939ba677d83c646ad4fd99e5b4bc38712f45ec" 3630 + dependencies = [ 3631 + "kernel32-sys", 3632 + "libc", 3633 + "winapi 0.2.8", 3634 + ] 3635 + 3636 + [[package]] 3637 + name = "rpassword" 3638 + version = "6.0.1" 3639 + source = "registry+https://github.com/rust-lang/crates.io-index" 3640 + checksum = "2bf099a1888612545b683d2661a1940089f6c2e5a8e38979b2159da876bfd956" 3641 + dependencies = [ 3642 + "libc", 3643 + "serde 1.0.159", 3644 + "serde_json 1.0.95", 3645 + "winapi 0.3.9", 3646 + ] 3647 + 3648 + [[package]] 3649 + name = "rubato" 3650 + version = "0.12.0" 3651 + source = "registry+https://github.com/rust-lang/crates.io-index" 3652 + checksum = "cd70209c27d5b08f5528bdc779ea3ffb418954e28987f9f9775c6eac41003f9c" 3653 + dependencies = [ 3654 + "num-complex", 3655 + "num-integer", 3656 + "num-traits 0.2.15", 3657 + "realfft", 3658 + ] 3659 + 3660 + [[package]] 3661 + name = "runas" 3662 + version = "0.2.1" 3663 + source = "registry+https://github.com/rust-lang/crates.io-index" 3664 + checksum = "a620b0994a180cdfa25c0439e6d58c0628272571501880d626ffff58e96a0799" 3665 + dependencies = [ 3666 + "cc", 3667 + "which 3.1.1", 3668 + ] 3669 + 3670 + [[package]] 3671 + name = "rust-pulsectl" 3672 + version = "0.2.12" 3673 + source = "git+https://github.com/open-trade/pulsectl#5e68f4c2b7c644fa321984688602d71e8ad0bba3" 3674 + dependencies = [ 3675 + "libpulse-binding", 3676 + ] 3677 + 3678 + [[package]] 3679 + name = "rustc-demangle" 3680 + version = "0.1.22" 3681 + source = "registry+https://github.com/rust-lang/crates.io-index" 3682 + checksum = "d4a36c42d1873f9a77c53bde094f9664d9891bc604a45b4798fd2c389ed12e5b" 3683 + 3684 + [[package]] 3685 + name = "rustc-hash" 3686 + version = "1.1.0" 3687 + source = "registry+https://github.com/rust-lang/crates.io-index" 3688 + checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2" 3689 + 3690 + [[package]] 3691 + name = "rustc_version" 3692 + version = "0.4.0" 3693 + source = "registry+https://github.com/rust-lang/crates.io-index" 3694 + checksum = "bfa0f585226d2e68097d4f95d113b15b83a82e819ab25717ec0590d9584ef366" 3695 + dependencies = [ 3696 + "semver", 3697 + ] 3698 + 3699 + [[package]] 3700 + name = "rustdesk" 3701 + version = "1.1.9" 3702 + dependencies = [ 3703 + "android_logger", 3704 + "arboard", 3705 + "async-trait", 3706 + "base64 0.13.1", 3707 + "cc", 3708 + "cfg-if 1.0.0", 3709 + "clap 3.2.23", 3710 + "clipboard", 3711 + "clipboard-master", 3712 + "cocoa", 3713 + "core-foundation", 3714 + "core-graphics", 3715 + "cpal", 3716 + "ctrlc", 3717 + "dasp", 3718 + "dispatch", 3719 + "enigo", 3720 + "flexi_logger", 3721 + "hbb_common", 3722 + "hound", 3723 + "include_dir", 3724 + "lazy_static", 3725 + "libc", 3726 + "libpulse-binding", 3727 + "libpulse-simple-binding", 3728 + "mac_address", 3729 + "machine-uid", 3730 + "magnum-opus", 3731 + "num_cpus", 3732 + "objc", 3733 + "parity-tokio-ipc", 3734 + "rdev", 3735 + "repng", 3736 + "rpassword 6.0.1", 3737 + "rubato", 3738 + "runas", 3739 + "rust-pulsectl", 3740 + "samplerate", 3741 + "sciter-rs", 3742 + "scrap", 3743 + "serde 1.0.159", 3744 + "serde_derive", 3745 + "serde_json 1.0.95", 3746 + "sha2", 3747 + "sys-locale", 3748 + "sysinfo", 3749 + "systray", 3750 + "tray-item", 3751 + "uuid", 3752 + "virtual_display", 3753 + "whoami", 3754 + "winapi 0.3.9", 3755 + "windows-service", 3756 + "winreg 0.10.1", 3757 + "winres", 3758 + ] 3759 + 3760 + [[package]] 3761 + name = "rustfft" 3762 + version = "6.1.0" 3763 + source = "registry+https://github.com/rust-lang/crates.io-index" 3764 + checksum = "e17d4f6cbdb180c9f4b2a26bbf01c4e647f1e1dea22fe8eb9db54198b32f9434" 3765 + dependencies = [ 3766 + "num-complex", 3767 + "num-integer", 3768 + "num-traits 0.2.15", 3769 + "primal-check", 3770 + "strength_reduce", 3771 + "transpose", 3772 + "version_check", 3773 + ] 3774 + 3775 + [[package]] 3776 + name = "rustix" 3777 + version = "0.37.11" 3778 + source = "registry+https://github.com/rust-lang/crates.io-index" 3779 + checksum = "85597d61f83914ddeba6a47b3b8ffe7365107221c2e557ed94426489fefb5f77" 3780 + dependencies = [ 3781 + "bitflags", 3782 + "errno", 3783 + "io-lifetimes", 3784 + "libc", 3785 + "linux-raw-sys", 3786 + "windows-sys 0.48.0", 3787 + ] 3788 + 3789 + [[package]] 3790 + name = "rustls" 3791 + version = "0.20.8" 3792 + source = "registry+https://github.com/rust-lang/crates.io-index" 3793 + checksum = "fff78fc74d175294f4e83b28343315ffcfb114b156f0185e9741cb5570f50e2f" 3794 + dependencies = [ 3795 + "ring", 3796 + "sct", 3797 + "webpki", 3798 + ] 3799 + 3800 + [[package]] 3801 + name = "rustls-native-certs" 3802 + version = "0.6.2" 3803 + source = "registry+https://github.com/rust-lang/crates.io-index" 3804 + checksum = "0167bac7a9f490495f3c33013e7722b53cb087ecbe082fb0c6387c96f634ea50" 3805 + dependencies = [ 3806 + "openssl-probe", 3807 + "rustls-pemfile 1.0.2", 3808 + "schannel", 3809 + "security-framework", 3810 + ] 3811 + 3812 + [[package]] 3813 + name = "rustls-pemfile" 3814 + version = "0.2.1" 3815 + source = "registry+https://github.com/rust-lang/crates.io-index" 3816 + checksum = "5eebeaeb360c87bfb72e84abdb3447159c0eaececf1bef2aecd65a8be949d1c9" 3817 + dependencies = [ 3818 + "base64 0.13.1", 3819 + ] 3820 + 3821 + [[package]] 3822 + name = "rustls-pemfile" 3823 + version = "1.0.2" 3824 + source = "registry+https://github.com/rust-lang/crates.io-index" 3825 + checksum = "d194b56d58803a43635bdc398cd17e383d6f71f9182b9a192c127ca42494a59b" 3826 + dependencies = [ 3827 + "base64 0.21.0", 3828 + ] 3829 + 3830 + [[package]] 3831 + name = "rustversion" 3832 + version = "1.0.12" 3833 + source = "registry+https://github.com/rust-lang/crates.io-index" 3834 + checksum = "4f3208ce4d8448b3f3e7d168a73f5e0c43a61e32930de3bceeccedb388b6bf06" 3835 + 3836 + [[package]] 3837 + name = "ryu" 3838 + version = "1.0.13" 3839 + source = "registry+https://github.com/rust-lang/crates.io-index" 3840 + checksum = "f91339c0467de62360649f8d3e185ca8de4224ff281f66000de5eb2a77a79041" 3841 + 3842 + [[package]] 3843 + name = "same-file" 3844 + version = "1.0.6" 3845 + source = "registry+https://github.com/rust-lang/crates.io-index" 3846 + checksum = "93fc1dc3aaa9bfed95e02e6eadabb4baf7e3078b0bd1b4d7b6b0b68378900502" 3847 + dependencies = [ 3848 + "winapi-util", 3849 + ] 3850 + 3851 + [[package]] 3852 + name = "samplerate" 3853 + version = "0.2.4" 3854 + source = "registry+https://github.com/rust-lang/crates.io-index" 3855 + checksum = "e032b2b24715c4f982f483ea3abdb3c9ba444d9f63e87b2843d6f998f5ba2698" 3856 + dependencies = [ 3857 + "libsamplerate-sys", 3858 + ] 3859 + 3860 + [[package]] 3861 + name = "schannel" 3862 + version = "0.1.21" 3863 + source = "registry+https://github.com/rust-lang/crates.io-index" 3864 + checksum = "713cfb06c7059f3588fb8044c0fad1d09e3c01d225e25b9220dbfdcf16dbb1b3" 3865 + dependencies = [ 3866 + "windows-sys 0.42.0", 3867 + ] 3868 + 3869 + [[package]] 3870 + name = "sciter-rs" 3871 + version = "0.5.57" 3872 + source = "git+https://github.com/open-trade/rust-sciter?branch=dyn#82025b9ba77d5ae14543009444033036dbe25917" 3873 + dependencies = [ 3874 + "lazy_static", 3875 + "libc", 3876 + "objc", 3877 + "objc-foundation", 3878 + ] 3879 + 3880 + [[package]] 3881 + name = "scopeguard" 3882 + version = "1.1.0" 3883 + source = "registry+https://github.com/rust-lang/crates.io-index" 3884 + checksum = "d29ab0c6d3fc0ee92fe66e2d99f700eab17a8d57d1c1d3b748380fb20baa78cd" 3885 + 3886 + [[package]] 3887 + name = "scrap" 3888 + version = "0.5.0" 3889 + dependencies = [ 3890 + "bindgen 0.59.2", 3891 + "block", 3892 + "cfg-if 1.0.0", 3893 + "dbus", 3894 + "docopt", 3895 + "gstreamer", 3896 + "gstreamer-app", 3897 + "gstreamer-video", 3898 + "libc", 3899 + "num_cpus", 3900 + "quest", 3901 + "repng", 3902 + "serde 1.0.159", 3903 + "target_build_utils", 3904 + "tracing", 3905 + "vcpkg", 3906 + "webm", 3907 + "winapi 0.3.9", 3908 + ] 3909 + 3910 + [[package]] 3911 + name = "scratch" 3912 + version = "1.0.5" 3913 + source = "registry+https://github.com/rust-lang/crates.io-index" 3914 + checksum = "1792db035ce95be60c3f8853017b3999209281c24e2ba5bc8e59bf97a0c590c1" 3915 + 3916 + [[package]] 3917 + name = "sct" 3918 + version = "0.7.0" 3919 + source = "registry+https://github.com/rust-lang/crates.io-index" 3920 + checksum = "d53dcdb7c9f8158937a7981b48accfd39a43af418591a5d008c7b22b5e1b7ca4" 3921 + dependencies = [ 3922 + "ring", 3923 + "untrusted", 3924 + ] 3925 + 3926 + [[package]] 3927 + name = "security-framework" 3928 + version = "2.8.2" 3929 + source = "registry+https://github.com/rust-lang/crates.io-index" 3930 + checksum = "a332be01508d814fed64bf28f798a146d73792121129962fdf335bb3c49a4254" 3931 + dependencies = [ 3932 + "bitflags", 3933 + "core-foundation", 3934 + "core-foundation-sys", 3935 + "libc", 3936 + "security-framework-sys", 3937 + ] 3938 + 3939 + [[package]] 3940 + name = "security-framework-sys" 3941 + version = "2.8.0" 3942 + source = "registry+https://github.com/rust-lang/crates.io-index" 3943 + checksum = "31c9bb296072e961fcbd8853511dd39c2d8be2deb1e17c6860b1d30732b323b4" 3944 + dependencies = [ 3945 + "core-foundation-sys", 3946 + "libc", 3947 + ] 3948 + 3949 + [[package]] 3950 + name = "semver" 3951 + version = "1.0.17" 3952 + source = "registry+https://github.com/rust-lang/crates.io-index" 3953 + checksum = "bebd363326d05ec3e2f532ab7660680f3b02130d780c299bca73469d521bc0ed" 3954 + 3955 + [[package]] 3956 + name = "serde" 3957 + version = "0.9.15" 3958 + source = "registry+https://github.com/rust-lang/crates.io-index" 3959 + checksum = "34b623917345a631dc9608d5194cc206b3fe6c3554cd1c75b937e55e285254af" 3960 + 3961 + [[package]] 3962 + name = "serde" 3963 + version = "1.0.159" 3964 + source = "registry+https://github.com/rust-lang/crates.io-index" 3965 + checksum = "3c04e8343c3daeec41f58990b9d77068df31209f2af111e059e9fe9646693065" 3966 + dependencies = [ 3967 + "serde_derive", 3968 + ] 3969 + 3970 + [[package]] 3971 + name = "serde_derive" 3972 + version = "1.0.159" 3973 + source = "registry+https://github.com/rust-lang/crates.io-index" 3974 + checksum = "4c614d17805b093df4b147b51339e7e44bf05ef59fba1e45d83500bcfb4d8585" 3975 + dependencies = [ 3976 + "proc-macro2", 3977 + "quote", 3978 + "syn 2.0.13", 3979 + ] 3980 + 3981 + [[package]] 3982 + name = "serde_json" 3983 + version = "0.9.10" 3984 + source = "registry+https://github.com/rust-lang/crates.io-index" 3985 + checksum = "ad8bcf487be7d2e15d3d543f04312de991d631cfe1b43ea0ade69e6a8a5b16a1" 3986 + dependencies = [ 3987 + "dtoa", 3988 + "itoa 0.3.4", 3989 + "num-traits 0.1.43", 3990 + "serde 0.9.15", 3991 + ] 3992 + 3993 + [[package]] 3994 + name = "serde_json" 3995 + version = "1.0.95" 3996 + source = "registry+https://github.com/rust-lang/crates.io-index" 3997 + checksum = "d721eca97ac802aa7777b701877c8004d950fc142651367300d21c1cc0194744" 3998 + dependencies = [ 3999 + "itoa 1.0.6", 4000 + "ryu", 4001 + "serde 1.0.159", 4002 + ] 4003 + 4004 + [[package]] 4005 + name = "serde_spanned" 4006 + version = "0.6.1" 4007 + source = "registry+https://github.com/rust-lang/crates.io-index" 4008 + checksum = "0efd8caf556a6cebd3b285caf480045fcc1ac04f6bd786b09a6f11af30c4fcf4" 4009 + dependencies = [ 4010 + "serde 1.0.159", 4011 + ] 4012 + 4013 + [[package]] 4014 + name = "sha2" 4015 + version = "0.10.6" 4016 + source = "registry+https://github.com/rust-lang/crates.io-index" 4017 + checksum = "82e6b795fe2e3b1e845bafcb27aa35405c4d47cdfc92af5fc8d3002f76cebdc0" 4018 + dependencies = [ 4019 + "cfg-if 1.0.0", 4020 + "cpufeatures", 4021 + "digest", 4022 + ] 4023 + 4024 + [[package]] 4025 + name = "shlex" 4026 + version = "1.1.0" 4027 + source = "registry+https://github.com/rust-lang/crates.io-index" 4028 + checksum = "43b2853a4d09f215c24cc5489c992ce46052d359b5109343cbafbf26bc62f8a3" 4029 + 4030 + [[package]] 4031 + name = "signal-hook-registry" 4032 + version = "1.4.1" 4033 + source = "registry+https://github.com/rust-lang/crates.io-index" 4034 + checksum = "d8229b473baa5980ac72ef434c4415e70c4b5e71b423043adb4ba059f89c99a1" 4035 + dependencies = [ 4036 + "libc", 4037 + ] 4038 + 4039 + [[package]] 4040 + name = "signature" 4041 + version = "1.6.4" 4042 + source = "registry+https://github.com/rust-lang/crates.io-index" 4043 + checksum = "74233d3b3b2f6d4b006dc19dee745e73e2a6bfb6f93607cd3b02bd5b00797d7c" 4044 + 4045 + [[package]] 4046 + name = "siphasher" 4047 + version = "0.2.3" 4048 + source = "registry+https://github.com/rust-lang/crates.io-index" 4049 + checksum = "0b8de496cf83d4ed58b6be86c3a275b8602f6ffe98d3024a869e124147a9a3ac" 4050 + 4051 + [[package]] 4052 + name = "slab" 4053 + version = "0.4.8" 4054 + source = "registry+https://github.com/rust-lang/crates.io-index" 4055 + checksum = "6528351c9bc8ab22353f9d776db39a20288e8d6c37ef8cfe3317cf875eecfc2d" 4056 + dependencies = [ 4057 + "autocfg 1.1.0", 4058 + ] 4059 + 4060 + [[package]] 4061 + name = "smallvec" 4062 + version = "1.10.0" 4063 + source = "registry+https://github.com/rust-lang/crates.io-index" 4064 + checksum = "a507befe795404456341dfab10cef66ead4c041f62b8b11bbb92bffe5d0953e0" 4065 + 4066 + [[package]] 4067 + name = "socket2" 4068 + version = "0.3.19" 4069 + source = "registry+https://github.com/rust-lang/crates.io-index" 4070 + checksum = "122e570113d28d773067fab24266b66753f6ea915758651696b6e35e49f88d6e" 4071 + dependencies = [ 4072 + "cfg-if 1.0.0", 4073 + "libc", 4074 + "winapi 0.3.9", 4075 + ] 4076 + 4077 + [[package]] 4078 + name = "socket2" 4079 + version = "0.4.9" 4080 + source = "registry+https://github.com/rust-lang/crates.io-index" 4081 + checksum = "64a4a911eed85daf18834cfaa86a79b7d266ff93ff5ba14005426219480ed662" 4082 + dependencies = [ 4083 + "libc", 4084 + "winapi 0.3.9", 4085 + ] 4086 + 4087 + [[package]] 4088 + name = "sodiumoxide" 4089 + version = "0.2.7" 4090 + source = "registry+https://github.com/rust-lang/crates.io-index" 4091 + checksum = "e26be3acb6c2d9a7aac28482586a7856436af4cfe7100031d219de2d2ecb0028" 4092 + dependencies = [ 4093 + "ed25519", 4094 + "libc", 4095 + "libsodium-sys", 4096 + "serde 1.0.159", 4097 + ] 4098 + 4099 + [[package]] 4100 + name = "spin" 4101 + version = "0.5.2" 4102 + source = "registry+https://github.com/rust-lang/crates.io-index" 4103 + checksum = "6e63cff320ae2c57904679ba7cb63280a3dc4613885beafb148ee7bf9aa9042d" 4104 + 4105 + [[package]] 4106 + name = "static_assertions" 4107 + version = "1.1.0" 4108 + source = "registry+https://github.com/rust-lang/crates.io-index" 4109 + checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f" 4110 + 4111 + [[package]] 4112 + name = "stdweb" 4113 + version = "0.1.3" 4114 + source = "registry+https://github.com/rust-lang/crates.io-index" 4115 + checksum = "ef5430c8e36b713e13b48a9f709cc21e046723fe44ce34587b73a830203b533e" 4116 + 4117 + [[package]] 4118 + name = "str-buf" 4119 + version = "1.0.6" 4120 + source = "registry+https://github.com/rust-lang/crates.io-index" 4121 + checksum = "9e08d8363704e6c71fc928674353e6b7c23dcea9d82d7012c8faf2a3a025f8d0" 4122 + 4123 + [[package]] 4124 + name = "strength_reduce" 4125 + version = "0.2.4" 4126 + source = "registry+https://github.com/rust-lang/crates.io-index" 4127 + checksum = "fe895eb47f22e2ddd4dabc02bce419d2e643c8e3b585c78158b349195bc24d82" 4128 + 4129 + [[package]] 4130 + name = "strsim" 4131 + version = "0.8.0" 4132 + source = "registry+https://github.com/rust-lang/crates.io-index" 4133 + checksum = "8ea5119cdb4c55b55d432abb513a0429384878c15dde60cc77b1c99de1a95a6a" 4134 + 4135 + [[package]] 4136 + name = "strsim" 4137 + version = "0.10.0" 4138 + source = "registry+https://github.com/rust-lang/crates.io-index" 4139 + checksum = "73473c0e59e6d5812c5dfe2a064a6444949f089e20eec9a2e5506596494e4623" 4140 + 4141 + [[package]] 4142 + name = "strum" 4143 + version = "0.18.0" 4144 + source = "registry+https://github.com/rust-lang/crates.io-index" 4145 + checksum = "57bd81eb48f4c437cadc685403cad539345bf703d78e63707418431cecd4522b" 4146 + 4147 + [[package]] 4148 + name = "strum_macros" 4149 + version = "0.18.0" 4150 + source = "registry+https://github.com/rust-lang/crates.io-index" 4151 + checksum = "87c85aa3f8ea653bfd3ddf25f7ee357ee4d204731f6aa9ad04002306f6e2774c" 4152 + dependencies = [ 4153 + "heck 0.3.3", 4154 + "proc-macro2", 4155 + "quote", 4156 + "syn 1.0.109", 4157 + ] 4158 + 4159 + [[package]] 4160 + name = "syn" 4161 + version = "1.0.109" 4162 + source = "registry+https://github.com/rust-lang/crates.io-index" 4163 + checksum = "72b64191b275b66ffe2469e8af2c1cfe3bafa67b529ead792a6d0160888b4237" 4164 + dependencies = [ 4165 + "proc-macro2", 4166 + "quote", 4167 + "unicode-ident", 4168 + ] 4169 + 4170 + [[package]] 4171 + name = "syn" 4172 + version = "2.0.13" 4173 + source = "registry+https://github.com/rust-lang/crates.io-index" 4174 + checksum = "4c9da457c5285ac1f936ebd076af6dac17a61cfe7826f2076b4d015cf47bc8ec" 4175 + dependencies = [ 4176 + "proc-macro2", 4177 + "quote", 4178 + "unicode-ident", 4179 + ] 4180 + 4181 + [[package]] 4182 + name = "synstructure" 4183 + version = "0.12.6" 4184 + source = "registry+https://github.com/rust-lang/crates.io-index" 4185 + checksum = "f36bdaa60a83aca3921b5259d5400cbf5e90fc51931376a9bd4a0eb79aa7210f" 4186 + dependencies = [ 4187 + "proc-macro2", 4188 + "quote", 4189 + "syn 1.0.109", 4190 + "unicode-xid", 4191 + ] 4192 + 4193 + [[package]] 4194 + name = "sys-locale" 4195 + version = "0.2.4" 4196 + source = "registry+https://github.com/rust-lang/crates.io-index" 4197 + checksum = "f8a11bd9c338fdba09f7881ab41551932ad42e405f61d01e8406baea71c07aee" 4198 + dependencies = [ 4199 + "js-sys", 4200 + "libc", 4201 + "wasm-bindgen", 4202 + "web-sys", 4203 + "windows-sys 0.45.0", 4204 + ] 4205 + 4206 + [[package]] 4207 + name = "sysinfo" 4208 + version = "0.23.13" 4209 + source = "registry+https://github.com/rust-lang/crates.io-index" 4210 + checksum = "3977ec2e0520829be45c8a2df70db2bf364714d8a748316a10c3c35d4d2b01c9" 4211 + dependencies = [ 4212 + "cfg-if 1.0.0", 4213 + "core-foundation-sys", 4214 + "libc", 4215 + "ntapi", 4216 + "once_cell", 4217 + "rayon", 4218 + "winapi 0.3.9", 4219 + ] 4220 + 4221 + [[package]] 4222 + name = "system-deps" 4223 + version = "1.3.2" 4224 + source = "registry+https://github.com/rust-lang/crates.io-index" 4225 + checksum = "0f3ecc17269a19353b3558b313bba738b25d82993e30d62a18406a24aba4649b" 4226 + dependencies = [ 4227 + "heck 0.3.3", 4228 + "pkg-config", 4229 + "strum", 4230 + "strum_macros", 4231 + "thiserror", 4232 + "toml 0.5.11", 4233 + "version-compare 0.0.10", 4234 + ] 4235 + 4236 + [[package]] 4237 + name = "system-deps" 4238 + version = "6.0.4" 4239 + source = "registry+https://github.com/rust-lang/crates.io-index" 4240 + checksum = "555fc8147af6256f3931a36bb83ad0023240ce9cf2b319dec8236fd1f220b05f" 4241 + dependencies = [ 4242 + "cfg-expr", 4243 + "heck 0.4.1", 4244 + "pkg-config", 4245 + "toml 0.7.3", 4246 + "version-compare 0.1.1", 4247 + ] 4248 + 4249 + [[package]] 4250 + name = "systray" 4251 + version = "0.4.1" 4252 + source = "git+https://github.com/liyue201/systray-rs#84cca4b4171661bc6c4d1ba5aaa2320ff8e085aa" 4253 + dependencies = [ 4254 + "glib 0.10.3", 4255 + "gtk 0.9.2", 4256 + "libappindicator 0.6.1", 4257 + "libc", 4258 + "log", 4259 + "winapi 0.3.9", 4260 + ] 4261 + 4262 + [[package]] 4263 + name = "target_build_utils" 4264 + version = "0.3.1" 4265 + source = "registry+https://github.com/rust-lang/crates.io-index" 4266 + checksum = "013d134ae4a25ee744ad6129db589018558f620ddfa44043887cdd45fa08e75c" 4267 + dependencies = [ 4268 + "phf", 4269 + "phf_codegen", 4270 + "serde_json 0.9.10", 4271 + ] 4272 + 4273 + [[package]] 4274 + name = "tempfile" 4275 + version = "3.5.0" 4276 + source = "registry+https://github.com/rust-lang/crates.io-index" 4277 + checksum = "b9fbec84f381d5795b08656e4912bec604d162bff9291d6189a78f4c8ab87998" 4278 + dependencies = [ 4279 + "cfg-if 1.0.0", 4280 + "fastrand", 4281 + "redox_syscall 0.3.5", 4282 + "rustix", 4283 + "windows-sys 0.45.0", 4284 + ] 4285 + 4286 + [[package]] 4287 + name = "termcolor" 4288 + version = "1.2.0" 4289 + source = "registry+https://github.com/rust-lang/crates.io-index" 4290 + checksum = "be55cf8942feac5c765c2c993422806843c9a9a45d4d5c407ad6dd2ea95eb9b6" 4291 + dependencies = [ 4292 + "winapi-util", 4293 + ] 4294 + 4295 + [[package]] 4296 + name = "termios" 4297 + version = "0.3.3" 4298 + source = "registry+https://github.com/rust-lang/crates.io-index" 4299 + checksum = "411c5bf740737c7918b8b1fe232dca4dc9f8e754b8ad5e20966814001ed0ac6b" 4300 + dependencies = [ 4301 + "libc", 4302 + ] 4303 + 4304 + [[package]] 4305 + name = "textwrap" 4306 + version = "0.11.0" 4307 + source = "registry+https://github.com/rust-lang/crates.io-index" 4308 + checksum = "d326610f408c7a4eb6f51c37c330e496b08506c9457c9d34287ecc38809fb060" 4309 + dependencies = [ 4310 + "unicode-width", 4311 + ] 4312 + 4313 + [[package]] 4314 + name = "textwrap" 4315 + version = "0.16.0" 4316 + source = "registry+https://github.com/rust-lang/crates.io-index" 4317 + checksum = "222a222a5bfe1bba4a77b45ec488a741b3cb8872e5e499451fd7d0129c9c7c3d" 4318 + 4319 + [[package]] 4320 + name = "thiserror" 4321 + version = "1.0.40" 4322 + source = "registry+https://github.com/rust-lang/crates.io-index" 4323 + checksum = "978c9a314bd8dc99be594bc3c175faaa9794be04a5a5e153caba6915336cebac" 4324 + dependencies = [ 4325 + "thiserror-impl", 4326 + ] 4327 + 4328 + [[package]] 4329 + name = "thiserror-impl" 4330 + version = "1.0.40" 4331 + source = "registry+https://github.com/rust-lang/crates.io-index" 4332 + checksum = "f9456a42c5b0d803c8cd86e73dd7cc9edd429499f37a3550d286d5e86720569f" 4333 + dependencies = [ 4334 + "proc-macro2", 4335 + "quote", 4336 + "syn 2.0.13", 4337 + ] 4338 + 4339 + [[package]] 4340 + name = "tiff" 4341 + version = "0.6.1" 4342 + source = "registry+https://github.com/rust-lang/crates.io-index" 4343 + checksum = "9a53f4706d65497df0c4349241deddf35f84cee19c87ed86ea8ca590f4464437" 4344 + dependencies = [ 4345 + "jpeg-decoder", 4346 + "miniz_oxide 0.4.4", 4347 + "weezl", 4348 + ] 4349 + 4350 + [[package]] 4351 + name = "time" 4352 + version = "0.3.9" 4353 + source = "registry+https://github.com/rust-lang/crates.io-index" 4354 + checksum = "c2702e08a7a860f005826c6815dcac101b19b5eb330c27fe4a5928fec1d20ddd" 4355 + dependencies = [ 4356 + "itoa 1.0.6", 4357 + "libc", 4358 + "num_threads", 4359 + "time-macros", 4360 + ] 4361 + 4362 + [[package]] 4363 + name = "time-macros" 4364 + version = "0.2.4" 4365 + source = "registry+https://github.com/rust-lang/crates.io-index" 4366 + checksum = "42657b1a6f4d817cda8e7a0ace261fe0cc946cf3a80314390b22cc61ae080792" 4367 + 4368 + [[package]] 4369 + name = "tinyvec" 4370 + version = "1.6.0" 4371 + source = "registry+https://github.com/rust-lang/crates.io-index" 4372 + checksum = "87cc5ceb3875bb20c2890005a4e226a4651264a5c75edb2421b52861a0a0cb50" 4373 + dependencies = [ 4374 + "tinyvec_macros", 4375 + ] 4376 + 4377 + [[package]] 4378 + name = "tinyvec_macros" 4379 + version = "0.1.1" 4380 + source = "registry+https://github.com/rust-lang/crates.io-index" 4381 + checksum = "1f3ccbac311fea05f86f61904b462b55fb3df8837a366dfc601a0161d0532f20" 4382 + 4383 + [[package]] 4384 + name = "tokio" 4385 + version = "1.27.0" 4386 + source = "registry+https://github.com/rust-lang/crates.io-index" 4387 + checksum = "d0de47a4eecbe11f498978a9b29d792f0d2692d1dd003650c24c76510e3bc001" 4388 + dependencies = [ 4389 + "autocfg 1.1.0", 4390 + "bytes", 4391 + "libc", 4392 + "mio 0.8.6", 4393 + "num_cpus", 4394 + "parking_lot 0.12.1", 4395 + "pin-project-lite", 4396 + "signal-hook-registry", 4397 + "socket2 0.4.9", 4398 + "tokio-macros", 4399 + "windows-sys 0.45.0", 4400 + ] 4401 + 4402 + [[package]] 4403 + name = "tokio-macros" 4404 + version = "2.0.0" 4405 + source = "registry+https://github.com/rust-lang/crates.io-index" 4406 + checksum = "61a573bdc87985e9d6ddeed1b3d864e8a302c847e40d647746df2f1de209d1ce" 4407 + dependencies = [ 4408 + "proc-macro2", 4409 + "quote", 4410 + "syn 2.0.13", 4411 + ] 4412 + 4413 + [[package]] 4414 + name = "tokio-socks" 4415 + version = "0.5.1-1" 4416 + source = "git+https://github.com/open-trade/tokio-socks#7034e79263ce25c348be072808d7601d82cd892d" 4417 + dependencies = [ 4418 + "bytes", 4419 + "either", 4420 + "futures-core", 4421 + "futures-sink", 4422 + "futures-util", 4423 + "pin-project", 4424 + "thiserror", 4425 + "tokio", 4426 + "tokio-util 0.7.7", 4427 + ] 4428 + 4429 + [[package]] 4430 + name = "tokio-util" 4431 + version = "0.6.10" 4432 + source = "registry+https://github.com/rust-lang/crates.io-index" 4433 + checksum = "36943ee01a6d67977dd3f84a5a1d2efeb4ada3a1ae771cadfaa535d9d9fc6507" 4434 + dependencies = [ 4435 + "bytes", 4436 + "futures-core", 4437 + "futures-io", 4438 + "futures-sink", 4439 + "log", 4440 + "pin-project-lite", 4441 + "slab", 4442 + "tokio", 4443 + ] 4444 + 4445 + [[package]] 4446 + name = "tokio-util" 4447 + version = "0.7.7" 4448 + source = "registry+https://github.com/rust-lang/crates.io-index" 4449 + checksum = "5427d89453009325de0d8f342c9490009f76e999cb7672d77e46267448f7e6b2" 4450 + dependencies = [ 4451 + "bytes", 4452 + "futures-core", 4453 + "futures-sink", 4454 + "pin-project-lite", 4455 + "tokio", 4456 + "tracing", 4457 + ] 4458 + 4459 + [[package]] 4460 + name = "toml" 4461 + version = "0.5.11" 4462 + source = "registry+https://github.com/rust-lang/crates.io-index" 4463 + checksum = "f4f7f0dd8d50a853a531c426359045b1998f04219d88799810762cd4ad314234" 4464 + dependencies = [ 4465 + "serde 1.0.159", 4466 + ] 4467 + 4468 + [[package]] 4469 + name = "toml" 4470 + version = "0.7.3" 4471 + source = "registry+https://github.com/rust-lang/crates.io-index" 4472 + checksum = "b403acf6f2bb0859c93c7f0d967cb4a75a7ac552100f9322faf64dc047669b21" 4473 + dependencies = [ 4474 + "serde 1.0.159", 4475 + "serde_spanned", 4476 + "toml_datetime", 4477 + "toml_edit", 4478 + ] 4479 + 4480 + [[package]] 4481 + name = "toml_datetime" 4482 + version = "0.6.1" 4483 + source = "registry+https://github.com/rust-lang/crates.io-index" 4484 + checksum = "3ab8ed2edee10b50132aed5f331333428b011c99402b5a534154ed15746f9622" 4485 + dependencies = [ 4486 + "serde 1.0.159", 4487 + ] 4488 + 4489 + [[package]] 4490 + name = "toml_edit" 4491 + version = "0.19.8" 4492 + source = "registry+https://github.com/rust-lang/crates.io-index" 4493 + checksum = "239410c8609e8125456927e6707163a3b1fdb40561e4b803bc041f466ccfdc13" 4494 + dependencies = [ 4495 + "indexmap", 4496 + "serde 1.0.159", 4497 + "serde_spanned", 4498 + "toml_datetime", 4499 + "winnow", 4500 + ] 4501 + 4502 + [[package]] 4503 + name = "tracing" 4504 + version = "0.1.37" 4505 + source = "registry+https://github.com/rust-lang/crates.io-index" 4506 + checksum = "8ce8c33a8d48bd45d624a6e523445fd21ec13d3653cd51f681abf67418f54eb8" 4507 + dependencies = [ 4508 + "cfg-if 1.0.0", 4509 + "pin-project-lite", 4510 + "tracing-attributes", 4511 + "tracing-core", 4512 + ] 4513 + 4514 + [[package]] 4515 + name = "tracing-attributes" 4516 + version = "0.1.23" 4517 + source = "registry+https://github.com/rust-lang/crates.io-index" 4518 + checksum = "4017f8f45139870ca7e672686113917c71c7a6e02d4924eda67186083c03081a" 4519 + dependencies = [ 4520 + "proc-macro2", 4521 + "quote", 4522 + "syn 1.0.109", 4523 + ] 4524 + 4525 + [[package]] 4526 + name = "tracing-core" 4527 + version = "0.1.30" 4528 + source = "registry+https://github.com/rust-lang/crates.io-index" 4529 + checksum = "24eb03ba0eab1fd845050058ce5e616558e8f8d8fca633e6b163fe25c797213a" 4530 + dependencies = [ 4531 + "once_cell", 4532 + ] 4533 + 4534 + [[package]] 4535 + name = "transpose" 4536 + version = "0.2.2" 4537 + source = "registry+https://github.com/rust-lang/crates.io-index" 4538 + checksum = "e6522d49d03727ffb138ae4cbc1283d3774f0d10aa7f9bf52e6784c45daf9b23" 4539 + dependencies = [ 4540 + "num-integer", 4541 + "strength_reduce", 4542 + ] 4543 + 4544 + [[package]] 4545 + name = "tray-item" 4546 + version = "0.7.1" 4547 + source = "registry+https://github.com/rust-lang/crates.io-index" 4548 + checksum = "0914b62e00e8f51241806cb9f9c4ea6b10c75d94cae02c89278de6f4b98c7d0f" 4549 + dependencies = [ 4550 + "cocoa", 4551 + "core-graphics", 4552 + "gtk 0.15.5", 4553 + "libappindicator 0.7.1", 4554 + "libc", 4555 + "objc", 4556 + "objc-foundation", 4557 + "objc_id", 4558 + "padlock", 4559 + "winapi 0.3.9", 4560 + ] 4561 + 4562 + [[package]] 4563 + name = "typenum" 4564 + version = "1.16.0" 4565 + source = "registry+https://github.com/rust-lang/crates.io-index" 4566 + checksum = "497961ef93d974e23eb6f433eb5fe1b7930b659f06d12dec6fc44a8f554c0bba" 4567 + 4568 + [[package]] 4569 + name = "unicode-ident" 4570 + version = "1.0.8" 4571 + source = "registry+https://github.com/rust-lang/crates.io-index" 4572 + checksum = "e5464a87b239f13a63a501f2701565754bae92d243d4bb7eb12f6d57d2269bf4" 4573 + 4574 + [[package]] 4575 + name = "unicode-segmentation" 4576 + version = "1.10.1" 4577 + source = "registry+https://github.com/rust-lang/crates.io-index" 4578 + checksum = "1dd624098567895118886609431a7c3b8f516e41d30e0643f03d94592a147e36" 4579 + 4580 + [[package]] 4581 + name = "unicode-width" 4582 + version = "0.1.10" 4583 + source = "registry+https://github.com/rust-lang/crates.io-index" 4584 + checksum = "c0edd1e5b14653f783770bce4a4dabb4a5108a5370a5f5d8cfe8710c361f6c8b" 4585 + 4586 + [[package]] 4587 + name = "unicode-xid" 4588 + version = "0.2.4" 4589 + source = "registry+https://github.com/rust-lang/crates.io-index" 4590 + checksum = "f962df74c8c05a667b5ee8bcf162993134c104e96440b663c8daa176dc772d8c" 4591 + 4592 + [[package]] 4593 + name = "untrusted" 4594 + version = "0.7.1" 4595 + source = "registry+https://github.com/rust-lang/crates.io-index" 4596 + checksum = "a156c684c91ea7d62626509bce3cb4e1d9ed5c4d978f7b4352658f96a4c26b4a" 4597 + 4598 + [[package]] 4599 + name = "uuid" 4600 + version = "1.3.0" 4601 + source = "registry+https://github.com/rust-lang/crates.io-index" 4602 + checksum = "1674845326ee10d37ca60470760d4288a6f80f304007d92e5c53bab78c9cfd79" 4603 + dependencies = [ 4604 + "getrandom", 4605 + ] 4606 + 4607 + [[package]] 4608 + name = "vcpkg" 4609 + version = "0.2.15" 4610 + source = "registry+https://github.com/rust-lang/crates.io-index" 4611 + checksum = "accd4ea62f7bb7a82fe23066fb0957d48ef677f6eeb8215f372f52e48bb32426" 4612 + 4613 + [[package]] 4614 + name = "vec_map" 4615 + version = "0.8.2" 4616 + source = "registry+https://github.com/rust-lang/crates.io-index" 4617 + checksum = "f1bddf1187be692e79c5ffeab891132dfb0f236ed36a43c7ed39f1165ee20191" 4618 + 4619 + [[package]] 4620 + name = "version-compare" 4621 + version = "0.0.10" 4622 + source = "registry+https://github.com/rust-lang/crates.io-index" 4623 + checksum = "d63556a25bae6ea31b52e640d7c41d1ab27faba4ccb600013837a3d0b3994ca1" 4624 + 4625 + [[package]] 4626 + name = "version-compare" 4627 + version = "0.1.1" 4628 + source = "registry+https://github.com/rust-lang/crates.io-index" 4629 + checksum = "579a42fc0b8e0c63b76519a339be31bed574929511fa53c1a3acae26eb258f29" 4630 + 4631 + [[package]] 4632 + name = "version_check" 4633 + version = "0.9.4" 4634 + source = "registry+https://github.com/rust-lang/crates.io-index" 4635 + checksum = "49874b5167b65d7193b8aba1567f5c7d93d001cafc34600cee003eda787e483f" 4636 + 4637 + [[package]] 4638 + name = "virtual_display" 4639 + version = "0.1.0" 4640 + dependencies = [ 4641 + "cc", 4642 + "hbb_common", 4643 + "lazy_static", 4644 + "serde 1.0.159", 4645 + "serde_derive", 4646 + "thiserror", 4647 + ] 4648 + 4649 + [[package]] 4650 + name = "walkdir" 4651 + version = "2.3.3" 4652 + source = "registry+https://github.com/rust-lang/crates.io-index" 4653 + checksum = "36df944cda56c7d8d8b7496af378e6b16de9284591917d307c9b4d313c44e698" 4654 + dependencies = [ 4655 + "same-file", 4656 + "winapi-util", 4657 + ] 4658 + 4659 + [[package]] 4660 + name = "wasi" 4661 + version = "0.11.0+wasi-snapshot-preview1" 4662 + source = "registry+https://github.com/rust-lang/crates.io-index" 4663 + checksum = "9c8d87e72b64a3b4db28d11ce29237c246188f4f51057d65a7eab63b7987e423" 4664 + 4665 + [[package]] 4666 + name = "wasm-bindgen" 4667 + version = "0.2.84" 4668 + source = "registry+https://github.com/rust-lang/crates.io-index" 4669 + checksum = "31f8dcbc21f30d9b8f2ea926ecb58f6b91192c17e9d33594b3df58b2007ca53b" 4670 + dependencies = [ 4671 + "cfg-if 1.0.0", 4672 + "wasm-bindgen-macro", 4673 + ] 4674 + 4675 + [[package]] 4676 + name = "wasm-bindgen-backend" 4677 + version = "0.2.84" 4678 + source = "registry+https://github.com/rust-lang/crates.io-index" 4679 + checksum = "95ce90fd5bcc06af55a641a86428ee4229e44e07033963a2290a8e241607ccb9" 4680 + dependencies = [ 4681 + "bumpalo", 4682 + "log", 4683 + "once_cell", 4684 + "proc-macro2", 4685 + "quote", 4686 + "syn 1.0.109", 4687 + "wasm-bindgen-shared", 4688 + ] 4689 + 4690 + [[package]] 4691 + name = "wasm-bindgen-macro" 4692 + version = "0.2.84" 4693 + source = "registry+https://github.com/rust-lang/crates.io-index" 4694 + checksum = "4c21f77c0bedc37fd5dc21f897894a5ca01e7bb159884559461862ae90c0b4c5" 4695 + dependencies = [ 4696 + "quote", 4697 + "wasm-bindgen-macro-support", 4698 + ] 4699 + 4700 + [[package]] 4701 + name = "wasm-bindgen-macro-support" 4702 + version = "0.2.84" 4703 + source = "registry+https://github.com/rust-lang/crates.io-index" 4704 + checksum = "2aff81306fcac3c7515ad4e177f521b5c9a15f2b08f4e32d823066102f35a5f6" 4705 + dependencies = [ 4706 + "proc-macro2", 4707 + "quote", 4708 + "syn 1.0.109", 4709 + "wasm-bindgen-backend", 4710 + "wasm-bindgen-shared", 4711 + ] 4712 + 4713 + [[package]] 4714 + name = "wasm-bindgen-shared" 4715 + version = "0.2.84" 4716 + source = "registry+https://github.com/rust-lang/crates.io-index" 4717 + checksum = "0046fef7e28c3804e5e38bfa31ea2a0f73905319b677e57ebe37e49358989b5d" 4718 + 4719 + [[package]] 4720 + name = "web-sys" 4721 + version = "0.3.61" 4722 + source = "registry+https://github.com/rust-lang/crates.io-index" 4723 + checksum = "e33b99f4b23ba3eec1a53ac264e35a755f00e966e0065077d6027c0f575b0b97" 4724 + dependencies = [ 4725 + "js-sys", 4726 + "wasm-bindgen", 4727 + ] 4728 + 4729 + [[package]] 4730 + name = "webm" 4731 + version = "1.0.2" 4732 + source = "registry+https://github.com/rust-lang/crates.io-index" 4733 + checksum = "ecb047148a12ef1fd8ab26302bca7e82036f005c3073b48e17cc1b44ec577136" 4734 + dependencies = [ 4735 + "webm-sys", 4736 + ] 4737 + 4738 + [[package]] 4739 + name = "webm-sys" 4740 + version = "1.0.3" 4741 + source = "registry+https://github.com/rust-lang/crates.io-index" 4742 + checksum = "0ded6ec82ccf51fe265b0b2b1579cac839574ed910c17baac58e807f8a9de7f3" 4743 + dependencies = [ 4744 + "cc", 4745 + ] 4746 + 4747 + [[package]] 4748 + name = "webpki" 4749 + version = "0.22.0" 4750 + source = "registry+https://github.com/rust-lang/crates.io-index" 4751 + checksum = "f095d78192e208183081cc07bc5515ef55216397af48b873e5edcd72637fa1bd" 4752 + dependencies = [ 4753 + "ring", 4754 + "untrusted", 4755 + ] 4756 + 4757 + [[package]] 4758 + name = "weezl" 4759 + version = "0.1.7" 4760 + source = "registry+https://github.com/rust-lang/crates.io-index" 4761 + checksum = "9193164d4de03a926d909d3bc7c30543cecb35400c02114792c2cae20d5e2dbb" 4762 + 4763 + [[package]] 4764 + name = "which" 4765 + version = "3.1.1" 4766 + source = "registry+https://github.com/rust-lang/crates.io-index" 4767 + checksum = "d011071ae14a2f6671d0b74080ae0cd8ebf3a6f8c9589a2cd45f23126fe29724" 4768 + dependencies = [ 4769 + "failure", 4770 + "libc", 4771 + ] 4772 + 4773 + [[package]] 4774 + name = "which" 4775 + version = "4.4.0" 4776 + source = "registry+https://github.com/rust-lang/crates.io-index" 4777 + checksum = "2441c784c52b289a054b7201fc93253e288f094e2f4be9058343127c4226a269" 4778 + dependencies = [ 4779 + "either", 4780 + "libc", 4781 + "once_cell", 4782 + ] 4783 + 4784 + [[package]] 4785 + name = "whoami" 4786 + version = "1.4.0" 4787 + source = "registry+https://github.com/rust-lang/crates.io-index" 4788 + checksum = "2c70234412ca409cc04e864e89523cb0fc37f5e1344ebed5a3ebf4192b6b9f68" 4789 + dependencies = [ 4790 + "wasm-bindgen", 4791 + "web-sys", 4792 + ] 4793 + 4794 + [[package]] 4795 + name = "widestring" 4796 + version = "0.4.3" 4797 + source = "registry+https://github.com/rust-lang/crates.io-index" 4798 + checksum = "c168940144dd21fd8046987c16a46a33d5fc84eec29ef9dcddc2ac9e31526b7c" 4799 + 4800 + [[package]] 4801 + name = "widestring" 4802 + version = "1.0.2" 4803 + source = "registry+https://github.com/rust-lang/crates.io-index" 4804 + checksum = "653f141f39ec16bba3c5abe400a0c60da7468261cc2cbf36805022876bc721a8" 4805 + 4806 + [[package]] 4807 + name = "winapi" 4808 + version = "0.2.8" 4809 + source = "registry+https://github.com/rust-lang/crates.io-index" 4810 + checksum = "167dc9d6949a9b857f3451275e911c3f44255842c1f7a76f33c55103a909087a" 4811 + 4812 + [[package]] 4813 + name = "winapi" 4814 + version = "0.3.9" 4815 + source = "registry+https://github.com/rust-lang/crates.io-index" 4816 + checksum = "5c839a674fcd7a98952e593242ea400abe93992746761e38641405d28b00f419" 4817 + dependencies = [ 4818 + "winapi-i686-pc-windows-gnu", 4819 + "winapi-x86_64-pc-windows-gnu", 4820 + ] 4821 + 4822 + [[package]] 4823 + name = "winapi-build" 4824 + version = "0.1.1" 4825 + source = "registry+https://github.com/rust-lang/crates.io-index" 4826 + checksum = "2d315eee3b34aca4797b2da6b13ed88266e6d612562a0c46390af8299fc699bc" 4827 + 4828 + [[package]] 4829 + name = "winapi-i686-pc-windows-gnu" 4830 + version = "0.4.0" 4831 + source = "registry+https://github.com/rust-lang/crates.io-index" 4832 + checksum = "ac3b87c63620426dd9b991e5ce0329eff545bccbbb34f3be09ff6fb6ab51b7b6" 4833 + 4834 + [[package]] 4835 + name = "winapi-util" 4836 + version = "0.1.5" 4837 + source = "registry+https://github.com/rust-lang/crates.io-index" 4838 + checksum = "70ec6ce85bb158151cae5e5c87f95a8e97d2c0c4b001223f33a334e3ce5de178" 4839 + dependencies = [ 4840 + "winapi 0.3.9", 4841 + ] 4842 + 4843 + [[package]] 4844 + name = "winapi-wsapoll" 4845 + version = "0.1.1" 4846 + source = "registry+https://github.com/rust-lang/crates.io-index" 4847 + checksum = "44c17110f57155602a80dca10be03852116403c9ff3cd25b079d666f2aa3df6e" 4848 + dependencies = [ 4849 + "winapi 0.3.9", 4850 + ] 4851 + 4852 + [[package]] 4853 + name = "winapi-x86_64-pc-windows-gnu" 4854 + version = "0.4.0" 4855 + source = "registry+https://github.com/rust-lang/crates.io-index" 4856 + checksum = "712e227841d057c1ee1cd2fb22fa7e5a5461ae8e48fa2ca79ec42cfc1931183f" 4857 + 4858 + [[package]] 4859 + name = "windows" 4860 + version = "0.48.0" 4861 + source = "registry+https://github.com/rust-lang/crates.io-index" 4862 + checksum = "e686886bc078bc1b0b600cac0147aadb815089b6e4da64016cbd754b6342700f" 4863 + dependencies = [ 4864 + "windows-targets 0.48.0", 4865 + ] 4866 + 4867 + [[package]] 4868 + name = "windows-service" 4869 + version = "0.4.0" 4870 + source = "registry+https://github.com/rust-lang/crates.io-index" 4871 + checksum = "0c643e10139d127d30d6d753398c8a6f0a43532e8370f6c9d29ebbff29b984ab" 4872 + dependencies = [ 4873 + "bitflags", 4874 + "err-derive", 4875 + "widestring 0.4.3", 4876 + "winapi 0.3.9", 4877 + ] 4878 + 4879 + [[package]] 4880 + name = "windows-sys" 4881 + version = "0.28.0" 4882 + source = "registry+https://github.com/rust-lang/crates.io-index" 4883 + checksum = "82ca39602d5cbfa692c4b67e3bcbb2751477355141c1ed434c94da4186836ff6" 4884 + dependencies = [ 4885 + "windows_aarch64_msvc 0.28.0", 4886 + "windows_i686_gnu 0.28.0", 4887 + "windows_i686_msvc 0.28.0", 4888 + "windows_x86_64_gnu 0.28.0", 4889 + "windows_x86_64_msvc 0.28.0", 4890 + ] 4891 + 4892 + [[package]] 4893 + name = "windows-sys" 4894 + version = "0.42.0" 4895 + source = "registry+https://github.com/rust-lang/crates.io-index" 4896 + checksum = "5a3e1820f08b8513f676f7ab6c1f99ff312fb97b553d30ff4dd86f9f15728aa7" 4897 + dependencies = [ 4898 + "windows_aarch64_gnullvm 0.42.2", 4899 + "windows_aarch64_msvc 0.42.2", 4900 + "windows_i686_gnu 0.42.2", 4901 + "windows_i686_msvc 0.42.2", 4902 + "windows_x86_64_gnu 0.42.2", 4903 + "windows_x86_64_gnullvm 0.42.2", 4904 + "windows_x86_64_msvc 0.42.2", 4905 + ] 4906 + 4907 + [[package]] 4908 + name = "windows-sys" 4909 + version = "0.45.0" 4910 + source = "registry+https://github.com/rust-lang/crates.io-index" 4911 + checksum = "75283be5efb2831d37ea142365f009c02ec203cd29a3ebecbc093d52315b66d0" 4912 + dependencies = [ 4913 + "windows-targets 0.42.2", 4914 + ] 4915 + 4916 + [[package]] 4917 + name = "windows-sys" 4918 + version = "0.48.0" 4919 + source = "registry+https://github.com/rust-lang/crates.io-index" 4920 + checksum = "677d2418bec65e3338edb076e806bc1ec15693c5d0104683f2efe857f61056a9" 4921 + dependencies = [ 4922 + "windows-targets 0.48.0", 4923 + ] 4924 + 4925 + [[package]] 4926 + name = "windows-targets" 4927 + version = "0.42.2" 4928 + source = "registry+https://github.com/rust-lang/crates.io-index" 4929 + checksum = "8e5180c00cd44c9b1c88adb3693291f1cd93605ded80c250a75d472756b4d071" 4930 + dependencies = [ 4931 + "windows_aarch64_gnullvm 0.42.2", 4932 + "windows_aarch64_msvc 0.42.2", 4933 + "windows_i686_gnu 0.42.2", 4934 + "windows_i686_msvc 0.42.2", 4935 + "windows_x86_64_gnu 0.42.2", 4936 + "windows_x86_64_gnullvm 0.42.2", 4937 + "windows_x86_64_msvc 0.42.2", 4938 + ] 4939 + 4940 + [[package]] 4941 + name = "windows-targets" 4942 + version = "0.48.0" 4943 + source = "registry+https://github.com/rust-lang/crates.io-index" 4944 + checksum = "7b1eb6f0cd7c80c79759c929114ef071b87354ce476d9d94271031c0497adfd5" 4945 + dependencies = [ 4946 + "windows_aarch64_gnullvm 0.48.0", 4947 + "windows_aarch64_msvc 0.48.0", 4948 + "windows_i686_gnu 0.48.0", 4949 + "windows_i686_msvc 0.48.0", 4950 + "windows_x86_64_gnu 0.48.0", 4951 + "windows_x86_64_gnullvm 0.48.0", 4952 + "windows_x86_64_msvc 0.48.0", 4953 + ] 4954 + 4955 + [[package]] 4956 + name = "windows-win" 4957 + version = "2.4.1" 4958 + source = "registry+https://github.com/rust-lang/crates.io-index" 4959 + checksum = "8d4243ec23afe4e9b4e668b3c0a0e973f1b8265f6a46223cfcbc16fd267480c0" 4960 + dependencies = [ 4961 + "winapi 0.3.9", 4962 + ] 4963 + 4964 + [[package]] 4965 + name = "windows_aarch64_gnullvm" 4966 + version = "0.42.2" 4967 + source = "registry+https://github.com/rust-lang/crates.io-index" 4968 + checksum = "597a5118570b68bc08d8d59125332c54f1ba9d9adeedeef5b99b02ba2b0698f8" 4969 + 4970 + [[package]] 4971 + name = "windows_aarch64_gnullvm" 4972 + version = "0.48.0" 4973 + source = "registry+https://github.com/rust-lang/crates.io-index" 4974 + checksum = "91ae572e1b79dba883e0d315474df7305d12f569b400fcf90581b06062f7e1bc" 4975 + 4976 + [[package]] 4977 + name = "windows_aarch64_msvc" 4978 + version = "0.28.0" 4979 + source = "registry+https://github.com/rust-lang/crates.io-index" 4980 + checksum = "52695a41e536859d5308cc613b4a022261a274390b25bd29dfff4bf08505f3c2" 4981 + 4982 + [[package]] 4983 + name = "windows_aarch64_msvc" 4984 + version = "0.42.2" 4985 + source = "registry+https://github.com/rust-lang/crates.io-index" 4986 + checksum = "e08e8864a60f06ef0d0ff4ba04124db8b0fb3be5776a5cd47641e942e58c4d43" 4987 + 4988 + [[package]] 4989 + name = "windows_aarch64_msvc" 4990 + version = "0.48.0" 4991 + source = "registry+https://github.com/rust-lang/crates.io-index" 4992 + checksum = "b2ef27e0d7bdfcfc7b868b317c1d32c641a6fe4629c171b8928c7b08d98d7cf3" 4993 + 4994 + [[package]] 4995 + name = "windows_i686_gnu" 4996 + version = "0.28.0" 4997 + source = "registry+https://github.com/rust-lang/crates.io-index" 4998 + checksum = "f54725ac23affef038fecb177de6c9bf065787c2f432f79e3c373da92f3e1d8a" 4999 + 5000 + [[package]] 5001 + name = "windows_i686_gnu" 5002 + version = "0.42.2" 5003 + source = "registry+https://github.com/rust-lang/crates.io-index" 5004 + checksum = "c61d927d8da41da96a81f029489353e68739737d3beca43145c8afec9a31a84f" 5005 + 5006 + [[package]] 5007 + name = "windows_i686_gnu" 5008 + version = "0.48.0" 5009 + source = "registry+https://github.com/rust-lang/crates.io-index" 5010 + checksum = "622a1962a7db830d6fd0a69683c80a18fda201879f0f447f065a3b7467daa241" 5011 + 5012 + [[package]] 5013 + name = "windows_i686_msvc" 5014 + version = "0.28.0" 5015 + source = "registry+https://github.com/rust-lang/crates.io-index" 5016 + checksum = "51d5158a43cc43623c0729d1ad6647e62fa384a3d135fd15108d37c683461f64" 5017 + 5018 + [[package]] 5019 + name = "windows_i686_msvc" 5020 + version = "0.42.2" 5021 + source = "registry+https://github.com/rust-lang/crates.io-index" 5022 + checksum = "44d840b6ec649f480a41c8d80f9c65108b92d89345dd94027bfe06ac444d1060" 5023 + 5024 + [[package]] 5025 + name = "windows_i686_msvc" 5026 + version = "0.48.0" 5027 + source = "registry+https://github.com/rust-lang/crates.io-index" 5028 + checksum = "4542c6e364ce21bf45d69fdd2a8e455fa38d316158cfd43b3ac1c5b1b19f8e00" 5029 + 5030 + [[package]] 5031 + name = "windows_x86_64_gnu" 5032 + version = "0.28.0" 5033 + source = "registry+https://github.com/rust-lang/crates.io-index" 5034 + checksum = "bc31f409f565611535130cfe7ee8e6655d3fa99c1c61013981e491921b5ce954" 5035 + 5036 + [[package]] 5037 + name = "windows_x86_64_gnu" 5038 + version = "0.42.2" 5039 + source = "registry+https://github.com/rust-lang/crates.io-index" 5040 + checksum = "8de912b8b8feb55c064867cf047dda097f92d51efad5b491dfb98f6bbb70cb36" 5041 + 5042 + [[package]] 5043 + name = "windows_x86_64_gnu" 5044 + version = "0.48.0" 5045 + source = "registry+https://github.com/rust-lang/crates.io-index" 5046 + checksum = "ca2b8a661f7628cbd23440e50b05d705db3686f894fc9580820623656af974b1" 5047 + 5048 + [[package]] 5049 + name = "windows_x86_64_gnullvm" 5050 + version = "0.42.2" 5051 + source = "registry+https://github.com/rust-lang/crates.io-index" 5052 + checksum = "26d41b46a36d453748aedef1486d5c7a85db22e56aff34643984ea85514e94a3" 5053 + 5054 + [[package]] 5055 + name = "windows_x86_64_gnullvm" 5056 + version = "0.48.0" 5057 + source = "registry+https://github.com/rust-lang/crates.io-index" 5058 + checksum = "7896dbc1f41e08872e9d5e8f8baa8fdd2677f29468c4e156210174edc7f7b953" 5059 + 5060 + [[package]] 5061 + name = "windows_x86_64_msvc" 5062 + version = "0.28.0" 5063 + source = "registry+https://github.com/rust-lang/crates.io-index" 5064 + checksum = "3f2b8c7cbd3bfdddd9ab98769f9746a7fad1bca236554cd032b78d768bc0e89f" 5065 + 5066 + [[package]] 5067 + name = "windows_x86_64_msvc" 5068 + version = "0.42.2" 5069 + source = "registry+https://github.com/rust-lang/crates.io-index" 5070 + checksum = "9aec5da331524158c6d1a4ac0ab1541149c0b9505fde06423b02f5ef0106b9f0" 5071 + 5072 + [[package]] 5073 + name = "windows_x86_64_msvc" 5074 + version = "0.48.0" 5075 + source = "registry+https://github.com/rust-lang/crates.io-index" 5076 + checksum = "1a515f5799fe4961cb532f983ce2b23082366b898e52ffbce459c86f67c8378a" 5077 + 5078 + [[package]] 5079 + name = "winnow" 5080 + version = "0.4.1" 5081 + source = "registry+https://github.com/rust-lang/crates.io-index" 5082 + checksum = "ae8970b36c66498d8ff1d66685dc86b91b29db0c7739899012f63a63814b4b28" 5083 + dependencies = [ 5084 + "memchr", 5085 + ] 5086 + 5087 + [[package]] 5088 + name = "winreg" 5089 + version = "0.6.2" 5090 + source = "registry+https://github.com/rust-lang/crates.io-index" 5091 + checksum = "b2986deb581c4fe11b621998a5e53361efe6b48a151178d0cd9eeffa4dc6acc9" 5092 + dependencies = [ 5093 + "winapi 0.3.9", 5094 + ] 5095 + 5096 + [[package]] 5097 + name = "winreg" 5098 + version = "0.10.1" 5099 + source = "registry+https://github.com/rust-lang/crates.io-index" 5100 + checksum = "80d0f4e272c85def139476380b12f9ac60926689dd2e01d4923222f40580869d" 5101 + dependencies = [ 5102 + "winapi 0.3.9", 5103 + ] 5104 + 5105 + [[package]] 5106 + name = "winres" 5107 + version = "0.1.12" 5108 + source = "registry+https://github.com/rust-lang/crates.io-index" 5109 + checksum = "b68db261ef59e9e52806f688020631e987592bd83619edccda9c47d42cde4f6c" 5110 + dependencies = [ 5111 + "toml 0.5.11", 5112 + ] 5113 + 5114 + [[package]] 5115 + name = "ws2_32-sys" 5116 + version = "0.2.1" 5117 + source = "registry+https://github.com/rust-lang/crates.io-index" 5118 + checksum = "d59cefebd0c892fa2dd6de581e937301d8552cb44489cdff035c6187cb63fa5e" 5119 + dependencies = [ 5120 + "winapi 0.2.8", 5121 + "winapi-build", 5122 + ] 5123 + 5124 + [[package]] 5125 + name = "x11" 5126 + version = "2.21.0" 5127 + source = "registry+https://github.com/rust-lang/crates.io-index" 5128 + checksum = "502da5464ccd04011667b11c435cb992822c2c0dbde1770c988480d312a0db2e" 5129 + dependencies = [ 5130 + "libc", 5131 + "pkg-config", 5132 + ] 5133 + 5134 + [[package]] 5135 + name = "x11-clipboard" 5136 + version = "0.5.3" 5137 + source = "registry+https://github.com/rust-lang/crates.io-index" 5138 + checksum = "473068b7b80ac86a18328824f1054e5e007898c47b5bbc281bd7abe32bc3653c" 5139 + dependencies = [ 5140 + "xcb", 5141 + ] 5142 + 5143 + [[package]] 5144 + name = "x11rb" 5145 + version = "0.9.0" 5146 + source = "registry+https://github.com/rust-lang/crates.io-index" 5147 + checksum = "6e99be55648b3ae2a52342f9a870c0e138709a3493261ce9b469afe6e4df6d8a" 5148 + dependencies = [ 5149 + "gethostname", 5150 + "nix 0.22.3", 5151 + "winapi 0.3.9", 5152 + "winapi-wsapoll", 5153 + ] 5154 + 5155 + [[package]] 5156 + name = "xcb" 5157 + version = "0.10.1" 5158 + source = "registry+https://github.com/rust-lang/crates.io-index" 5159 + checksum = "771e2b996df720cd1c6dd9ff90f62d91698fd3610cc078388d0564bdd6622a9c" 5160 + dependencies = [ 5161 + "libc", 5162 + "log", 5163 + "quick-xml", 5164 + ] 5165 + 5166 + [[package]] 5167 + name = "zstd" 5168 + version = "0.9.2+zstd.1.5.1" 5169 + source = "registry+https://github.com/rust-lang/crates.io-index" 5170 + checksum = "2390ea1bf6c038c39674f22d95f0564725fc06034a47129179810b2fc58caa54" 5171 + dependencies = [ 5172 + "zstd-safe", 5173 + ] 5174 + 5175 + [[package]] 5176 + name = "zstd-safe" 5177 + version = "4.1.3+zstd.1.5.1" 5178 + source = "registry+https://github.com/rust-lang/crates.io-index" 5179 + checksum = "e99d81b99fb3c2c2c794e3fe56c305c63d5173a16a46b5850b07c935ffc7db79" 5180 + dependencies = [ 5181 + "libc", 5182 + "zstd-sys", 5183 + ] 5184 + 5185 + [[package]] 5186 + name = "zstd-sys" 5187 + version = "1.6.2+zstd.1.5.1" 5188 + source = "registry+https://github.com/rust-lang/crates.io-index" 5189 + checksum = "2daf2f248d9ea44454bfcb2516534e8b8ad2fc91bf818a1885495fc42bc8ac9f" 5190 + dependencies = [ 5191 + "cc", 5192 + "libc", 5193 + ]
-35
pkgs/applications/networking/remote/rustdesk/cargo.patch
··· 1 - diff --git a/Cargo.lock b/Cargo.lock 2 - index fb17c7e..ef157e5 100644 3 - --- a/Cargo.lock 4 - +++ b/Cargo.lock 5 - @@ -2332,10 +2332,10 @@ dependencies = [ 6 - [[package]] 7 - name = "magnum-opus" 8 - version = "0.4.0" 9 - -source = "git+https://github.com/open-trade/magnum-opus#3c3d0b86ae95c84930bebffe4bcb03b3bd83342b" 10 - +source = "git+https://github.com/TheRadioGuy/magnum-opus#171e1d021004626f7444d1e39b98f50bc3cb2604" 11 - dependencies = [ 12 - - "bindgen", 13 - - "target_build_utils", 14 - + "libc", 15 - + "opusic-sys", 16 - ] 17 - 18 - [[package]] 19 - @@ -2796,6 +2796,16 @@ version = "0.1.5" 20 - source = "registry+https://github.com/rust-lang/crates.io-index" 21 - checksum = "ff011a302c396a5197692431fc1948019154afc178baf7d8e37367442a4601cf" 22 - 23 - +[[package]] 24 - +name = "opusic-sys" 25 - +version = "0.3.6" 26 - +source = "registry+https://github.com/rust-lang/crates.io-index" 27 - +checksum = "5eace752ce07a037241dba8f02c654799f051e431b27028056bcb480e83b54f5" 28 - +dependencies = [ 29 - + "cmake", 30 - + "libc", 31 - +] 32 - + 33 - [[package]] 34 - name = "os_str_bytes" 35 - version = "6.0.0" 36 1 diff --git a/Cargo.toml b/Cargo.toml 37 2 index 1b715bd..960e8da 100644 38 3 --- a/Cargo.toml
+59 -39
pkgs/applications/networking/remote/rustdesk/default.nix
··· 1 1 { lib 2 - , stdenv 3 2 , fetchFromGitHub 4 3 , makeDesktopItem 5 4 , copyDesktopItems ··· 8 7 , cmake 9 8 , yasm 10 9 , nasm 11 - , zip 12 10 , pkg-config 13 11 , clang 14 12 , gtk3 ··· 43 41 ./fix-for-rust-1.65.diff 44 42 ]; 45 43 46 - cargoSha256 = "sha256-1OMWEk+DerltF7kwdo4d04rbgIFLHBRq3vZaL7jtrdE="; 44 + LIBCLANG_PATH = "${llvmPackages.libclang.lib}/lib"; 47 45 48 - LIBCLANG_PATH="${llvmPackages.libclang.lib}/lib"; 46 + cargoLock = { 47 + lockFile = ./Cargo.lock; 48 + outputHashes = { 49 + "confy-0.4.0" = "sha256-e91cvEixhpPzIthAxzTa3fDY6eCsHUy/eZQAqs7QTDo="; 50 + "parity-tokio-ipc-0.7.3-1" = "sha256-eULJePtBu0iBI3It/bPH0h82Obsb1PJALgwYwrnCFYI="; 51 + "rdev-0.5.0-2" = "sha256-7CEZ2wIM4QAPfY1tGKqXfHplTaxHnccVqFRPjY21Svo="; 52 + "tokio-socks-0.5.1-1" = "sha256-45QQ6FrhGU9uEhbKXTKd/mY6MDumO6p46NmlakdyDQk="; 53 + "libappindicator-0.6.1" = "sha256-JGnnZrcwbh8WJ6+/4bYhfD3HvgF2C7XaaGb6TaMRWdw="; 54 + "magnum-opus-0.4.0" = "sha256-U5uuN4YolOYDnFNbtPpwYefcBDTUUyioui0UCcW8dyo="; 55 + "rust-pulsectl-0.2.12" = "sha256-8jXTspWvjONFcvw9/Z8C43g4BuGZ3rsG32tvLMQbtbM="; 56 + "sciter-rs-0.5.57" = "sha256-ZZnZDhMjK0LjgmK0da1yvB0uoKueLhhhQtzmjoN+1R0="; 57 + "systray-0.4.1" = "sha256-p1PMr/8oS6zHx4+Ng4zCqt0xZ57cq3wAu6/agyWq5Jw="; 58 + }; 59 + }; 49 60 50 61 # Change magnus-opus version to upstream so that it does not use 51 62 # vcpkg for libopus since it does not work. ··· 55 66 56 67 # Manually simulate a vcpkg installation so that it can link the libaries 57 68 # properly. 58 - postUnpack = let 59 - vcpkg_target = "x64-linux"; 69 + postUnpack = 70 + let 71 + vcpkg_target = "x64-linux"; 60 72 61 - updates_vcpkg_file = writeText "update_vcpkg_rustdesk" 62 - '' 63 - Package : libyuv 64 - Architecture : ${vcpkg_target} 65 - Version : 1.0 66 - Status : is installed 73 + updates_vcpkg_file = writeText "update_vcpkg_rustdesk" 74 + '' 75 + Package : libyuv 76 + Architecture : ${vcpkg_target} 77 + Version : 1.0 78 + Status : is installed 67 79 68 - Package : libvpx 69 - Architecture : ${vcpkg_target} 70 - Version : 1.0 71 - Status : is installed 72 - ''; 73 - in '' 74 - export VCPKG_ROOT="$TMP/vcpkg"; 80 + Package : libvpx 81 + Architecture : ${vcpkg_target} 82 + Version : 1.0 83 + Status : is installed 84 + ''; 85 + in 86 + '' 87 + export VCPKG_ROOT="$TMP/vcpkg"; 75 88 76 - mkdir -p $VCPKG_ROOT/.vcpkg-root 77 - mkdir -p $VCPKG_ROOT/installed/${vcpkg_target}/lib 78 - mkdir -p $VCPKG_ROOT/installed/vcpkg/updates 79 - ln -s ${updates_vcpkg_file} $VCPKG_ROOT/installed/vcpkg/status 80 - mkdir -p $VCPKG_ROOT/installed/vcpkg/info 81 - touch $VCPKG_ROOT/installed/vcpkg/info/libyuv_1.0_${vcpkg_target}.list 82 - touch $VCPKG_ROOT/installed/vcpkg/info/libvpx_1.0_${vcpkg_target}.list 89 + mkdir -p $VCPKG_ROOT/.vcpkg-root 90 + mkdir -p $VCPKG_ROOT/installed/${vcpkg_target}/lib 91 + mkdir -p $VCPKG_ROOT/installed/vcpkg/updates 92 + ln -s ${updates_vcpkg_file} $VCPKG_ROOT/installed/vcpkg/status 93 + mkdir -p $VCPKG_ROOT/installed/vcpkg/info 94 + touch $VCPKG_ROOT/installed/vcpkg/info/libyuv_1.0_${vcpkg_target}.list 95 + touch $VCPKG_ROOT/installed/vcpkg/info/libvpx_1.0_${vcpkg_target}.list 83 96 84 - ln -s ${libvpx.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/ 85 - ln -s ${libyuv.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/ 86 - ''; 97 + ln -s ${libvpx.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/ 98 + ln -s ${libyuv.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/ 99 + ''; 87 100 88 101 nativeBuildInputs = [ pkg-config cmake makeWrapper copyDesktopItems yasm nasm clang wrapGAppsHook ]; 89 102 buildInputs = [ alsa-lib pulseaudio libXfixes libxcb xdotool gtk3 libvpx libopus libXtst libyuv ]; ··· 91 104 # Checks require an active X display. 92 105 doCheck = false; 93 106 94 - desktopItems = [ (makeDesktopItem { 95 - name = "rustdesk"; 96 - exec = meta.mainProgram; 97 - icon = "rustdesk"; 98 - desktopName = "RustDesk"; 99 - comment = meta.description; 100 - genericName = "Remote Desktop"; 101 - categories = ["Network"]; 102 - }) ]; 107 + desktopItems = [ 108 + (makeDesktopItem { 109 + name = "rustdesk"; 110 + exec = meta.mainProgram; 111 + icon = "rustdesk"; 112 + desktopName = "RustDesk"; 113 + comment = meta.description; 114 + genericName = "Remote Desktop"; 115 + categories = [ "Network" ]; 116 + }) 117 + ]; 118 + 119 + postPatch = '' 120 + rm Cargo.lock 121 + ln -s ${./Cargo.lock} Cargo.lock 122 + ''; 103 123 104 124 # Add static ui resources and libsciter to same folder as binary so that it 105 125 # can find them. ··· 122 142 description = "Yet another remote desktop software"; 123 143 homepage = "https://rustdesk.com"; 124 144 license = licenses.gpl3Only; 125 - maintainers = with maintainers; [ leixb ]; 145 + maintainers = with maintainers; [ ocfox leixb ]; 126 146 platforms = [ "x86_64-linux" ]; 127 147 mainProgram = "rustdesk"; 128 148 };
+5
pkgs/applications/virtualization/docker/default.nix
··· 254 254 meta = with lib; { 255 255 homepage = "https://www.docker.com/"; 256 256 description = "An open source project to pack, ship and run any application as a lightweight container"; 257 + longDescription = '' 258 + Docker is a platform designed to help developers build, share, and run modern applications. 259 + 260 + To enable the docker daemon on NixOS, set the `virtualisation.docker.enable` option to `true`. 261 + ''; 257 262 license = licenses.asl20; 258 263 maintainers = with maintainers; [ offline tailhook vdemeester periklis mikroskeem maxeaubrey ]; 259 264 };
+6 -11
pkgs/build-support/fetchmavenartifact/default.nix
··· 37 37 assert (repos != []) || (url != "") || (urls != []); 38 38 39 39 let 40 - name_ = 41 - lib.concatStrings [ 42 - (lib.replaceStrings ["."] ["_"] groupId) "_" 43 - (lib.replaceStrings ["."] ["_"] artifactId) "-" 44 - version 45 - ]; 46 - suffix = if isNull classifier then "" else "-${classifier}"; 40 + pname = (lib.replaceStrings [ "." ] [ "_" ] groupId) + "_" + (lib.replaceStrings [ "." ] [ "_" ] artifactId); 41 + suffix = lib.optionalString (classifier != null) "-${classifier}"; 47 42 filename = "${artifactId}-${version}${suffix}.jar"; 48 43 mkJarUrl = repoUrl: 49 44 lib.concatStringsSep "/" [ ··· 59 54 else map mkJarUrl repos; 60 55 jar = 61 56 fetchurl ( 62 - builtins.removeAttrs args ["groupId" "artifactId" "version" "classifier" "repos" "url" ] 63 - // { urls = urls_; name = "${name_}.jar"; } 57 + builtins.removeAttrs args [ "groupId" "artifactId" "version" "classifier" "repos" "url" ] 58 + // { urls = urls_; name = "${pname}-${version}.jar"; } 64 59 ); 65 60 in 66 61 stdenv.mkDerivation { 67 - name = name_; 68 - phases = "installPhase fixupPhase"; 62 + inherit pname version; 63 + dontUnpack = true; 69 64 # By moving the jar to $out/share/java we make it discoverable by java 70 65 # packages packages that mention this derivation in their buildInputs. 71 66 installPhase = ''
+5 -1
pkgs/development/compilers/llvm/rocm/llvm.nix
··· 24 24 , targetDir ? "llvm" 25 25 , targetProjects ? [ ] 26 26 , targetRuntimes ? [ ] 27 + # "NATIVE" resolves into x86 or aarch64 depending on stdenv 28 + , llvmTargetsToBuild ? [ "NATIVE" ] 27 29 , extraPatches ? [ ] 28 30 , extraNativeBuildInputs ? [ ] 29 31 , extraBuildInputs ? [ ] ··· 46 48 if stdenv.isx86_64 then "X86" 47 49 else if stdenv.isAarch64 then "AArch64" 48 50 else throw "Unsupported ROCm LLVM platform"; 51 + inferNativeTarget = t: if t == "NATIVE" then llvmNativeTarget else t; 52 + llvmTargetsToBuild' = [ "AMDGPU" ] ++ builtins.map inferNativeTarget llvmTargetsToBuild; 49 53 in stdenv.mkDerivation (finalAttrs: { 50 54 pname = "rocm-llvm-${targetName}"; 51 55 version = "5.4.4"; ··· 98 102 sourceRoot = "${finalAttrs.src.name}/${targetDir}"; 99 103 100 104 cmakeFlags = [ 101 - "-DLLVM_TARGETS_TO_BUILD=AMDGPU;${llvmNativeTarget}" 105 + "-DLLVM_TARGETS_TO_BUILD=${builtins.concatStringsSep ";" llvmTargetsToBuild'}" 102 106 ] ++ lib.optionals (finalAttrs.passthru.isLLVM && targetProjects != [ ]) [ 103 107 "-DLLVM_ENABLE_PROJECTS=${lib.concatStringsSep ";" targetProjects}" 104 108 ] ++ lib.optionals ((finalAttrs.passthru.isLLVM || targetDir == "runtimes") && targetRuntimes != [ ]) [
+4 -4
pkgs/development/libraries/libiio/default.nix
··· 37 37 flex 38 38 bison 39 39 pkg-config 40 - ]; 40 + python 41 + ] ++ lib.optional python.isPy3k python.pkgs.setuptools; 41 42 42 43 buildInputs = [ 43 - python 44 44 libxml2 45 45 libusb1 46 - ] ++ lib.optional python.isPy3k python.pkgs.setuptools 47 - ++ lib.optional avahiSupport avahi 46 + ] ++ lib.optional avahiSupport avahi 48 47 ++ lib.optional stdenv.isLinux libaio 49 48 ++ lib.optionals stdenv.isDarwin [ CFNetwork CoreServices ]; 50 49 51 50 cmakeFlags = [ 52 51 "-DUDEV_RULES_INSTALL_DIR=${placeholder "out"}/lib/udev/rules.d" 52 + "-DPython_EXECUTABLE=${python.pythonForBuild.interpreter}" 53 53 "-DPYTHON_BINDINGS=on" 54 54 # osx framework is disabled, 55 55 # the linux-like directory structure is used for proper output splitting
+19 -10
pkgs/development/libraries/ndpi/default.nix
··· 1 1 { lib 2 2 , stdenv 3 - , fetchFromGitHub 4 - , which 5 3 , autoconf 6 4 , automake 5 + , fetchFromGitHub 6 + , json_c 7 + , libpcap 7 8 , libtool 8 - , libpcap 9 - , json_c 10 - , pkg-config }: 9 + , pkg-config 10 + , which 11 + }: 11 12 12 13 stdenv.mkDerivation rec { 13 14 pname = "ndpi"; 14 - version = "4.2"; 15 + version = "4.6"; 15 16 16 17 src = fetchFromGitHub { 17 18 owner = "ntop"; 18 19 repo = "nDPI"; 19 - rev = version; 20 - sha256 = "sha256-ZWWuyPGl+hbrfXdtPvCBqMReuJ4FiGx+qiI7qCz6wtQ="; 20 + rev = "refs/tags/${version}"; 21 + hash = "sha256-S0lVh5FZewPbYG/1ikI2RroCSC7OI8Xmfeq73hYCHnY="; 21 22 }; 22 23 23 24 configureScript = "./autogen.sh"; 24 25 25 - nativeBuildInputs = [ which autoconf automake libtool pkg-config ]; 26 + nativeBuildInputs = [ 27 + autoconf 28 + automake 29 + libtool 30 + pkg-config 31 + which 32 + ]; 33 + 26 34 buildInputs = [ 35 + json_c 27 36 libpcap 28 - json_c 29 37 ]; 30 38 31 39 meta = with lib; { ··· 34 42 nDPI is a library for deep-packet inspection based on OpenDPI. 35 43 ''; 36 44 homepage = "https://www.ntop.org/products/deep-packet-inspection/ndpi/"; 45 + changelog = "https://github.com/ntop/nDPI/blob/${version}/CHANGELOG.md"; 37 46 license = with licenses; [ lgpl3Plus bsd3 ]; 38 47 maintainers = with maintainers; [ takikawa ]; 39 48 mainProgram = "ndpiReader";
+5 -7
pkgs/development/libraries/rapidfuzz-cpp/default.nix
··· 5 5 , catch2_3 6 6 }: 7 7 8 - stdenv.mkDerivation rec { 8 + stdenv.mkDerivation (finalAttrs: { 9 9 pname = "rapidfuzz-cpp"; 10 10 version = "1.10.4"; 11 11 12 12 src = fetchFromGitHub { 13 13 owner = "maxbachmann"; 14 14 repo = "rapidfuzz-cpp"; 15 - rev = "v${version}"; 15 + rev = "v${finalAttrs.version}"; 16 16 hash = "sha256-I7MdeLs+J5a57ypgUJIW0/pSFPzK4nZA4ZrVRdKX7J4="; 17 17 }; 18 18 ··· 20 20 cmake 21 21 ]; 22 22 23 - cmakeFlags = lib.optionals doCheck [ 23 + cmakeFlags = lib.optionals finalAttrs.finalPackage.doCheck [ 24 24 "-DRAPIDFUZZ_BUILD_TESTING=ON" 25 25 ]; 26 26 ··· 33 33 catch2_3 34 34 ]; 35 35 36 - doCheck = stdenv.buildPlatform.canExecute stdenv.hostPlatform; 37 - 38 36 meta = { 39 37 description = "Rapid fuzzy string matching in C++ using the Levenshtein Distance"; 40 38 homepage = "https://github.com/maxbachmann/rapidfuzz-cpp"; 41 - changelog = "https://github.com/maxbachmann/rapidfuzz-cpp/blob/${src.rev}/CHANGELOG.md"; 39 + changelog = "https://github.com/maxbachmann/rapidfuzz-cpp/blob/${finalAttrs.src.rev}/CHANGELOG.md"; 42 40 license = lib.licenses.mit; 43 41 maintainers = with lib.maintainers; [ dotlambda ]; 44 42 platforms = lib.platforms.unix; 45 43 }; 46 - } 44 + })
+3 -1
pkgs/development/libraries/science/math/tensorrt/extension.nix
··· 27 27 defaultBuild = { "tensorrt" = if allBuilds ? ${computeName tensorRTDefaultVersion} 28 28 then allBuilds.${computeName tensorRTDefaultVersion} 29 29 else throw "tensorrt-${tensorRTDefaultVersion} does not support your cuda version ${cudaVersion}"; }; 30 - in allBuilds // defaultBuild; 30 + in { 31 + inherit buildTensorRTPackage; 32 + } // allBuilds // defaultBuild; 31 33 32 34 tensorRTVersions = { 33 35 "8.4.0" = [
+11 -2
pkgs/development/ocaml-modules/torch/default.nix
··· 2 2 , stdenv 3 3 , buildDunePackage 4 4 , fetchFromGitHub 5 + , fetchpatch 5 6 , cmdliner 6 7 , ctypes 7 8 , dune-configurator ··· 24 25 25 26 src = fetchFromGitHub { 26 27 owner = "LaurentMazare"; 27 - repo = "ocaml-${pname}"; 28 - rev = version; 28 + repo = "ocaml-${pname}"; 29 + rev = version; 29 30 hash = "sha256-z/9NUBjeFWE63Z/e8OyzDiy8hrn6qzjaiBH8G9MPeos="; 30 31 }; 32 + 33 + patches = [ 34 + # Pytorch 2.0 support. Drop when it reaches a release 35 + (fetchpatch { 36 + url = "https://github.com/LaurentMazare/ocaml-torch/commit/ef7ef30cafecb09e45ec1ed8ce4bedae5947cfa5.patch"; 37 + hash = "sha256-smdwKy40iIISp/25L2J4az6KmqFS1soeChBElUyhl5A="; 38 + }) 39 + ]; 31 40 32 41 buildInputs = [ dune-configurator ]; 33 42
+2
pkgs/development/python-modules/bottle/default.nix
··· 27 27 "test_delete_cookie" 28 28 "test_error" 29 29 "test_error_in_generator_callback" 30 + # timing sensitive 31 + "test_ims" 30 32 ]; 31 33 32 34 __darwinAllowLocalNetworking = true;
+253
pkgs/development/python-modules/openai-triton/default.nix
··· 1 + { lib 2 + , buildPythonPackage 3 + , python 4 + , fetchpatch 5 + , fetchFromGitHub 6 + , addOpenGLRunpath 7 + , cmake 8 + , cudaPackages 9 + , llvmPackages 10 + , pybind11 11 + , gtest 12 + , zlib 13 + , ncurses 14 + , libxml2 15 + , lit 16 + , filelock 17 + , torchWithRocm 18 + , pytest 19 + , pytestCheckHook 20 + , pythonRelaxDepsHook 21 + , pkgsTargetTarget 22 + }: 23 + 24 + let 25 + pname = "triton"; 26 + version = "2.0.0"; 27 + 28 + inherit (cudaPackages) cuda_cudart backendStdenv; 29 + 30 + # A time may come we'll want to be cross-friendly 31 + # 32 + # Short explanation: we need pkgsTargetTarget, because we use string 33 + # interpolation instead of buildInputs. 34 + # 35 + # Long explanation: OpenAI/triton downloads and vendors a copy of NVidia's 36 + # ptxas compiler. We're not running this ptxas on the build machine, but on 37 + # the user's machine, i.e. our Target platform. The second "Target" in 38 + # pkgsTargetTarget maybe doesn't matter, because ptxas compiles programs to 39 + # be executed on the GPU. 40 + # Cf. https://nixos.org/manual/nixpkgs/unstable/#sec-cross-infra 41 + ptxas = "${pkgsTargetTarget.cudaPackages.cuda_nvcc}/bin/ptxas"; 42 + 43 + llvm = (llvmPackages.llvm.override { 44 + llvmTargetsToBuild = [ "NATIVE" "NVPTX" ]; 45 + # Upstream CI sets these too: 46 + # targetProjects = [ "mlir" ]; 47 + extraCMakeFlags = [ 48 + "-DLLVM_INSTALL_UTILS=ON" 49 + ]; 50 + }); 51 + in 52 + buildPythonPackage { 53 + inherit pname version; 54 + 55 + format = "setuptools"; 56 + 57 + src = fetchFromGitHub { 58 + owner = "openai"; 59 + repo = pname; 60 + rev = "v${version}"; 61 + hash = "sha256-9GZzugab+Pdt74Dj6zjlEzjj4BcJ69rzMJmqcVMxsKU="; 62 + }; 63 + 64 + patches = [ 65 + # Prerequisite for llvm15 patch 66 + (fetchpatch { 67 + url = "https://github.com/openai/triton/commit/2aba985daaa70234823ea8f1161da938477d3e02.patch"; 68 + hash = "sha256-LGv0+Ut2WYPC4Ksi4803Hwmhi3FyQOF9zElJc/JCobk="; 69 + }) 70 + (fetchpatch { 71 + url = "https://github.com/openai/triton/commit/e3941f9d09cdd31529ba4a41018cfc0096aafea6.patch"; 72 + hash = "sha256-A+Gor6qzFlGQhVVhiaaYOzqqx8yO2MdssnQS6TIfUWg="; 73 + }) 74 + 75 + # Source: https://github.com/openai/triton/commit/fc7a8e35819bda632bdcf1cf75fd9abe4d4e077a.patch 76 + # The original patch adds ptxas binary, so we include our own clean copy 77 + # Drop with the next update 78 + ./llvm15.patch 79 + 80 + # TODO: there have been commits upstream aimed at removing the "torch" 81 + # circular dependency, but the patches fail to apply on the release 82 + # revision. Keeping the link for future reference 83 + # Also cf. https://github.com/openai/triton/issues/1374 84 + 85 + # (fetchpatch { 86 + # url = "https://github.com/openai/triton/commit/fc7c0b0e437a191e421faa61494b2ff4870850f1.patch"; 87 + # hash = "sha256-f0shIqHJkVvuil2Yku7vuqWFn7VCRKFSFjYRlwx25ig="; 88 + # }) 89 + ]; 90 + 91 + postPatch = '' 92 + substituteInPlace python/setup.py \ 93 + --replace \ 94 + '= get_thirdparty_packages(triton_cache_path)' \ 95 + '= os.environ["cmakeFlags"].split()' 96 + '' 97 + # Wiring triton=2.0.0 with llcmPackages_rocm.llvm=5.4.3 98 + # Revisit when updating either triton or llvm 99 + + '' 100 + substituteInPlace CMakeLists.txt \ 101 + --replace "nvptx" "NVPTX" \ 102 + --replace "LLVM 11" "LLVM" 103 + sed -i '/AddMLIR/a set(MLIR_TABLEGEN_EXE "${llvmPackages.mlir}/bin/mlir-tblgen")' CMakeLists.txt 104 + sed -i '/AddMLIR/a set(MLIR_INCLUDE_DIR ''${MLIR_INCLUDE_DIRS})' CMakeLists.txt 105 + find -iname '*.td' -exec \ 106 + sed -i \ 107 + -e '\|include "mlir/IR/OpBase.td"|a include "mlir/IR/AttrTypeBase.td"' \ 108 + -e 's|include "mlir/Dialect/StandardOps/IR/Ops.td"|include "mlir/Dialect/Func/IR/FuncOps.td"|' \ 109 + '{}' ';' 110 + substituteInPlace unittest/CMakeLists.txt --replace "include(GoogleTest)" "find_package(GTest REQUIRED)" 111 + sed -i 's/^include.*$//' unittest/CMakeLists.txt 112 + sed -i '/LINK_LIBS/i NVPTXInfo' lib/Target/PTX/CMakeLists.txt 113 + sed -i '/LINK_LIBS/i NVPTXCodeGen' lib/Target/PTX/CMakeLists.txt 114 + '' 115 + # TritonMLIRIR already links MLIRIR. Not transitive? 116 + # + '' 117 + # echo "target_link_libraries(TritonPTX PUBLIC MLIRIR)" >> lib/Target/PTX/CMakeLists.txt 118 + # '' 119 + # Already defined in llvm, when built with -DLLVM_INSTALL_UTILS 120 + + '' 121 + substituteInPlace bin/CMakeLists.txt \ 122 + --replace "add_subdirectory(FileCheck)" "" 123 + 124 + rm cmake/FindLLVM.cmake 125 + '' 126 + + 127 + ( 128 + let 129 + # Bash was getting weird without linting, 130 + # but basically upstream contains [cc, ..., "-lcuda", ...] 131 + # and we replace it with [..., "-lcuda", "-L/run/opengl-driver/lib", "-L$stubs", ...] 132 + old = [ "-lcuda" ]; 133 + new = [ "-lcuda" "-L${addOpenGLRunpath.driverLink}" "-L${cuda_cudart}/lib/stubs/" ]; 134 + 135 + quote = x: ''"${x}"''; 136 + oldStr = lib.concatMapStringsSep ", " quote old; 137 + newStr = lib.concatMapStringsSep ", " quote new; 138 + in 139 + '' 140 + substituteInPlace python/triton/compiler.py \ 141 + --replace '${oldStr}' '${newStr}' 142 + '' 143 + ) 144 + # Triton seems to be looking up cuda.h 145 + + '' 146 + sed -i 's|cu_include_dir = os.path.join.*$|cu_include_dir = "${cuda_cudart}/include"|' python/triton/compiler.py 147 + ''; 148 + 149 + nativeBuildInputs = [ 150 + cmake 151 + pythonRelaxDepsHook 152 + 153 + # Requires torch (circular dependency) and probably needs GPUs: 154 + # pytestCheckHook 155 + 156 + # Note for future: 157 + # These *probably* should go in depsTargetTarget 158 + # ...but we cannot test cross right now anyway 159 + # because we only support cudaPackages on x86_64-linux atm 160 + lit 161 + llvm 162 + llvmPackages.mlir 163 + ]; 164 + 165 + buildInputs = [ 166 + gtest 167 + libxml2.dev 168 + ncurses 169 + pybind11 170 + zlib 171 + ]; 172 + 173 + propagatedBuildInputs = [ 174 + filelock 175 + ]; 176 + 177 + # Avoid GLIBCXX mismatch with other cuda-enabled python packages 178 + preConfigure = '' 179 + export CC="${backendStdenv.cc}/bin/cc"; 180 + export CXX="${backendStdenv.cc}/bin/c++"; 181 + 182 + # Upstream's setup.py tries to write cache somewhere in ~/ 183 + export HOME=$TMPDIR 184 + 185 + # Upstream's github actions patch setup.cfg to write base-dir. May be redundant 186 + echo " 187 + [build_ext] 188 + base-dir=$PWD" >> python/setup.cfg 189 + 190 + # The rest (including buildPhase) is relative to ./python/ 191 + cd python/ 192 + 193 + # Work around download_and_copy_ptxas() 194 + dst_cuda="$PWD/triton/third_party/cuda/bin" 195 + mkdir -p "$dst_cuda" 196 + ln -s "${ptxas}" "$dst_cuda/" 197 + ''; 198 + 199 + # CMake is run by setup.py instead 200 + dontUseCmakeConfigure = true; 201 + cmakeFlags = [ 202 + "-DMLIR_DIR=${llvmPackages.mlir}/lib/cmake/mlir" 203 + ]; 204 + 205 + postFixup = 206 + let 207 + ptxasDestination = "$out/${python.sitePackages}/triton/third_party/cuda/bin/ptxas"; 208 + in 209 + # Setuptools (?) strips runpath and +x flags. Let's just restore the symlink 210 + '' 211 + rm -f ${ptxasDestination} 212 + ln -s ${ptxas} ${ptxasDestination} 213 + ''; 214 + 215 + checkInputs = [ 216 + cmake # ctest 217 + ]; 218 + dontUseSetuptoolsCheck = true; 219 + preCheck = 220 + # build/temp* refers to build_ext.build_temp (looked up in the build logs) 221 + '' 222 + (cd /build/source/python/build/temp* ; ctest) 223 + '' # For pytestCheckHook 224 + + '' 225 + cd test/unit 226 + ''; 227 + pythonImportsCheck = [ 228 + # Circular dependency on torch 229 + # "triton" 230 + # "triton.language" 231 + ]; 232 + 233 + # Ultimately, torch is our test suite: 234 + passthru.tests = { 235 + inherit torchWithRocm; 236 + }; 237 + 238 + pythonRemoveDeps = [ 239 + # Circular dependency, cf. https://github.com/openai/triton/issues/1374 240 + "torch" 241 + 242 + # CLI tools without dist-info 243 + "cmake" 244 + "lit" 245 + ]; 246 + meta = with lib; { 247 + description = "Development repository for the Triton language and compiler"; 248 + homepage = "https://github.com/openai/triton/"; 249 + platforms = lib.platforms.unix; 250 + license = licenses.mit; 251 + maintainers = with maintainers; [ SomeoneSerge ]; 252 + }; 253 + }
+4617
pkgs/development/python-modules/openai-triton/llvm15.patch
··· 1 + From fc7a8e35819bda632bdcf1cf75fd9abe4d4e077a Mon Sep 17 00:00:00 2001 2 + From: Christian Sigg <chsigg@users.noreply.github.com> 3 + Date: Thu, 16 Feb 2023 15:40:53 +0100 4 + Subject: [PATCH] Rebase Triton to LLVM-15. (#1070) 5 + 6 + This PR rebases Triton from LLVM-14 to LLVM-15. Most changes are 7 + mechanical, except for the analysis framework changes. 8 + --- 9 + CMakeLists.txt | 6 +- 10 + bin/CMakeLists.txt | 2 +- 11 + bin/FileCheck/FileCheck.cpp | 3 + 12 + bin/triton-opt.cpp | 6 +- 13 + bin/triton-translate.cpp | 7 +- 14 + include/triton/Analysis/Alias.h | 21 +- 15 + include/triton/Analysis/Allocation.h | 2 + 16 + include/triton/Analysis/AxisInfo.h | 56 ++- 17 + include/triton/Analysis/Utility.h | 6 +- 18 + include/triton/Conversion/Passes.td | 4 +- 19 + include/triton/Dialect/Triton/IR/Dialect.h | 7 +- 20 + .../triton/Dialect/Triton/IR/TritonDialect.td | 8 +- 21 + include/triton/Dialect/Triton/IR/TritonOps.td | 12 +- 22 + .../triton/Dialect/Triton/IR/TritonTypes.td | 2 + 23 + .../Dialect/Triton/Transforms/Passes.td | 3 +- 24 + include/triton/Dialect/TritonGPU/IR/Dialect.h | 4 +- 25 + .../Dialect/TritonGPU/IR/TritonGPUAttrDefs.td | 7 + 26 + .../Dialect/TritonGPU/IR/TritonGPUDialect.td | 2 +- 27 + .../Dialect/TritonGPU/IR/TritonGPUOps.td | 13 +- 28 + lib/Analysis/Alias.cpp | 14 +- 29 + lib/Analysis/Allocation.cpp | 30 +- 30 + lib/Analysis/AxisInfo.cpp | 79 ++-- 31 + lib/Analysis/CMakeLists.txt | 2 +- 32 + lib/Analysis/Membar.cpp | 2 +- 33 + lib/Analysis/Utility.cpp | 54 +++ 34 + .../TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp | 3 - 35 + lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h | 10 +- 36 + .../TritonGPUToLLVM/DotOpToLLVM.cpp | 5 - 37 + .../TritonGPUToLLVM/ElementwiseOpToLLVM.cpp | 2 - 38 + .../TritonGPUToLLVM/LoadStoreOpToLLVM.cpp | 5 +- 39 + .../TritonGPUToLLVM/ReduceOpToLLVM.cpp | 2 - 40 + .../TritonGPUToLLVM/TritonGPUToLLVM.cpp | 7 +- 41 + .../TritonGPUToLLVM/TritonGPUToLLVMBase.h | 26 +- 42 + .../TritonGPUToLLVM/TritonGPUToLLVMPass.cpp | 52 +-- 43 + lib/Conversion/TritonGPUToLLVM/Utility.h | 5 +- 44 + .../TritonToTritonGPUPass.cpp | 69 ++-- 45 + lib/Dialect/Triton/IR/CMakeLists.txt | 10 +- 46 + lib/Dialect/Triton/IR/Ops.cpp | 34 +- 47 + lib/Dialect/Triton/Transforms/Combine.cpp | 6 +- 48 + lib/Dialect/Triton/Transforms/Combine.td | 2 +- 49 + lib/Dialect/TritonGPU/IR/Dialect.cpp | 27 +- 50 + lib/Dialect/TritonGPU/Transforms/Coalesce.cpp | 20 +- 51 + lib/Dialect/TritonGPU/Transforms/Combine.cpp | 2 +- 52 + lib/Dialect/TritonGPU/Transforms/Combine.td | 1 + 53 + .../Transforms/DecomposeConversions.cpp | 2 +- 54 + lib/Dialect/TritonGPU/Transforms/Pipeline.cpp | 10 +- 55 + .../Transforms/ReorderInstructions.cpp | 2 +- 56 + .../Transforms/TritonGPUConversion.cpp | 12 +- 57 + .../Transforms/UpdateMmaForVolta.cpp | 6 +- 58 + lib/Dialect/TritonGPU/Transforms/Utility.cpp | 2 +- 59 + lib/Target/LLVMIR/CMakeLists.txt | 3 +- 60 + lib/Target/PTX/PTXTranslation.cpp | 3 + 61 + python/setup.py | 15 +- 62 + python/src/triton.cc | 85 +++-- 63 + python/test/unit/language/test_core.py | 2 +- 64 + python/triton/compiler.py | 4 +- 65 + test/Analysis/test-alias.mlir | 24 +- 66 + test/Analysis/test-alignment.mlir | 344 +++++++++--------- 67 + test/Analysis/test-allocation.mlir | 32 +- 68 + test/Analysis/test-membar.mlir | 38 +- 69 + test/Conversion/triton_ops.mlir | 10 +- 70 + test/Conversion/triton_to_tritongpu.mlir | 6 +- 71 + test/Conversion/tritongpu_to_llvm.mlir | 94 ++--- 72 + test/Target/tritongpu_to_llvmir.mlir | 4 +- 73 + test/Target/tritongpu_to_ptx.mlir | 2 +- 74 + test/Triton/combine.mlir | 40 +- 75 + test/Triton/vecadd.mlir | 4 +- 76 + test/TritonGPU/coalesce.mlir | 2 +- 77 + test/TritonGPU/combine.mlir | 38 +- 78 + test/TritonGPU/loop-pipeline.mlir | 22 +- 79 + test/TritonGPU/matmul.mlir | 4 +- 80 + test/TritonGPU/prefetch.mlir | 4 +- 81 + test/TritonGPU/update-mma-for-volta.mlir | 4 +- 82 + test/lib/Analysis/TestAlias.cpp | 29 +- 83 + test/lib/Analysis/TestAllocation.cpp | 5 +- 84 + test/lib/Analysis/TestAxisInfo.cpp | 51 +-- 85 + test/lib/Analysis/TestMembar.cpp | 7 +- 86 + 78 files changed, 808 insertions(+), 742 deletions(-) 87 + 88 + diff --git a/CMakeLists.txt b/CMakeLists.txt 89 + index d0d361fc7c..b281a28400 100644 90 + --- a/CMakeLists.txt 91 + +++ b/CMakeLists.txt 92 + @@ -1,4 +1,7 @@ 93 + cmake_minimum_required(VERSION 3.6) 94 + + 95 + +cmake_policy(SET CMP0116 OLD) 96 + + 97 + include(ExternalProject) 98 + 99 + set(CMAKE_CXX_STANDARD 17) 100 + @@ -155,7 +158,6 @@ if(TRITON_BUILD_PYTHON_MODULE) 101 + endif() 102 + endif() 103 + 104 + - 105 + # # Triton 106 + # file(GLOB_RECURSE LIBTRITON_SRC lib/*.cc) 107 + # if (WIN32 AND TRITON_BUILD_PYTHON_MODULE) 108 + @@ -212,7 +214,7 @@ if(TRITON_BUILD_PYTHON_MODULE) 109 + # optimizations 110 + MLIRPass 111 + MLIRTransforms 112 + - MLIRLLVMIR 113 + + MLIRLLVMDialect 114 + MLIRSupport 115 + MLIRTargetLLVMIRExport 116 + MLIRExecutionEngine 117 + diff --git a/bin/CMakeLists.txt b/bin/CMakeLists.txt 118 + index 906f635f8b..695b3479fd 100644 119 + --- a/bin/CMakeLists.txt 120 + +++ b/bin/CMakeLists.txt 121 + @@ -48,7 +48,7 @@ llvm_update_compile_flags(triton-translate) 122 + # MLIR core 123 + MLIROptLib 124 + MLIRIR 125 + - MLIRLLVMIR 126 + + MLIRLLVMDialect 127 + MLIRPass 128 + MLIRSupport 129 + MLIRTransforms 130 + diff --git a/bin/FileCheck/FileCheck.cpp b/bin/FileCheck/FileCheck.cpp 131 + index 819efc3541..9ac6f1b277 100644 132 + --- a/bin/FileCheck/FileCheck.cpp 133 + +++ b/bin/FileCheck/FileCheck.cpp 134 + @@ -19,6 +19,7 @@ 135 + #include "llvm/Support/CommandLine.h" 136 + #include "llvm/Support/InitLLVM.h" 137 + #include "llvm/Support/Process.h" 138 + +#include "llvm/Support/SourceMgr.h" 139 + #include "llvm/Support/WithColor.h" 140 + #include "llvm/Support/raw_ostream.h" 141 + #include <cmath> 142 + @@ -360,6 +361,8 @@ static std::string GetCheckTypeAbbreviation(Check::FileCheckType Ty) { 143 + return "bad-not"; 144 + case Check::CheckBadCount: 145 + return "bad-count"; 146 + + case Check::CheckMisspelled: 147 + + return "misspelled"; 148 + case Check::CheckNone: 149 + llvm_unreachable("invalid FileCheckType"); 150 + } 151 + diff --git a/bin/triton-opt.cpp b/bin/triton-opt.cpp 152 + index 9f3b53b7ae..f96232e1b0 100644 153 + --- a/bin/triton-opt.cpp 154 + +++ b/bin/triton-opt.cpp 155 + @@ -8,7 +8,7 @@ 156 + 157 + #include "mlir/IR/Dialect.h" 158 + #include "mlir/InitAllPasses.h" 159 + -#include "mlir/Support/MlirOptMain.h" 160 + +#include "mlir/Tools/mlir-opt/MlirOptMain.h" 161 + 162 + namespace mlir { 163 + namespace test { 164 + @@ -33,8 +33,8 @@ int main(int argc, char **argv) { 165 + // TODO: register Triton & TritonGPU passes 166 + mlir::DialectRegistry registry; 167 + registry.insert<mlir::triton::TritonDialect, 168 + - mlir::triton::gpu::TritonGPUDialect, mlir::math::MathDialect, 169 + - mlir::arith::ArithmeticDialect, mlir::StandardOpsDialect, 170 + + mlir::triton::gpu::TritonGPUDialect, mlir::func::FuncDialect, 171 + + mlir::math::MathDialect, mlir::arith::ArithmeticDialect, 172 + mlir::scf::SCFDialect, mlir::gpu::GPUDialect>(); 173 + 174 + return mlir::asMainReturnCode(mlir::MlirOptMain( 175 + diff --git a/bin/triton-translate.cpp b/bin/triton-translate.cpp 176 + index 05ba15e453..56b5d65857 100644 177 + --- a/bin/triton-translate.cpp 178 + +++ b/bin/triton-translate.cpp 179 + @@ -3,7 +3,7 @@ 180 + #include "mlir/IR/AsmState.h" 181 + #include "mlir/IR/BuiltinOps.h" 182 + #include "mlir/IR/Dialect.h" 183 + -#include "mlir/Parser.h" 184 + +#include "mlir/Parser/Parser.h" 185 + #include "mlir/Pass/Pass.h" 186 + #include "mlir/Pass/PassManager.h" 187 + #include "mlir/Support/FileUtilities.h" 188 + @@ -38,7 +38,7 @@ OwningOpRef<ModuleOp> loadMLIRModule(llvm::StringRef inputFilename, 189 + mlir::DialectRegistry registry; 190 + registry.insert<TritonDialect, triton::gpu::TritonGPUDialect, 191 + mlir::math::MathDialect, arith::ArithmeticDialect, 192 + - StandardOpsDialect, scf::SCFDialect>(); 193 + + scf::SCFDialect>(); 194 + 195 + context.appendDialectRegistry(registry); 196 + 197 + @@ -50,7 +50,8 @@ OwningOpRef<ModuleOp> loadMLIRModule(llvm::StringRef inputFilename, 198 + context.loadAllAvailableDialects(); 199 + context.allowUnregisteredDialects(); 200 + 201 + - OwningOpRef<ModuleOp> module(parseSourceFile(sourceMgr, &context)); 202 + + OwningOpRef<ModuleOp> module = 203 + + parseSourceFile<ModuleOp>(sourceMgr, &context); 204 + if (!module) { 205 + llvm::errs() << "Parse MLIR file failed."; 206 + return nullptr; 207 + diff --git a/include/triton/Analysis/Alias.h b/include/triton/Analysis/Alias.h 208 + index fa6b906fc9..631df518bc 100644 209 + --- a/include/triton/Analysis/Alias.h 210 + +++ b/include/triton/Analysis/Alias.h 211 + @@ -2,7 +2,7 @@ 212 + #define TRITON_ANALYSIS_ALIAS_H 213 + 214 + #include "mlir/Analysis/AliasAnalysis.h" 215 + -#include "mlir/Analysis/DataFlowAnalysis.h" 216 + +#include "mlir/Analysis/DataFlow/SparseAnalysis.h" 217 + #include "llvm/ADT/DenseSet.h" 218 + 219 + namespace mlir { 220 + @@ -21,7 +21,7 @@ class AliasInfo { 221 + } 222 + 223 + /// The pessimistic value state of a value without alias 224 + - static AliasInfo getPessimisticValueState(MLIRContext *context) { 225 + + static AliasInfo getPessimisticValueState(MLIRContext *context = nullptr) { 226 + return AliasInfo(); 227 + } 228 + static AliasInfo getPessimisticValueState(Value value) { return AliasInfo(); } 229 + @@ -29,6 +29,10 @@ class AliasInfo { 230 + /// The union of both arguments 231 + static AliasInfo join(const AliasInfo &lhs, const AliasInfo &rhs); 232 + 233 + + void print(raw_ostream &os) const { 234 + + llvm::interleaveComma(allocs, os, [&](Value alloc) { alloc.print(os); }); 235 + + } 236 + + 237 + private: 238 + /// The set of allocated values that are aliased by this lattice. 239 + /// For now, we only consider aliased value produced by the following 240 + @@ -58,9 +62,13 @@ class AliasInfo { 241 + //===----------------------------------------------------------------------===// 242 + // Shared Memory Alias Analysis 243 + //===----------------------------------------------------------------------===// 244 + -class SharedMemoryAliasAnalysis : public ForwardDataFlowAnalysis<AliasInfo> { 245 + +class SharedMemoryAliasAnalysis 246 + + : public dataflow::SparseDataFlowAnalysis<dataflow::Lattice<AliasInfo>> { 247 + public: 248 + - using ForwardDataFlowAnalysis<AliasInfo>::ForwardDataFlowAnalysis; 249 + + using dataflow::SparseDataFlowAnalysis< 250 + + dataflow::Lattice<AliasInfo>>::SparseDataFlowAnalysis; 251 + + using dataflow::SparseDataFlowAnalysis< 252 + + dataflow::Lattice<AliasInfo>>::getLatticeElement; 253 + 254 + /// XXX(Keren): Compatible interface with MLIR AliasAnalysis for future use. 255 + /// Given two values, returns their aliasing behavior. 256 + @@ -70,9 +78,10 @@ class SharedMemoryAliasAnalysis : public ForwardDataFlowAnalysis<AliasInfo> { 257 + ModRefResult getModRef(Operation *op, Value location); 258 + 259 + /// Computes if the alloc set of the results are changed. 260 + - ChangeResult 261 + + void 262 + visitOperation(Operation *op, 263 + - ArrayRef<LatticeElement<AliasInfo> *> operands) override; 264 + + ArrayRef<const dataflow::Lattice<AliasInfo> *> operands, 265 + + ArrayRef<dataflow::Lattice<AliasInfo> *> results) override; 266 + }; 267 + 268 + } // namespace mlir 269 + diff --git a/include/triton/Analysis/Allocation.h b/include/triton/Analysis/Allocation.h 270 + index b7c136d602..89b77034cc 100644 271 + --- a/include/triton/Analysis/Allocation.h 272 + +++ b/include/triton/Analysis/Allocation.h 273 + @@ -188,6 +188,8 @@ class Allocation { 274 + friend class triton::AllocationAnalysis; 275 + }; 276 + 277 + +template <typename T> Interval(T, T) -> Interval<T>; 278 + + 279 + } // namespace mlir 280 + 281 + #endif // TRITON_ANALYSIS_ALLOCATION_H 282 + diff --git a/include/triton/Analysis/AxisInfo.h b/include/triton/Analysis/AxisInfo.h 283 + index fdfbd8fbb3..7083b9c43b 100644 284 + --- a/include/triton/Analysis/AxisInfo.h 285 + +++ b/include/triton/Analysis/AxisInfo.h 286 + @@ -1,9 +1,10 @@ 287 + #ifndef TRITON_ANALYSIS_AXISINFO_H 288 + #define TRITON_ANALYSIS_AXISINFO_H 289 + 290 + -#include "mlir/Analysis/DataFlowAnalysis.h" 291 + +#include "mlir/Analysis/DataFlow/SparseAnalysis.h" 292 + #include "llvm/Support/raw_ostream.h" 293 + 294 + +#include "mlir/Support/LLVM.h" 295 + #include "triton/Analysis/Utility.h" 296 + #include "triton/Dialect/Triton/IR/Dialect.h" 297 + #include "triton/Dialect/TritonGPU/IR/Dialect.h" 298 + @@ -62,7 +63,7 @@ class AxisInfo { 299 + } 300 + 301 + /// The pessimistic value state of the contiguity is unknown. 302 + - static AxisInfo getPessimisticValueState(MLIRContext *context) { 303 + + static AxisInfo getPessimisticValueState(MLIRContext *context = nullptr) { 304 + return AxisInfo(); 305 + } 306 + static AxisInfo getPessimisticValueState(Value value); 307 + @@ -70,6 +71,22 @@ class AxisInfo { 308 + /// The gcd of both arguments for each dimension 309 + static AxisInfo join(const AxisInfo &lhs, const AxisInfo &rhs); 310 + 311 + + void print(raw_ostream &os) const { 312 + + auto print = [&](StringRef name, DimVectorT vec) { 313 + + os << name << " = ["; 314 + + llvm::interleaveComma(vec, os); 315 + + os << "]"; 316 + + }; 317 + + print("contiguity", contiguity); 318 + + print(", divisibility", divisibility); 319 + + print(", constancy", constancy); 320 + + os << ", constant_value = "; 321 + + if (constantValue) 322 + + os << *constantValue; 323 + + else 324 + + os << "<none>"; 325 + + } 326 + + 327 + private: 328 + /// The _contiguity_ information maps the `d`-th 329 + /// dimension to the length of the shortest 330 + @@ -147,7 +164,8 @@ class AxisInfoVisitor { 331 + } 332 + 333 + virtual AxisInfo 334 + - getAxisInfo(Operation *op, ArrayRef<LatticeElement<AxisInfo> *> operands) = 0; 335 + + getAxisInfo(Operation *op, 336 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) = 0; 337 + 338 + virtual bool match(Operation *op) = 0; 339 + }; 340 + @@ -157,15 +175,16 @@ template <typename OpTy> class AxisInfoVisitorImpl : public AxisInfoVisitor { 341 + public: 342 + using AxisInfoVisitor::AxisInfoVisitor; 343 + 344 + - AxisInfo getAxisInfo(Operation *op, 345 + - ArrayRef<LatticeElement<AxisInfo> *> operands) final { 346 + + AxisInfo 347 + + getAxisInfo(Operation *op, 348 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) final { 349 + return getAxisInfo(cast<OpTy>(op), operands); 350 + } 351 + 352 + bool match(Operation *op) final { return isa<OpTy>(op); } 353 + 354 + - virtual AxisInfo getAxisInfo(OpTy op, 355 + - ArrayRef<LatticeElement<AxisInfo> *> operands) { 356 + + virtual AxisInfo 357 + + getAxisInfo(OpTy op, ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) { 358 + llvm_unreachable("Unimplemented getAxisInfo"); 359 + } 360 + }; 361 + @@ -176,8 +195,9 @@ class BinaryOpVisitorImpl : public AxisInfoVisitorImpl<OpTy> { 362 + public: 363 + using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl; 364 + 365 + - AxisInfo getAxisInfo(OpTy op, 366 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 367 + + AxisInfo 368 + + getAxisInfo(OpTy op, 369 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 370 + auto lhsInfo = operands[0]->getValue(); 371 + auto rhsInfo = operands[1]->getValue(); 372 + auto rank = lhsInfo.getRank(); 373 + @@ -230,7 +250,8 @@ class AxisInfoVisitorList { 374 + (visitors.emplace_back(std::make_unique<Ts>()), ...); 375 + } 376 + 377 + - AxisInfo apply(Operation *op, ArrayRef<LatticeElement<AxisInfo> *> operands) { 378 + + AxisInfo apply(Operation *op, 379 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) { 380 + for (auto &visitor : visitors) 381 + if (visitor->match(op)) 382 + return visitor->getAxisInfo(op, operands); 383 + @@ -241,16 +262,19 @@ class AxisInfoVisitorList { 384 + std::vector<std::unique_ptr<AxisInfoVisitor>> visitors; 385 + }; 386 + 387 + -class AxisInfoAnalysis : public ForwardDataFlowAnalysis<AxisInfo> { 388 + +class AxisInfoAnalysis 389 + + : public dataflow::SparseDataFlowAnalysis<dataflow::Lattice<AxisInfo>> { 390 + private: 391 + AxisInfoVisitorList visitors; 392 + 393 + public: 394 + - AxisInfoAnalysis(MLIRContext *context); 395 + + AxisInfoAnalysis(DataFlowSolver &solver); 396 + + using dataflow::SparseDataFlowAnalysis< 397 + + dataflow::Lattice<AxisInfo>>::getLatticeElement; 398 + 399 + - ChangeResult 400 + - visitOperation(Operation *op, 401 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override; 402 + + void visitOperation(Operation *op, 403 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands, 404 + + ArrayRef<dataflow::Lattice<AxisInfo> *> results) override; 405 + 406 + unsigned getPtrContiguity(Value ptr); 407 + 408 + @@ -261,4 +285,4 @@ class AxisInfoAnalysis : public ForwardDataFlowAnalysis<AxisInfo> { 409 + 410 + } // namespace mlir 411 + 412 + -#endif 413 + \ No newline at end of file 414 + +#endif 415 + diff --git a/include/triton/Analysis/Utility.h b/include/triton/Analysis/Utility.h 416 + index c5ac137dc1..ee7fadb59d 100644 417 + --- a/include/triton/Analysis/Utility.h 418 + +++ b/include/triton/Analysis/Utility.h 419 + @@ -1,6 +1,7 @@ 420 + #ifndef TRITON_ANALYSIS_UTILITY_H 421 + #define TRITON_ANALYSIS_UTILITY_H 422 + 423 + +#include "mlir/Analysis/DataFlowFramework.h" 424 + #include "mlir/Analysis/SliceAnalysis.h" 425 + #include "triton/Dialect/TritonGPU/IR/Dialect.h" 426 + #include <algorithm> 427 + @@ -12,7 +13,7 @@ namespace mlir { 428 + class ReduceOpHelper { 429 + public: 430 + explicit ReduceOpHelper(triton::ReduceOp op) : op(op) { 431 + - srcTy = op.operand().getType().cast<RankedTensorType>(); 432 + + srcTy = op.getOperand().getType().cast<RankedTensorType>(); 433 + } 434 + 435 + ArrayRef<int64_t> getSrcShape() { return srcTy.getShape(); } 436 + @@ -103,6 +104,9 @@ SetVector<Operation *> 437 + multiRootGetSlice(Operation *op, TransitiveFilter backwardFilter = nullptr, 438 + TransitiveFilter forwardFilter = nullptr); 439 + 440 + +// Create a basic DataFlowSolver with constant and dead code analysis included. 441 + +std::unique_ptr<DataFlowSolver> createDataFlowSolver(); 442 + + 443 + } // namespace mlir 444 + 445 + #endif // TRITON_ANALYSIS_UTILITY_H 446 + diff --git a/include/triton/Conversion/Passes.td b/include/triton/Conversion/Passes.td 447 + index 70bb20b78e..be00eb2dac 100644 448 + --- a/include/triton/Conversion/Passes.td 449 + +++ b/include/triton/Conversion/Passes.td 450 + @@ -12,7 +12,6 @@ def ConvertTritonToTritonGPU: Pass<"convert-triton-to-tritongpu", "mlir::ModuleO 451 + 452 + let dependentDialects = ["mlir::arith::ArithmeticDialect", 453 + "mlir::math::MathDialect", 454 + - "mlir::StandardOpsDialect", 455 + // TODO: Does this pass depend on SCF? 456 + "mlir::scf::SCFDialect", 457 + "mlir::triton::TritonDialect", 458 + @@ -41,8 +40,7 @@ def ConvertTritonGPUToLLVM : Pass<"convert-triton-gpu-to-llvm", "mlir::ModuleOp" 459 + "mlir::tensor::TensorDialect", 460 + "mlir::triton::TritonDialect", 461 + "mlir::triton::gpu::TritonGPUDialect", 462 + - "mlir::NVVM::NVVMDialect", 463 + - "mlir::StandardOpsDialect"]; 464 + + "mlir::NVVM::NVVMDialect"]; 465 + 466 + let options = [ 467 + Option<"computeCapability", "compute-capability", 468 + diff --git a/include/triton/Dialect/Triton/IR/Dialect.h b/include/triton/Dialect/Triton/IR/Dialect.h 469 + index e8012a51df..15869e262e 100644 470 + --- a/include/triton/Dialect/Triton/IR/Dialect.h 471 + +++ b/include/triton/Dialect/Triton/IR/Dialect.h 472 + @@ -1,14 +1,15 @@ 473 + #ifndef TRITON_DIALECT_TRITON_IR_DIALECT_H_ 474 + #define TRITON_DIALECT_TRITON_IR_DIALECT_H_ 475 + 476 + +#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" 477 + +#include "mlir/Dialect/ControlFlow/IR/ControlFlow.h" 478 + +#include "mlir/Dialect/Func/IR/FuncOps.h" 479 + #include "mlir/Dialect/Math/IR/Math.h" 480 + -#include "mlir/Dialect/SCF/SCF.h" 481 + -#include "mlir/Dialect/StandardOps/IR/Ops.h" 482 + +#include "mlir/Dialect/SCF/IR/SCF.h" 483 + #include "mlir/Dialect/Tensor/IR/Tensor.h" 484 + #include "mlir/IR/BuiltinOps.h" 485 + #include "mlir/IR/Dialect.h" 486 + #include "mlir/Interfaces/ControlFlowInterfaces.h" 487 + - 488 + #include "triton/Dialect/Triton/IR/Dialect.h.inc" 489 + #include "triton/Dialect/Triton/IR/OpsEnums.h.inc" 490 + #include "triton/Dialect/Triton/IR/Traits.h" 491 + diff --git a/include/triton/Dialect/Triton/IR/TritonDialect.td b/include/triton/Dialect/Triton/IR/TritonDialect.td 492 + index 07b069e14f..d98ce73884 100644 493 + --- a/include/triton/Dialect/Triton/IR/TritonDialect.td 494 + +++ b/include/triton/Dialect/Triton/IR/TritonDialect.td 495 + @@ -25,12 +25,9 @@ def Triton_Dialect : Dialect { 496 + let dependentDialects = [ 497 + "arith::ArithmeticDialect", 498 + "math::MathDialect", 499 + - "StandardOpsDialect", 500 + "scf::SCFDialect", 501 + - 502 + - // Since LLVM 15 503 + - // "cf::ControlFlowDialect", 504 + - // "func::FuncDialect" 505 + + "cf::ControlFlowDialect", 506 + + "func::FuncDialect" 507 + ]; 508 + 509 + let extraClassDeclaration = [{ 510 + @@ -38,6 +35,7 @@ def Triton_Dialect : Dialect { 511 + }]; 512 + 513 + let hasConstantMaterializer = 1; 514 + + let useDefaultTypePrinterParser = 1; 515 + } 516 + 517 + include "triton/Dialect/Triton/IR/TritonTypes.td" 518 + diff --git a/include/triton/Dialect/Triton/IR/TritonOps.td b/include/triton/Dialect/Triton/IR/TritonOps.td 519 + index 779e0b648c..0a69211179 100644 520 + --- a/include/triton/Dialect/Triton/IR/TritonOps.td 521 + +++ b/include/triton/Dialect/Triton/IR/TritonOps.td 522 + @@ -141,11 +141,7 @@ def TT_LoadOp : TT_Op<"load", 523 + "triton::EvictionPolicy":$evict, "bool":$isVolatile)>, 524 + ]; 525 + 526 + - // let assemblyFormat = "operands attr-dict `:` type($result)"; 527 + - let parser = [{ return mlir::triton::parseLoadOp(parser, result); }]; 528 + - 529 + - let printer = [{ return mlir::triton::printLoadOp(p, *this); }]; 530 + - 531 + + let hasCustomAssemblyFormat = 1; 532 + let hasCanonicalizer = 1; 533 + } 534 + 535 + @@ -170,11 +166,7 @@ def TT_StoreOp : TT_Op<"store", 536 + "triton::EvictionPolicy":$evict)>, 537 + ]; 538 + 539 + - // let assemblyFormat = "operands attr-dict `:` type($value)"; 540 + - let parser = [{ return mlir::triton::parseStoreOp(parser, result); }]; 541 + - 542 + - let printer = [{ return mlir::triton::printStoreOp(p, *this); }]; 543 + - 544 + + let hasCustomAssemblyFormat = 1; 545 + let hasCanonicalizer = 1; 546 + } 547 + 548 + diff --git a/include/triton/Dialect/Triton/IR/TritonTypes.td b/include/triton/Dialect/Triton/IR/TritonTypes.td 549 + index 66d2a7b9a9..2fe2fd077d 100644 550 + --- a/include/triton/Dialect/Triton/IR/TritonTypes.td 551 + +++ b/include/triton/Dialect/Triton/IR/TritonTypes.td 552 + @@ -1,6 +1,7 @@ 553 + #ifndef TRITON_TYPES 554 + #define TRITON_TYPES 555 + 556 + +include "mlir/IR/AttrTypeBase.td" 557 + include "triton/Dialect/Triton/IR/TritonDialect.td" 558 + 559 + // 560 + @@ -58,6 +59,7 @@ def TT_Ptr : TritonTypeDef<"Pointer", "ptr"> { 561 + }]> 562 + ]; 563 + 564 + + let hasCustomAssemblyFormat = 1; 565 + let skipDefaultBuilders = 1; 566 + } 567 + def TT_PtrTensor : TensorOf<[TT_Ptr]>; 568 + diff --git a/include/triton/Dialect/Triton/Transforms/Passes.td b/include/triton/Dialect/Triton/Transforms/Passes.td 569 + index 8f77aed774..a25cdc5680 100644 570 + --- a/include/triton/Dialect/Triton/Transforms/Passes.td 571 + +++ b/include/triton/Dialect/Triton/Transforms/Passes.td 572 + @@ -16,8 +16,7 @@ def TritonCombineOps : Pass</*cli-arg*/"triton-combine", /*Op*/"mlir::ModuleOp"> 573 + 574 + let constructor = "mlir::triton::createCombineOpsPass()"; 575 + 576 + - let dependentDialects = ["mlir::arith::ArithmeticDialect", 577 + - /*SelectOp*/"mlir::StandardOpsDialect"]; 578 + + let dependentDialects = ["mlir::arith::ArithmeticDialect"]; 579 + } 580 + 581 + #endif 582 + diff --git a/include/triton/Dialect/TritonGPU/IR/Dialect.h b/include/triton/Dialect/TritonGPU/IR/Dialect.h 583 + index b4c8daec7b..dfc5f53ab1 100644 584 + --- a/include/triton/Dialect/TritonGPU/IR/Dialect.h 585 + +++ b/include/triton/Dialect/TritonGPU/IR/Dialect.h 586 + @@ -1,19 +1,17 @@ 587 + #ifndef TRITON_DIALECT_TRITONGPU_IR_DIALECT_H_ 588 + #define TRITON_DIALECT_TRITONGPU_IR_DIALECT_H_ 589 + 590 + -#include "mlir/Dialect/GPU/GPUDialect.h" 591 + +#include "mlir/Dialect/GPU/IR/GPUDialect.h" 592 + #include "mlir/Dialect/Tensor/IR/Tensor.h" 593 + #include "mlir/IR/BuiltinOps.h" 594 + #include "mlir/IR/Dialect.h" 595 + 596 + // TritonGPU depends on Triton 597 + #include "triton/Dialect/Triton/IR/Dialect.h" 598 + - 599 + #include "triton/Dialect/TritonGPU/IR/Dialect.h.inc" 600 + #include "triton/Dialect/TritonGPU/IR/Traits.h" 601 + 602 + #define GET_ATTRDEF_CLASSES 603 + -#include "triton/Dialect/Triton/IR/AttrInterfaces.h.inc" 604 + #include "triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.h.inc" 605 + 606 + #define GET_OP_CLASSES 607 + diff --git a/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td b/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td 608 + index 0242c3cc17..af2aeb03a8 100644 609 + --- a/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td 610 + +++ b/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td 611 + @@ -1,6 +1,7 @@ 612 + #ifndef TRITONGPU_ATTRDEFS 613 + #define TRITONGPU_ATTRDEFS 614 + 615 + +include "mlir/IR/AttrTypeBase.td" 616 + include "triton/Dialect/TritonGPU/IR/TritonGPUDialect.td" 617 + include "triton/Dialect/Triton/IR/TritonInterfaces.td" 618 + 619 + @@ -136,6 +137,7 @@ A_{3, 2} A_{3, 3} A_{3, 0} A_{3, 1} ... [phase 1] / 620 + ]; 621 + 622 + let extraClassDeclaration = extraBaseClassDeclaration; 623 + + let hasCustomAssemblyFormat = 1; 624 + } 625 + 626 + //===----------------------------------------------------------------------===// 627 + @@ -273,6 +275,7 @@ for 628 + // ArrayRefParameter<"unsigned">:$sizePerCTA 629 + ); 630 + 631 + + let hasCustomAssemblyFormat = 1; 632 + } 633 + 634 + //===----------------------------------------------------------------------===// 635 + @@ -422,6 +425,7 @@ For example, the matrix L corresponding to blockTileSize=[32,16] is: 636 + static constexpr int numBitsToHoldMmaV1ID{5}; 637 + }]; 638 + 639 + + let hasCustomAssemblyFormat = 1; 640 + } 641 + 642 + def SliceEncodingAttr : DistributedEncoding<"SliceEncoding"> { 643 + @@ -456,6 +460,8 @@ def SliceEncodingAttr : DistributedEncoding<"SliceEncoding"> { 644 + template<class T> 645 + SmallVector<T> paddedShape(ArrayRef<T> shape) const; 646 + }]; 647 + + 648 + + let hasCustomAssemblyFormat = 1; 649 + } 650 + 651 + def DotOperandEncodingAttr : DistributedEncoding<"DotOperandEncoding"> { 652 + @@ -492,6 +498,7 @@ section 9.7.13.4.1 for more details. 653 + 654 + ]; 655 + 656 + + let hasCustomAssemblyFormat = 1; 657 + let extraClassDeclaration = extraBaseClassDeclaration; 658 + } 659 + 660 + diff --git a/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td b/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td 661 + index 87ec1d36c6..6489a721b4 100644 662 + --- a/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td 663 + +++ b/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td 664 + @@ -30,7 +30,7 @@ def TritonGPU_Dialect : Dialect { 665 + } 666 + }]; 667 + 668 + - 669 + + let useDefaultAttributePrinterParser = 1; 670 + } 671 + 672 + #endif 673 + diff --git a/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td b/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td 674 + index 510f8d0183..7aba11dc75 100644 675 + --- a/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td 676 + +++ b/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td 677 + @@ -59,7 +59,7 @@ def TTG_AsyncCommitGroupOp : TTG_Op<"async_commit_group"> { 678 + // This is needed because these ops don't 679 + // handle encodings 680 + // e.g., https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/Dialect/Arith/IR/ArithOps.td#L111 681 + -def TTG_CmpIOp : TTG_Op<"cmpi", [NoSideEffect, Elementwise, 682 + +def TTG_CmpIOp : TTG_Op<"cmpi", [NoSideEffect, Elementwise, 683 + SameOperandsAndResultShape, 684 + SameOperandsAndResultEncoding]> { 685 + let summary = "integer comparison operation"; 686 + @@ -73,7 +73,7 @@ def TTG_CmpIOp : TTG_Op<"cmpi", [NoSideEffect, Elementwise, 687 + let results = (outs TT_BoolLike:$result); 688 + } 689 + 690 + -def TTG_CmpFOp : TTG_Op<"cmpf", [NoSideEffect, Elementwise, 691 + +def TTG_CmpFOp : TTG_Op<"cmpf", [NoSideEffect, Elementwise, 692 + SameOperandsAndResultShape, 693 + SameOperandsAndResultEncoding]> { 694 + let summary = "floating-point comparison operation"; 695 + @@ -88,8 +88,8 @@ def TTG_CmpFOp : TTG_Op<"cmpf", [NoSideEffect, Elementwise, 696 + } 697 + 698 + // TODO: migrate to arith::SelectOp on LLVM16 699 + -def TTG_SelectOp : TTG_Op<"select", [NoSideEffect, Elementwise, 700 + - SameOperandsAndResultShape, 701 + +def TTG_SelectOp : TTG_Op<"select", [NoSideEffect, Elementwise, 702 + + SameOperandsAndResultShape, 703 + SameOperandsAndResultEncoding]> { 704 + let summary = "select operation"; 705 + 706 + @@ -188,10 +188,7 @@ def TTG_InsertSliceAsyncOp : TTG_Op<"insert_slice_async", 707 + } 708 + }]; 709 + 710 + - // The custom parser could be replaced with oilist in LLVM-16 711 + - let parser = [{ return parseInsertSliceAsyncOp(parser, result); }]; 712 + - 713 + - let printer = [{ return printInsertSliceAsyncOp(p, *this); }]; 714 + + let hasCustomAssemblyFormat = 1; 715 + } 716 + 717 + def TTG_AllocTensorOp : TTG_Op<"alloc_tensor", [MemoryEffects<[MemAlloc]>, // Allocate shared memory 718 + diff --git a/lib/Analysis/Alias.cpp b/lib/Analysis/Alias.cpp 719 + index a39e4de9aa..208fdd4afc 100644 720 + --- a/lib/Analysis/Alias.cpp 721 + +++ b/lib/Analysis/Alias.cpp 722 + @@ -18,8 +18,9 @@ AliasInfo AliasInfo::join(const AliasInfo &lhs, const AliasInfo &rhs) { 723 + return ret; 724 + } 725 + 726 + -ChangeResult SharedMemoryAliasAnalysis::visitOperation( 727 + - Operation *op, ArrayRef<LatticeElement<AliasInfo> *> operands) { 728 + +void SharedMemoryAliasAnalysis::visitOperation( 729 + + Operation *op, ArrayRef<const dataflow::Lattice<AliasInfo> *> operands, 730 + + ArrayRef<dataflow::Lattice<AliasInfo> *> results) { 731 + AliasInfo aliasInfo; 732 + bool pessimistic = true; 733 + if (maybeSharedAllocationOp(op)) { 734 + @@ -44,14 +45,11 @@ ChangeResult SharedMemoryAliasAnalysis::visitOperation( 735 + } 736 + 737 + if (pessimistic) { 738 + - return markAllPessimisticFixpoint(op->getResults()); 739 + + return markAllPessimisticFixpoint(results); 740 + } 741 + // Join all lattice elements 742 + - ChangeResult result = ChangeResult::NoChange; 743 + - for (Value value : op->getResults()) { 744 + - result |= getLatticeElement(value).join(aliasInfo); 745 + - } 746 + - return result; 747 + + for (auto *result : results) 748 + + propagateIfChanged(result, result->join(aliasInfo)); 749 + } 750 + 751 + AliasResult SharedMemoryAliasAnalysis::alias(Value lhs, Value rhs) { 752 + diff --git a/lib/Analysis/Allocation.cpp b/lib/Analysis/Allocation.cpp 753 + index 712c08c475..b4de8dcd9d 100644 754 + --- a/lib/Analysis/Allocation.cpp 755 + +++ b/lib/Analysis/Allocation.cpp 756 + @@ -1,4 +1,5 @@ 757 + #include "triton/Analysis/Allocation.h" 758 + +#include "mlir/Analysis/DataFlowFramework.h" 759 + #include "mlir/Analysis/Liveness.h" 760 + #include "mlir/Analysis/SliceAnalysis.h" 761 + #include "mlir/Dialect/Tensor/IR/Tensor.h" 762 + @@ -33,10 +34,8 @@ constexpr int kPtrBitWidth = 64; 763 + 764 + static std::pair<SmallVector<unsigned>, SmallVector<unsigned>> 765 + getCvtOrder(const Attribute &srcLayout, const Attribute &dstLayout) { 766 + - auto srcBlockedLayout = srcLayout.dyn_cast<BlockedEncodingAttr>(); 767 + auto srcMmaLayout = srcLayout.dyn_cast<MmaEncodingAttr>(); 768 + auto srcDotLayout = srcLayout.dyn_cast<DotOperandEncodingAttr>(); 769 + - auto dstBlockedLayout = dstLayout.dyn_cast<BlockedEncodingAttr>(); 770 + auto dstMmaLayout = dstLayout.dyn_cast<MmaEncodingAttr>(); 771 + auto dstDotLayout = dstLayout.dyn_cast<DotOperandEncodingAttr>(); 772 + assert(!(srcMmaLayout && dstMmaLayout) && 773 + @@ -224,14 +223,12 @@ class AllocationAnalysis { 774 + } 775 + 776 + void getValueAlias(Value value, SharedMemoryAliasAnalysis &analysis) { 777 + - LatticeElement<AliasInfo> *latticeElement = 778 + - analysis.lookupLatticeElement(value); 779 + - if (latticeElement) { 780 + - auto &info = latticeElement->getValue(); 781 + - if (!info.getAllocs().empty()) { 782 + - for (auto alloc : info.getAllocs()) { 783 + - allocation->addAlias(value, alloc); 784 + - } 785 + + dataflow::Lattice<AliasInfo> *latticeElement = 786 + + analysis.getLatticeElement(value); 787 + + if (latticeElement && !latticeElement->isUninitialized()) { 788 + + AliasInfo &info = latticeElement->getValue(); 789 + + for (auto alloc : info.getAllocs()) { 790 + + allocation->addAlias(value, alloc); 791 + } 792 + } 793 + } 794 + @@ -244,14 +241,19 @@ class AllocationAnalysis { 795 + getScratchValueSize(op); 796 + }); 797 + // Get the alias values 798 + - SharedMemoryAliasAnalysis aliasAnalysis(operation->getContext()); 799 + - aliasAnalysis.run(operation); 800 + + std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver(); 801 + + SharedMemoryAliasAnalysis *aliasAnalysis = 802 + + solver->load<SharedMemoryAliasAnalysis>(); 803 + + if (failed(solver->initializeAndRun(operation))) { 804 + + // TODO: return error instead of bailing out.. 805 + + llvm_unreachable("failed to run SharedMemoryAliasAnalysis"); 806 + + } 807 + operation->walk<WalkOrder::PreOrder>([&](Operation *op) { 808 + for (auto operand : op->getOperands()) { 809 + - getValueAlias(operand, aliasAnalysis); 810 + + getValueAlias(operand, *aliasAnalysis); 811 + } 812 + for (auto value : op->getResults()) { 813 + - getValueAlias(value, aliasAnalysis); 814 + + getValueAlias(value, *aliasAnalysis); 815 + } 816 + }); 817 + } 818 + diff --git a/lib/Analysis/AxisInfo.cpp b/lib/Analysis/AxisInfo.cpp 819 + index 0b7142b04d..4af46c3fbb 100644 820 + --- a/lib/Analysis/AxisInfo.cpp 821 + +++ b/lib/Analysis/AxisInfo.cpp 822 + @@ -1,4 +1,4 @@ 823 + -#include "mlir/Analysis/DataFlowAnalysis.h" 824 + +#include "mlir/Analysis/DataFlowFramework.h" 825 + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" 826 + #include "llvm/Support/raw_ostream.h" 827 + 828 + @@ -52,7 +52,7 @@ AxisInfo AxisInfo::getPessimisticValueState(Value value) { 829 + BlockArgument blockArg = value.dyn_cast<BlockArgument>(); 830 + if (blockArg && blockArg.getOwner()->isEntryBlock()) { 831 + Operation *op = blockArg.getOwner()->getParentOp(); 832 + - if (FuncOp fun = dyn_cast<FuncOp>(op)) { 833 + + if (func::FuncOp fun = dyn_cast<func::FuncOp>(op)) { 834 + Attribute attr = 835 + fun.getArgAttr(blockArg.getArgNumber(), "tt.divisibility"); 836 + if (attr) 837 + @@ -136,8 +136,9 @@ class CastOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> { 838 + public: 839 + using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl; 840 + 841 + - AxisInfo getAxisInfo(OpTy op, 842 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 843 + + AxisInfo 844 + + getAxisInfo(OpTy op, 845 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 846 + return operands[0]->getValue(); 847 + } 848 + }; 849 + @@ -147,8 +148,9 @@ class MakeRangeOpAxisInfoVisitor final 850 + public: 851 + using AxisInfoVisitorImpl<triton::MakeRangeOp>::AxisInfoVisitorImpl; 852 + 853 + - AxisInfo getAxisInfo(triton::MakeRangeOp op, 854 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 855 + + AxisInfo 856 + + getAxisInfo(triton::MakeRangeOp op, 857 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 858 + auto start = op.start(); 859 + auto end = op.end(); 860 + return AxisInfo(/*contiguity=*/{end - start}, 861 + @@ -162,8 +164,9 @@ class ConstantOpAxisInfoVisitor final 862 + public: 863 + using AxisInfoVisitorImpl<arith::ConstantOp>::AxisInfoVisitorImpl; 864 + 865 + - AxisInfo getAxisInfo(arith::ConstantOp op, 866 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 867 + + AxisInfo 868 + + getAxisInfo(arith::ConstantOp op, 869 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 870 + auto intAttr = op.getValue().dyn_cast<IntegerAttr>(); 871 + auto boolAttr = op.getValue().dyn_cast<BoolAttr>(); 872 + if (intAttr || boolAttr) { 873 + @@ -416,8 +419,9 @@ class SplatOpAxisInfoVisitor final 874 + public: 875 + using AxisInfoVisitorImpl<triton::SplatOp>::AxisInfoVisitorImpl; 876 + 877 + - AxisInfo getAxisInfo(triton::SplatOp op, 878 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 879 + + AxisInfo 880 + + getAxisInfo(triton::SplatOp op, 881 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 882 + Type _retTy = *op->result_type_begin(); 883 + TensorType retTy = _retTy.cast<TensorType>(); 884 + AxisInfo opInfo = operands[0]->getValue(); 885 + @@ -439,8 +443,9 @@ class ExpandDimsOpAxisInfoVisitor final 886 + public: 887 + using AxisInfoVisitorImpl<triton::ExpandDimsOp>::AxisInfoVisitorImpl; 888 + 889 + - AxisInfo getAxisInfo(triton::ExpandDimsOp op, 890 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 891 + + AxisInfo 892 + + getAxisInfo(triton::ExpandDimsOp op, 893 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 894 + AxisInfo opInfo = operands[0]->getValue(); 895 + AxisInfo::DimVectorT contiguity = opInfo.getContiguity(); 896 + AxisInfo::DimVectorT divisibility = opInfo.getDivisibility(); 897 + @@ -458,8 +463,9 @@ class BroadcastOpAxisInfoVisitor final 898 + public: 899 + using AxisInfoVisitorImpl<triton::BroadcastOp>::AxisInfoVisitorImpl; 900 + 901 + - AxisInfo getAxisInfo(triton::BroadcastOp op, 902 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 903 + + AxisInfo 904 + + getAxisInfo(triton::BroadcastOp op, 905 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 906 + Type _retTy = *op->result_type_begin(); 907 + Type _opTy = *op->operand_type_begin(); 908 + TensorType retTy = _retTy.cast<TensorType>(); 909 + @@ -486,8 +492,9 @@ class CmpOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> { 910 + public: 911 + using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl; 912 + 913 + - AxisInfo getAxisInfo(OpTy op, 914 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 915 + + AxisInfo 916 + + getAxisInfo(OpTy op, 917 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 918 + auto resTy = op.getResult().getType().template dyn_cast<RankedTensorType>(); 919 + if (!resTy) 920 + return AxisInfo(); 921 + @@ -596,8 +603,9 @@ class SelectOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> { 922 + public: 923 + using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl; 924 + 925 + - AxisInfo getAxisInfo(OpTy op, 926 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 927 + + AxisInfo 928 + + getAxisInfo(OpTy op, 929 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 930 + auto resTy = op.getResult().getType().template dyn_cast<RankedTensorType>(); 931 + if (!resTy) 932 + return AxisInfo(); 933 + @@ -757,8 +765,9 @@ class MaxMinOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> { 934 + public: 935 + using AxisInfoVisitorImpl<OpTy>::AxisInfoVisitorImpl; 936 + 937 + - AxisInfo getAxisInfo(OpTy op, 938 + - ArrayRef<LatticeElement<AxisInfo> *> operands) override { 939 + + AxisInfo 940 + + getAxisInfo(OpTy op, 941 + + ArrayRef<const dataflow::Lattice<AxisInfo> *> operands) override { 942 + auto lhsInfo = operands[0]->getValue(); 943 + auto rhsInfo = operands[1]->getValue(); 944 + std::optional<int64_t> constantValue; 945 + @@ -786,8 +795,8 @@ class MaxMinOpAxisInfoVisitor final : public AxisInfoVisitorImpl<OpTy> { 946 + // AxisInfoAnalysis 947 + //===----------------------------------------------------------------------===// 948 + 949 + -AxisInfoAnalysis::AxisInfoAnalysis(MLIRContext *context) 950 + - : ForwardDataFlowAnalysis<AxisInfo>(context) { 951 + +AxisInfoAnalysis::AxisInfoAnalysis(DataFlowSolver &solver) 952 + + : dataflow::SparseDataFlowAnalysis<dataflow::Lattice<AxisInfo>>(solver) { 953 + // UnrealizedConversionCast: 954 + // This is needed by TritonGPUToLLVM, to get AxisInfo when the graph is 955 + // in the process of a PartialConversion, where UnrealizedConversionCast 956 + @@ -819,7 +828,7 @@ AxisInfoAnalysis::AxisInfoAnalysis(MLIRContext *context) 957 + visitors.append<LogicalOpAxisInfoVisitor<arith::AndIOp>, 958 + LogicalOpAxisInfoVisitor<arith::OrIOp>, 959 + LogicalOpAxisInfoVisitor<arith::XOrIOp>>(); 960 + - visitors.append<SelectOpAxisInfoVisitor<mlir::SelectOp>, 961 + + visitors.append<SelectOpAxisInfoVisitor<mlir::arith::SelectOp>, 962 + SelectOpAxisInfoVisitor<triton::gpu::SelectOp>>(); 963 + visitors.append<ShLIOpAxisInfoVisitor, ShROpAxisInfoVisitor<arith::ShRUIOp>, 964 + ShROpAxisInfoVisitor<arith::ShRSIOp>>(); 965 + @@ -829,11 +838,12 @@ AxisInfoAnalysis::AxisInfoAnalysis(MLIRContext *context) 966 + MaxMinOpAxisInfoVisitor<arith::MinUIOp>>(); 967 + } 968 + 969 + -ChangeResult AxisInfoAnalysis::visitOperation( 970 + - Operation *op, ArrayRef<LatticeElement<AxisInfo> *> operands) { 971 + +void AxisInfoAnalysis::visitOperation( 972 + + Operation *op, ArrayRef<const dataflow::Lattice<AxisInfo> *> operands, 973 + + ArrayRef<dataflow::Lattice<AxisInfo> *> results) { 974 + AxisInfo curr = visitors.apply(op, operands); 975 + if (curr.getRank() == 0) { 976 + - return markAllPessimisticFixpoint(op->getResults()); 977 + + return markAllPessimisticFixpoint(results); 978 + } 979 + // override with hint 980 + auto newContiguity = curr.getContiguity(); 981 + @@ -854,11 +864,8 @@ ChangeResult AxisInfoAnalysis::visitOperation( 982 + curr = mlir::AxisInfo(newContiguity, newDivisibility, newConstancy, 983 + curr.getConstantValue()); 984 + // join all lattice elements 985 + - ChangeResult result = ChangeResult::NoChange; 986 + - for (Value value : op->getResults()) { 987 + - result |= getLatticeElement(value).join(curr); 988 + - } 989 + - return result; 990 + + for (auto *result : results) 991 + + propagateIfChanged(result, result->join(curr)); 992 + } 993 + 994 + unsigned AxisInfoAnalysis::getPtrContiguity(Value ptr) { 995 + @@ -884,7 +891,10 @@ unsigned AxisInfoAnalysis::getPtrAlignment(Value ptr) { 996 + auto tensorTy = ptr.getType().dyn_cast<RankedTensorType>(); 997 + if (!tensorTy) 998 + return 1; 999 + - auto axisInfo = lookupLatticeElement(ptr)->getValue(); 1000 + + dataflow::Lattice<AxisInfo> *latticeElement = getLatticeElement(ptr); 1001 + + if (!latticeElement || latticeElement->isUninitialized()) 1002 + + return 1; 1003 + + auto axisInfo = latticeElement->getValue(); 1004 + auto layout = tensorTy.getEncoding(); 1005 + auto order = triton::gpu::getOrder(layout); 1006 + auto maxMultipleBytes = axisInfo.getDivisibility(order[0]); 1007 + @@ -900,8 +910,11 @@ unsigned AxisInfoAnalysis::getMaskAlignment(Value mask) { 1008 + auto tensorTy = mask.getType().dyn_cast<RankedTensorType>(); 1009 + if (!tensorTy) 1010 + return 1; 1011 + + dataflow::Lattice<AxisInfo> *latticeElement = getLatticeElement(mask); 1012 + + if (!latticeElement || latticeElement->isUninitialized()) 1013 + + return 1; 1014 + + auto maskAxis = latticeElement->getValue(); 1015 + auto maskOrder = triton::gpu::getOrder(tensorTy.getEncoding()); 1016 + - auto maskAxis = lookupLatticeElement(mask)->getValue(); 1017 + auto alignment = std::max<unsigned>(maskAxis.getConstancy(maskOrder[0]), 1); 1018 + return alignment; 1019 + } 1020 + diff --git a/lib/Analysis/CMakeLists.txt b/lib/Analysis/CMakeLists.txt 1021 + index afbc692510..1f761f845c 100644 1022 + --- a/lib/Analysis/CMakeLists.txt 1023 + +++ b/lib/Analysis/CMakeLists.txt 1024 + @@ -8,7 +8,7 @@ add_mlir_library(TritonAnalysis 1025 + DEPENDS 1026 + TritonTableGen 1027 + TritonGPUAttrDefsIncGen 1028 + - 1029 + + 1030 + LINK_LIBS PUBLIC 1031 + MLIRAnalysis 1032 + ) 1033 + diff --git a/lib/Analysis/Membar.cpp b/lib/Analysis/Membar.cpp 1034 + index acc885e827..910274b2ac 100644 1035 + --- a/lib/Analysis/Membar.cpp 1036 + +++ b/lib/Analysis/Membar.cpp 1037 + @@ -2,7 +2,7 @@ 1038 + #include "triton/Analysis/Alias.h" 1039 + #include "triton/Dialect/TritonGPU/IR/Dialect.h" 1040 + 1041 + -#include "mlir/Dialect/GPU/GPUDialect.h" 1042 + +#include "mlir/Dialect/GPU/IR/GPUDialect.h" 1043 + #include "mlir/Dialect/Tensor/IR/Tensor.h" 1044 + 1045 + namespace mlir { 1046 + diff --git a/lib/Analysis/Utility.cpp b/lib/Analysis/Utility.cpp 1047 + index d9e917e731..6ea52df272 100644 1048 + --- a/lib/Analysis/Utility.cpp 1049 + +++ b/lib/Analysis/Utility.cpp 1050 + @@ -1,5 +1,8 @@ 1051 + #include "triton/Analysis/Utility.h" 1052 + +#include "mlir/Analysis/DataFlow/ConstantPropagationAnalysis.h" 1053 + +#include "mlir/Analysis/DataFlow/DeadCodeAnalysis.h" 1054 + #include "mlir/IR/Dialect.h" 1055 + +#include "mlir/IR/Matchers.h" 1056 + #include "triton/Dialect/Triton/IR/Dialect.h" 1057 + #include "triton/Dialect/TritonGPU/IR/Dialect.h" 1058 + #include <deque> 1059 + @@ -325,4 +328,55 @@ SetVector<Operation *> multiRootGetSlice(Operation *op, 1060 + return multiRootTopologicalSort(slice); 1061 + } 1062 + 1063 + +namespace { 1064 + +// Copied from TestDeadCodeAnalysis.cpp, because some dead code analysis 1065 + +// interacts with constant propagation, but SparseConstantPropagation 1066 + +// doesn't seem to be sufficient. 1067 + +struct ConstantAnalysis : public DataFlowAnalysis { 1068 + + using DataFlowAnalysis::DataFlowAnalysis; 1069 + + 1070 + + LogicalResult initialize(Operation *top) override { 1071 + + WalkResult result = top->walk([&](Operation *op) { 1072 + + if (failed(visit(op))) 1073 + + return WalkResult::interrupt(); 1074 + + return WalkResult::advance(); 1075 + + }); 1076 + + return success(!result.wasInterrupted()); 1077 + + } 1078 + + 1079 + + LogicalResult visit(ProgramPoint point) override { 1080 + + Operation *op = point.get<Operation *>(); 1081 + + Attribute value; 1082 + + if (matchPattern(op, m_Constant(&value))) { 1083 + + auto *constant = getOrCreate<dataflow::Lattice<dataflow::ConstantValue>>( 1084 + + op->getResult(0)); 1085 + + propagateIfChanged(constant, constant->join(dataflow::ConstantValue( 1086 + + value, op->getDialect()))); 1087 + + return success(); 1088 + + } 1089 + + setAllToUnknownConstants(op->getResults()); 1090 + + for (Region &region : op->getRegions()) 1091 + + setAllToUnknownConstants(region.getArguments()); 1092 + + return success(); 1093 + + } 1094 + + 1095 + + /// Set all given values as not constants. 1096 + + void setAllToUnknownConstants(ValueRange values) { 1097 + + dataflow::ConstantValue unknownConstant(nullptr, nullptr); 1098 + + for (Value value : values) { 1099 + + auto *constant = 1100 + + getOrCreate<dataflow::Lattice<dataflow::ConstantValue>>(value); 1101 + + propagateIfChanged(constant, constant->join(unknownConstant)); 1102 + + } 1103 + + } 1104 + +}; 1105 + +} // namespace 1106 + + 1107 + +std::unique_ptr<DataFlowSolver> createDataFlowSolver() { 1108 + + auto solver = std::make_unique<DataFlowSolver>(); 1109 + + solver->load<dataflow::DeadCodeAnalysis>(); 1110 + + solver->load<ConstantAnalysis>(); 1111 + + return solver; 1112 + +} 1113 + + 1114 + } // namespace mlir 1115 + diff --git a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp 1116 + index 6a46265bd7..e352eb3698 100644 1117 + --- a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp 1118 + +++ b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp 1119 + @@ -159,9 +159,6 @@ struct ConvertLayoutOpConversion 1120 + Value smemBase) const { 1121 + auto accumNumCTAsEachRep = product<unsigned>(numCTAsEachRep); 1122 + auto layout = type.getEncoding(); 1123 + - auto blockedLayout = layout.dyn_cast<BlockedEncodingAttr>(); 1124 + - auto sliceLayout = layout.dyn_cast<SliceEncodingAttr>(); 1125 + - auto mmaLayout = layout.dyn_cast<MmaEncodingAttr>(); 1126 + auto rank = type.getRank(); 1127 + auto sizePerThread = getSizePerThread(layout); 1128 + auto accumSizePerThread = product<unsigned>(sizePerThread); 1129 + diff --git a/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h b/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h 1130 + index 4b89965aa9..1d9e00519b 100644 1131 + --- a/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h 1132 + +++ b/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h 1133 + @@ -7,10 +7,8 @@ 1134 + #include "mlir/Conversion/LLVMCommon/LoweringOptions.h" 1135 + #include "mlir/Conversion/LLVMCommon/Pattern.h" 1136 + #include "mlir/Conversion/MathToLLVM/MathToLLVM.h" 1137 + -#include "mlir/Conversion/SCFToStandard/SCFToStandard.h" 1138 + -#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h" 1139 + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" 1140 + -#include "mlir/Dialect/GPU/GPUDialect.h" 1141 + +#include "mlir/Dialect/GPU/IR/GPUDialect.h" 1142 + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" 1143 + #include "mlir/Dialect/Tensor/IR/Tensor.h" 1144 + #include "mlir/IR/Matchers.h" 1145 + @@ -422,9 +420,9 @@ struct MMA16816ConversionHelper { 1146 + MMA16816ConversionHelper(Type dotOperand, MmaEncodingAttr mmaLayout, 1147 + Value thread, ConversionPatternRewriter &rewriter, 1148 + TypeConverter *typeConverter, Location loc) 1149 + - : mmaLayout(mmaLayout), thread(thread), helper(mmaLayout), 1150 + - rewriter(rewriter), typeConverter(typeConverter), loc(loc), 1151 + - ctx(mmaLayout.getContext()), wpt(mmaLayout.getWarpsPerCTA()) { 1152 + + : mmaLayout(mmaLayout), wpt(mmaLayout.getWarpsPerCTA()), thread(thread), 1153 + + helper(mmaLayout), rewriter(rewriter), typeConverter(typeConverter), 1154 + + loc(loc), ctx(mmaLayout.getContext()) { 1155 + helper.deduceMmaType(dotOperand); 1156 + 1157 + Value _32 = i32_val(32); 1158 + diff --git a/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp 1159 + index 0f8070ca9f..e4bd47c411 100644 1160 + --- a/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp 1161 + +++ b/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp 1162 + @@ -115,8 +115,6 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern<triton::DotOp> { 1163 + auto DTensorTy = D.getType().cast<RankedTensorType>(); 1164 + auto AShape = ATensorTy.getShape(); 1165 + auto BShape = BTensorTy.getShape(); 1166 + - auto DShape = DTensorTy.getShape(); 1167 + - auto wpt = mmaLayout.getWarpsPerCTA(); 1168 + 1169 + bool isARow = ALayout.getIsMMAv1Row().cast<BoolAttr>().getValue(); 1170 + bool isBRow = BLayout.getIsMMAv1Row().cast<BoolAttr>().getValue(); 1171 + @@ -221,7 +219,6 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern<triton::DotOp> { 1172 + ConversionPatternRewriter &rewriter) const { 1173 + auto *ctx = rewriter.getContext(); 1174 + auto loc = op.getLoc(); 1175 + - auto threadId = getThreadId(rewriter, loc); 1176 + 1177 + auto A = op.a(); 1178 + auto B = op.b(); 1179 + @@ -230,12 +227,10 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern<triton::DotOp> { 1180 + 1181 + auto aTensorTy = A.getType().cast<RankedTensorType>(); 1182 + auto bTensorTy = B.getType().cast<RankedTensorType>(); 1183 + - auto cTensorTy = C.getType().cast<RankedTensorType>(); 1184 + auto dTensorTy = D.getType().cast<RankedTensorType>(); 1185 + 1186 + auto aShape = aTensorTy.getShape(); 1187 + auto bShape = bTensorTy.getShape(); 1188 + - auto cShape = cTensorTy.getShape(); 1189 + 1190 + BlockedEncodingAttr dLayout = 1191 + dTensorTy.getEncoding().cast<BlockedEncodingAttr>(); 1192 + diff --git a/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp 1193 + index deb71b9597..0b9e67674b 100644 1194 + --- a/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp 1195 + +++ b/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp 1196 + @@ -61,7 +61,6 @@ struct FpToFpOpConversion 1197 + convertFp16x4ToFp8x4(Location loc, ConversionPatternRewriter &rewriter, 1198 + const Value &v0, const Value &v1, const Value &v2, 1199 + const Value &v3) { 1200 + - auto ctx = rewriter.getContext(); 1201 + auto fp16x2VecTy = vec_ty(f16_ty, 2); 1202 + Value fp16x2Vec0 = undef(fp16x2VecTy); 1203 + Value fp16x2Vec1 = undef(fp16x2VecTy); 1204 + @@ -153,7 +152,6 @@ struct FpToFpOpConversion 1205 + convertBf16x4ToFp8x4(Location loc, ConversionPatternRewriter &rewriter, 1206 + const Value &v0, const Value &v1, const Value &v2, 1207 + const Value &v3) { 1208 + - auto ctx = rewriter.getContext(); 1209 + auto bf16x2VecTy = vec_ty(i16_ty, 2); 1210 + Value bf16x2Vec0 = undef(bf16x2VecTy); 1211 + Value bf16x2Vec1 = undef(bf16x2VecTy); 1212 + diff --git a/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp 1213 + index 9a8b4702bc..bae675f0cb 100644 1214 + --- a/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp 1215 + +++ b/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp 1216 + @@ -109,7 +109,8 @@ struct LoadOpConversion 1217 + DenseElementsAttr constAttr; 1218 + int64_t splatVal = 0; 1219 + if (other && valueElemTy.isa<IntegerType>() && 1220 + - matchPattern(other, m_Constant(&constAttr)) && constAttr.isSplat()) { 1221 + + matchPattern(other, m_Constant(&constAttr)) && constAttr.isSplat() && 1222 + + constAttr.getElementType().isa<IntegerType>()) { 1223 + otherIsSplatConstInt = true; 1224 + splatVal = constAttr.getSplatValue<APInt>().getSExtValue(); 1225 + } 1226 + @@ -333,7 +334,6 @@ struct StoreOpConversion 1227 + elem = rewriter.create<LLVM::SExtOp>(loc, type::i8Ty(ctx), elem); 1228 + elem = bitcast(elem, valueElemTy); 1229 + 1230 + - Type u32Ty = typeConverter->convertType(type::u32Ty(ctx)); 1231 + llWord = insert_element(wordTy, llWord, elem, i32_val(elemIdx)); 1232 + } 1233 + llWord = bitcast(llWord, valArgTy); 1234 + @@ -387,7 +387,6 @@ struct AtomicCASOpConversion 1235 + ConversionPatternRewriter &rewriter) const override { 1236 + auto loc = op.getLoc(); 1237 + MLIRContext *ctx = rewriter.getContext(); 1238 + - Value ptr = op.ptr(); 1239 + 1240 + Value llPtr = adaptor.ptr(); 1241 + Value llCmp = adaptor.cmp(); 1242 + diff --git a/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp 1243 + index 69abd889be..1c973dc196 100644 1244 + --- a/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp 1245 + +++ b/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp 1246 + @@ -286,7 +286,6 @@ struct ReduceOpConversion 1247 + auto srcTy = op.operand().getType().cast<RankedTensorType>(); 1248 + auto srcLayout = srcTy.getEncoding(); 1249 + auto srcShape = srcTy.getShape(); 1250 + - auto srcRank = srcTy.getRank(); 1251 + auto order = getOrder(srcLayout); 1252 + 1253 + auto threadsPerWarp = triton::gpu::getThreadsPerWarp(srcLayout); 1254 + @@ -351,7 +350,6 @@ struct ReduceOpConversion 1255 + 1256 + Value zero = i32_val(0); 1257 + Value laneZero = icmp_eq(laneIdAxis, zero); 1258 + - Value warpZero = icmp_eq(warpIdAxis, zero); 1259 + 1260 + for (auto it : accs) { 1261 + const SmallVector<unsigned> &key = it.first; 1262 + diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp 1263 + index 5b77150b1a..78cfa076bd 100644 1264 + --- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp 1265 + +++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp 1266 + @@ -11,11 +11,11 @@ using ::mlir::LLVM::getStructFromElements; 1267 + using ::mlir::triton::gpu::getElemsPerThread; 1268 + using ::mlir::triton::gpu::SharedEncodingAttr; 1269 + 1270 + -struct ReturnOpConversion : public ConvertOpToLLVMPattern<::mlir::ReturnOp> { 1271 + - using ConvertOpToLLVMPattern<ReturnOp>::ConvertOpToLLVMPattern; 1272 + +struct ReturnOpConversion : public ConvertOpToLLVMPattern<func::ReturnOp> { 1273 + + using ConvertOpToLLVMPattern<func::ReturnOp>::ConvertOpToLLVMPattern; 1274 + 1275 + LogicalResult 1276 + - matchAndRewrite(ReturnOp op, OpAdaptor adaptor, 1277 + + matchAndRewrite(func::ReturnOp op, OpAdaptor adaptor, 1278 + ConversionPatternRewriter &rewriter) const override { 1279 + unsigned numArguments = op.getNumOperands(); 1280 + 1281 + @@ -476,7 +476,6 @@ struct ExtractSliceOpConversion 1282 + 1283 + auto llvmElemTy = getTypeConverter()->convertType(srcTy.getElementType()); 1284 + auto elemPtrTy = ptr_ty(llvmElemTy, 3); 1285 + - auto resTy = op.getType().dyn_cast<RankedTensorType>(); 1286 + smemObj = SharedMemoryObject(gep(elemPtrTy, smemObj.base, offset), 1287 + strideVals, offsetVals); 1288 + auto retVal = getStructFromSharedMemoryObject(loc, smemObj, rewriter); 1289 + diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h 1290 + index bb10d5b24a..00e399f848 100644 1291 + --- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h 1292 + +++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h 1293 + @@ -4,6 +4,7 @@ 1294 + // TODO: refactor so that it doesn't fail if Allocation.h 1295 + // is included after utility.h (due to conflict in `store` macro 1296 + // and <atomic> 1297 + +#include "mlir/Dialect/Func/IR/FuncOps.h" 1298 + #include "triton/Analysis/Allocation.h" 1299 + 1300 + // 1301 + @@ -39,15 +40,15 @@ void vprintf_array(Value thread, ArrayRef<Value> arr, std::string info, 1302 + // TODO(Superjomn): remove the code when MLIR v15.0 is included. 1303 + // All the rights are reserved by the LLVM community. 1304 + 1305 + -struct FuncOpConversionBase : public ConvertOpToLLVMPattern<FuncOp> { 1306 + +struct FuncOpConversionBase : public ConvertOpToLLVMPattern<func::FuncOp> { 1307 + private: 1308 + /// Only retain those attributes that are not constructed by 1309 + /// `LLVMFuncOp::build`. If `filterArgAttrs` is set, also filter out argument 1310 + /// attributes. 1311 + - static void filterFuncAttributes(ArrayRef<NamedAttribute> attrs, 1312 + - bool filterArgAttrs, 1313 + + static void filterFuncAttributes(func::FuncOp op, bool filterArgAttrs, 1314 + SmallVectorImpl<NamedAttribute> &result) { 1315 + - for (const auto &attr : attrs) { 1316 + + 1317 + + for (const auto &attr : op->getAttrs()) { 1318 + if (attr.getName() == SymbolTable::getSymbolAttrName() || 1319 + attr.getName() == FunctionOpInterface::getTypeAttrName() || 1320 + attr.getName() == "std.varargs" || 1321 + @@ -65,27 +66,27 @@ struct FuncOpConversionBase : public ConvertOpToLLVMPattern<FuncOp> { 1322 + } 1323 + 1324 + protected: 1325 + - using ConvertOpToLLVMPattern<FuncOp>::ConvertOpToLLVMPattern; 1326 + + using ConvertOpToLLVMPattern<func::FuncOp>::ConvertOpToLLVMPattern; 1327 + 1328 + // Convert input FuncOp to LLVMFuncOp by using the LLVMTypeConverter provided 1329 + // to this legalization pattern. 1330 + LLVM::LLVMFuncOp 1331 + - convertFuncOpToLLVMFuncOp(FuncOp funcOp, 1332 + + convertFuncOpToLLVMFuncOp(func::FuncOp funcOp, 1333 + ConversionPatternRewriter &rewriter) const { 1334 + // Convert the original function arguments. They are converted using the 1335 + // LLVMTypeConverter provided to this legalization pattern. 1336 + auto varargsAttr = funcOp->getAttrOfType<BoolAttr>("func.varargs"); 1337 + TypeConverter::SignatureConversion result(funcOp.getNumArguments()); 1338 + auto llvmType = getTypeConverter()->convertFunctionSignature( 1339 + - funcOp.getType(), varargsAttr && varargsAttr.getValue(), result); 1340 + + funcOp.getFunctionType(), varargsAttr && varargsAttr.getValue(), 1341 + + result); 1342 + if (!llvmType) 1343 + return nullptr; 1344 + 1345 + // Propagate argument/result attributes to all converted arguments/result 1346 + // obtained after converting a given original argument/result. 1347 + SmallVector<NamedAttribute, 4> attributes; 1348 + - filterFuncAttributes(funcOp->getAttrs(), /*filterArgAttrs=*/true, 1349 + - attributes); 1350 + + filterFuncAttributes(funcOp, /*filterArgAttrs=*/true, attributes); 1351 + if (ArrayAttr resAttrDicts = funcOp.getAllResultAttrs()) { 1352 + assert(!resAttrDicts.empty() && "expected array to be non-empty"); 1353 + auto newResAttrDicts = 1354 + @@ -131,7 +132,7 @@ struct FuncOpConversionBase : public ConvertOpToLLVMPattern<FuncOp> { 1355 + } 1356 + auto newFuncOp = rewriter.create<LLVM::LLVMFuncOp>( 1357 + funcOp.getLoc(), funcOp.getName(), llvmType, linkage, 1358 + - /*dsoLocal*/ false, attributes); 1359 + + /*dsoLocal*/ false, LLVM::CConv::C, attributes); 1360 + rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(), 1361 + newFuncOp.end()); 1362 + if (failed(rewriter.convertRegionTypes(&newFuncOp.getBody(), *typeConverter, 1363 + @@ -191,8 +192,8 @@ class ConvertTritonGPUOpToLLVMPatternBase { 1364 + const Allocation *allocation, 1365 + Value smem, 1366 + IndexCacheInfo indexCacheInfo) 1367 + - : converter(&typeConverter), indexCacheInfo(indexCacheInfo), 1368 + - allocation(allocation), smem(smem) {} 1369 + + : converter(&typeConverter), allocation(allocation), smem(smem), 1370 + + indexCacheInfo(indexCacheInfo) {} 1371 + 1372 + LLVMTypeConverter *getTypeConverter() const { return converter; } 1373 + 1374 + @@ -861,7 +862,6 @@ class ConvertTritonGPUOpToLLVMPatternBase { 1375 + ArrayRef<int64_t> shape) const { 1376 + auto parent = sliceLayout.getParent(); 1377 + unsigned dim = sliceLayout.getDim(); 1378 + - size_t rank = shape.size(); 1379 + auto parentIndices = 1380 + emitIndices(loc, rewriter, parent, sliceLayout.paddedShape(shape)); 1381 + unsigned numIndices = parentIndices.size(); 1382 + diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp 1383 + index ff1af09835..6f66af4e34 100644 1384 + --- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp 1385 + +++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp 1386 + @@ -1,10 +1,11 @@ 1387 + #include "triton/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.h" 1388 + 1389 + +#include "mlir/Analysis/DataFlowFramework.h" 1390 + #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" 1391 + +#include "mlir/Conversion/ControlFlowToLLVM//ControlFlowToLLVM.h" 1392 + #include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h" 1393 + #include "mlir/Conversion/MathToLLVM/MathToLLVM.h" 1394 + -#include "mlir/Conversion/SCFToStandard/SCFToStandard.h" 1395 + -#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h" 1396 + +#include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" 1397 + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" 1398 + #include "mlir/Dialect/LLVMIR/NVVMDialect.h" 1399 + #include "mlir/Pass/Pass.h" 1400 + @@ -40,7 +41,6 @@ class TritonLLVMConversionTarget : public ConversionTarget { 1401 + addIllegalDialect<triton::TritonDialect>(); 1402 + addIllegalDialect<triton::gpu::TritonGPUDialect>(); 1403 + addIllegalDialect<mlir::gpu::GPUDialect>(); 1404 + - addIllegalDialect<mlir::StandardOpsDialect>(); 1405 + addLegalOp<mlir::UnrealizedConversionCastOp>(); 1406 + } 1407 + }; 1408 + @@ -51,7 +51,7 @@ class TritonLLVMFunctionConversionTarget : public ConversionTarget { 1409 + : ConversionTarget(ctx) { 1410 + addLegalDialect<LLVM::LLVMDialect>(); 1411 + addLegalDialect<NVVM::NVVMDialect>(); 1412 + - addIllegalOp<mlir::FuncOp>(); 1413 + + addIllegalOp<mlir::func::FuncOp>(); 1414 + addLegalOp<mlir::UnrealizedConversionCastOp>(); 1415 + } 1416 + }; 1417 + @@ -69,7 +69,7 @@ struct FuncOpConversion : public FuncOpConversionBase { 1418 + : FuncOpConversionBase(converter, benefit), numWarps(numWarps) {} 1419 + 1420 + LogicalResult 1421 + - matchAndRewrite(FuncOp funcOp, OpAdaptor adaptor, 1422 + + matchAndRewrite(func::FuncOp funcOp, OpAdaptor adaptor, 1423 + ConversionPatternRewriter &rewriter) const override { 1424 + auto newFuncOp = convertFuncOpToLLVMFuncOp(funcOp, rewriter); 1425 + if (!newFuncOp) 1426 + @@ -133,7 +133,8 @@ class ConvertTritonGPUToLLVM 1427 + decomposeBlockedToDotOperand(mod); 1428 + 1429 + // Step 2 1430 + - decomposeInsertSliceAsyncOp(mod); 1431 + + if (failed(decomposeInsertSliceAsyncOp(mod))) 1432 + + return signalPassFailure(); 1433 + 1434 + // Step 3 1435 + Allocation allocation(mod); 1436 + @@ -142,7 +143,7 @@ class ConvertTritonGPUToLLVM 1437 + 1438 + // Step 4 1439 + RewritePatternSet scf_patterns(context); 1440 + - mlir::populateLoopToStdConversionPatterns(scf_patterns); 1441 + + mlir::populateSCFToControlFlowConversionPatterns(scf_patterns); 1442 + mlir::ConversionTarget scf_target(*context); 1443 + scf_target.addIllegalOp<scf::ForOp, scf::IfOp, scf::ParallelOp, 1444 + scf::WhileOp, scf::ExecuteRegionOp>(); 1445 + @@ -159,8 +160,10 @@ class ConvertTritonGPUToLLVM 1446 + return signalPassFailure(); 1447 + 1448 + // Step 6 - get axis and shared memory info 1449 + - AxisInfoAnalysis axisInfoAnalysis(mod.getContext()); 1450 + - axisInfoAnalysis.run(mod); 1451 + + std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver(); 1452 + + AxisInfoAnalysis *axisInfoAnalysis = solver->load<AxisInfoAnalysis>(); 1453 + + if (failed(solver->initializeAndRun(mod))) 1454 + + return signalPassFailure(); 1455 + initSharedMemory(allocation.getSharedMemorySize(), typeConverter); 1456 + mod->setAttr("triton_gpu.shared", 1457 + mlir::IntegerAttr::get(mlir::IntegerType::get(context, 32), 1458 + @@ -178,38 +181,39 @@ class ConvertTritonGPUToLLVM 1459 + 1460 + // Normal conversions 1461 + populateTritonGPUToLLVMPatterns(typeConverter, patterns, numWarps, 1462 + - axisInfoAnalysis, &allocation, smem, 1463 + + *axisInfoAnalysis, &allocation, smem, 1464 + indexCacheInfo, /*benefit=*/10); 1465 + // ConvertLayoutOp 1466 + populateConvertLayoutOpToLLVMPatterns(typeConverter, patterns, numWarps, 1467 + - axisInfoAnalysis, &allocation, smem, 1468 + + *axisInfoAnalysis, &allocation, smem, 1469 + indexCacheInfo, /*benefit=*/10); 1470 + // DotOp 1471 + populateDotOpToLLVMPatterns(typeConverter, patterns, numWarps, 1472 + - axisInfoAnalysis, &allocation, smem, 1473 + + *axisInfoAnalysis, &allocation, smem, 1474 + /*benefit=*/10); 1475 + // ElementwiseOp 1476 + populateElementwiseOpToLLVMPatterns(typeConverter, patterns, numWarps, 1477 + - axisInfoAnalysis, &allocation, smem, 1478 + + *axisInfoAnalysis, &allocation, smem, 1479 + /*benefit=*/10); 1480 + // LoadStoreOp 1481 + populateLoadStoreOpToLLVMPatterns(typeConverter, patterns, numWarps, 1482 + - axisInfoAnalysis, &allocation, smem, 1483 + + *axisInfoAnalysis, &allocation, smem, 1484 + indexCacheInfo, /*benefit=*/10); 1485 + // ReduceOp 1486 + populateReduceOpToLLVMPatterns(typeConverter, patterns, numWarps, 1487 + - axisInfoAnalysis, &allocation, smem, 1488 + + *axisInfoAnalysis, &allocation, smem, 1489 + indexCacheInfo, /*benefit=*/10); 1490 + // ViewOp 1491 + populateViewOpToLLVMPatterns(typeConverter, patterns, numWarps, 1492 + - axisInfoAnalysis, &allocation, smem, 1493 + + *axisInfoAnalysis, &allocation, smem, 1494 + /*benefit=*/10); 1495 + 1496 + // Add arith/math's patterns to help convert scalar expression to LLVM. 1497 + mlir::arith::populateArithmeticToLLVMConversionPatterns(typeConverter, 1498 + patterns); 1499 + mlir::populateMathToLLVMConversionPatterns(typeConverter, patterns); 1500 + - mlir::populateStdToLLVMConversionPatterns(typeConverter, patterns); 1501 + + mlir::cf::populateControlFlowToLLVMConversionPatterns(typeConverter, 1502 + + patterns); 1503 + mlir::populateGpuToNVVMConversionPatterns(typeConverter, patterns); 1504 + 1505 + if (failed(applyPartialConversion(mod, target, std::move(patterns)))) 1506 + @@ -306,9 +310,11 @@ class ConvertTritonGPUToLLVM 1507 + }); 1508 + } 1509 + 1510 + - void decomposeInsertSliceAsyncOp(ModuleOp mod) const { 1511 + - AxisInfoAnalysis axisInfoAnalysis(mod.getContext()); 1512 + - axisInfoAnalysis.run(mod); 1513 + + LogicalResult decomposeInsertSliceAsyncOp(ModuleOp mod) const { 1514 + + std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver(); 1515 + + AxisInfoAnalysis *axisInfoAnalysis = solver->load<AxisInfoAnalysis>(); 1516 + + if (failed(solver->initializeAndRun(mod))) 1517 + + return failure(); 1518 + // TODO(Keren): This is a hacky knob that may cause performance regression 1519 + // when decomposition has been performed. We should remove this knob once we 1520 + // have thorough analysis on async wait. Currently, we decompose 1521 + @@ -342,7 +348,7 @@ class ConvertTritonGPUToLLVM 1522 + auto resSharedLayout = 1523 + dstTy.getEncoding().dyn_cast<triton::gpu::SharedEncodingAttr>(); 1524 + auto resElemTy = dstTy.getElementType(); 1525 + - unsigned inVec = axisInfoAnalysis.getPtrContiguity(src); 1526 + + unsigned inVec = axisInfoAnalysis->getPtrContiguity(src); 1527 + unsigned outVec = resSharedLayout.getVec(); 1528 + unsigned minVec = std::min(outVec, inVec); 1529 + auto maxBitWidth = 1530 + @@ -400,11 +406,11 @@ class ConvertTritonGPUToLLVM 1531 + } else if (decomposed) { 1532 + // Wait for all previous async ops 1533 + OpBuilder builder(asyncWaitOp); 1534 + - auto newAsyncWaitOp = 1535 + - builder.create<triton::gpu::AsyncWaitOp>(asyncWaitOp.getLoc(), 0); 1536 + + builder.create<triton::gpu::AsyncWaitOp>(asyncWaitOp.getLoc(), 0); 1537 + asyncWaitOp.erase(); 1538 + } 1539 + }); 1540 + + return success(); 1541 + } 1542 + }; 1543 + 1544 + diff --git a/lib/Conversion/TritonGPUToLLVM/Utility.h b/lib/Conversion/TritonGPUToLLVM/Utility.h 1545 + index d35dac28c5..11976908cf 100644 1546 + --- a/lib/Conversion/TritonGPUToLLVM/Utility.h 1547 + +++ b/lib/Conversion/TritonGPUToLLVM/Utility.h 1548 + @@ -220,10 +220,7 @@ struct SharedMemoryObject { 1549 + ConversionPatternRewriter &rewriter) 1550 + : base(base) { 1551 + strides = getStridesFromShapeAndOrder(shape, order, loc, rewriter); 1552 + - 1553 + - for (auto idx : order) { 1554 + - offsets.emplace_back(i32_val(0)); 1555 + - } 1556 + + offsets.append(order.size(), i32_val(0)); 1557 + } 1558 + 1559 + SmallVector<Value> getElems() const { 1560 + diff --git a/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp b/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp 1561 + index fe42202c34..5f230f787f 100644 1562 + --- a/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp 1563 + +++ b/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp 1564 + @@ -1,10 +1,10 @@ 1565 + #include "triton/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.h" 1566 + 1567 + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" 1568 + -#include "mlir/Dialect/GPU/GPUDialect.h" 1569 + +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" 1570 + +#include "mlir/Dialect/GPU/IR/GPUDialect.h" 1571 + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" 1572 + #include "mlir/Dialect/LLVMIR/NVVMDialect.h" 1573 + -#include "mlir/Dialect/StandardOps/IR/Ops.h" 1574 + #include "mlir/Pass/Pass.h" 1575 + #include "mlir/Transforms/DialectConversion.h" 1576 + #include "triton/Dialect/Triton/IR/Dialect.h" 1577 + @@ -59,10 +59,13 @@ class ArithConstantPattern : public OpConversionPattern<arith::ConstantOp> { 1578 + Type retType = getTypeConverter()->convertType(op.getType()); 1579 + auto value = adaptor.getValue().dyn_cast<DenseElementsAttr>(); 1580 + assert(value); 1581 + - rewriter.replaceOpWithNewOp<arith::ConstantOp>( 1582 + - op, retType, 1583 + - value.reshape(retType) // This is a hack. We just want to add encoding 1584 + - ); 1585 + + if (value.getElementType().isInteger(1) && value.isSplat()) 1586 + + // Workaround until https://reviews.llvm.org/D133743 is included. 1587 + + value = DenseElementsAttr::get(retType, value.getSplatValue<bool>()); 1588 + + else 1589 + + // This is a hack. We just want to add encoding 1590 + + value = value.reshape(retType); 1591 + + rewriter.replaceOpWithNewOp<arith::ConstantOp>(op, retType, value); 1592 + return success(); 1593 + } 1594 + }; 1595 + @@ -127,12 +130,12 @@ void populateArithmeticPatternsAndLegality( 1596 + } 1597 + 1598 + // this shouldn't exist if mlir's SelectOp checked encodings properly 1599 + -class StdSelectPattern : public OpConversionPattern<SelectOp> { 1600 + +class StdSelectPattern : public OpConversionPattern<arith::SelectOp> { 1601 + public: 1602 + - using OpConversionPattern<SelectOp>::OpConversionPattern; 1603 + + using OpConversionPattern<arith::SelectOp>::OpConversionPattern; 1604 + 1605 + LogicalResult 1606 + - matchAndRewrite(SelectOp op, typename SelectOp::Adaptor adaptor, 1607 + + matchAndRewrite(arith::SelectOp op, OpAdaptor adaptor, 1608 + ConversionPatternRewriter &rewriter) const override { 1609 + Type retType = this->getTypeConverter()->convertType(op.getType()); 1610 + rewriter.replaceOpWithNewOp<triton::gpu::SelectOp>( 1611 + @@ -148,8 +151,8 @@ void populateStdPatternsAndLegality(TritonGPUTypeConverter &typeConverter, 1612 + MLIRContext *context = patterns.getContext(); 1613 + // Rewrite rule 1614 + patterns.add<StdSelectPattern>(typeConverter, context); 1615 + - target.addLegalOp<ReturnOp>(); // this is ok because all functions are inlined 1616 + - // by the frontend 1617 + + target.addLegalOp<func::ReturnOp>(); // this is ok because all functions are 1618 + + // inlined by the frontend 1619 + } 1620 + 1621 + void populateMathPatternsAndLegality(TritonGPUTypeConverter &typeConverter, 1622 + @@ -455,18 +458,19 @@ struct TritonPrintfPattern : public OpConversionPattern<triton::PrintfOp> { 1623 + void populateTritonPatterns(TritonGPUTypeConverter &typeConverter, 1624 + RewritePatternSet &patterns) { 1625 + MLIRContext *context = patterns.getContext(); 1626 + - patterns.add< // TODO: view should have custom pattern that views the layout 1627 + - TritonGenericPattern<triton::ViewOp>, 1628 + - TritonGenericPattern<triton::BitcastOp>, 1629 + - TritonGenericPattern<triton::FpToFpOp>, 1630 + - TritonGenericPattern<triton::IntToPtrOp>, 1631 + - TritonGenericPattern<triton::PtrToIntOp>, 1632 + - TritonGenericPattern<triton::SplatOp>, TritonBroadcastPattern, 1633 + - TritonGenericPattern<triton::AddPtrOp>, TritonCatPattern, 1634 + - TritonReducePattern, TritonTransPattern, TritonExpandDimsPattern, 1635 + - TritonMakeRangePattern, TritonDotPattern, TritonLoadPattern, 1636 + - TritonStorePattern, TritonExtElemwisePattern, TritonPrintfPattern, 1637 + - TritonAtomicRMWPattern>(typeConverter, context); 1638 + + patterns 1639 + + .insert< // TODO: view should have custom pattern that views the layout 1640 + + TritonGenericPattern<triton::ViewOp>, 1641 + + TritonGenericPattern<triton::BitcastOp>, 1642 + + TritonGenericPattern<triton::FpToFpOp>, 1643 + + TritonGenericPattern<triton::IntToPtrOp>, 1644 + + TritonGenericPattern<triton::PtrToIntOp>, 1645 + + TritonGenericPattern<triton::SplatOp>, TritonBroadcastPattern, 1646 + + TritonGenericPattern<triton::AddPtrOp>, TritonCatPattern, 1647 + + TritonReducePattern, TritonTransPattern, TritonExpandDimsPattern, 1648 + + TritonMakeRangePattern, TritonDotPattern, TritonLoadPattern, 1649 + + TritonStorePattern, TritonExtElemwisePattern, TritonPrintfPattern, 1650 + + TritonAtomicRMWPattern>(typeConverter, context); 1651 + } 1652 + 1653 + // 1654 + @@ -623,29 +627,28 @@ void populateSCFPatterns(TritonGPUTypeConverter &typeConverter, 1655 + 1656 + // CF 1657 + 1658 + -class CFBranchPattern : public OpConversionPattern<BranchOp> { 1659 + +class CFBranchPattern : public OpConversionPattern<cf::BranchOp> { 1660 + public: 1661 + - using OpConversionPattern<BranchOp>::OpConversionPattern; 1662 + + using OpConversionPattern<cf::BranchOp>::OpConversionPattern; 1663 + 1664 + LogicalResult 1665 + - matchAndRewrite(BranchOp op, BranchOp::Adaptor adaptor, 1666 + + matchAndRewrite(cf::BranchOp op, cf::BranchOp::Adaptor adaptor, 1667 + ConversionPatternRewriter &rewriter) const override { 1668 + - auto converter = getTypeConverter(); 1669 + - auto newOp = rewriter.replaceOpWithNewOp<BranchOp>(op, op.getSuccessor(), 1670 + - adaptor.getOperands()); 1671 + + auto newOp = rewriter.replaceOpWithNewOp<cf::BranchOp>( 1672 + + op, op.getSuccessor(), adaptor.getOperands()); 1673 + return success(); 1674 + } 1675 + }; 1676 + 1677 + -class CFCondBranchPattern : public OpConversionPattern<CondBranchOp> { 1678 + +class CFCondBranchPattern : public OpConversionPattern<cf::CondBranchOp> { 1679 + public: 1680 + - using OpConversionPattern<CondBranchOp>::OpConversionPattern; 1681 + + using OpConversionPattern<cf::CondBranchOp>::OpConversionPattern; 1682 + 1683 + LogicalResult 1684 + - matchAndRewrite(CondBranchOp op, CondBranchOp::Adaptor adaptor, 1685 + + matchAndRewrite(cf::CondBranchOp op, cf::CondBranchOp::Adaptor adaptor, 1686 + ConversionPatternRewriter &rewriter) const override { 1687 + auto converter = getTypeConverter(); 1688 + - auto newOp = rewriter.replaceOpWithNewOp<CondBranchOp>( 1689 + + auto newOp = rewriter.replaceOpWithNewOp<cf::CondBranchOp>( 1690 + op, adaptor.getCondition(), op.getTrueDest(), 1691 + adaptor.getTrueDestOperands(), op.getFalseDest(), 1692 + adaptor.getFalseDestOperands()); 1693 + diff --git a/lib/Dialect/Triton/IR/CMakeLists.txt b/lib/Dialect/Triton/IR/CMakeLists.txt 1694 + index 2d679b21fd..705554ba6b 100644 1695 + --- a/lib/Dialect/Triton/IR/CMakeLists.txt 1696 + +++ b/lib/Dialect/Triton/IR/CMakeLists.txt 1697 + @@ -10,11 +10,7 @@ add_mlir_dialect_library(TritonIR 1698 + 1699 + LINK_LIBS PUBLIC 1700 + MLIRIR 1701 + - MLIRArithmetic 1702 + - MLIRSCF 1703 + - 1704 + - # Since LLVM 15 1705 + - # MLIRFunc 1706 + - # else 1707 + - MLIRStandard 1708 + + MLIRArithmeticDialect 1709 + + MLIRSCFDialect 1710 + + MLIRFuncDialect 1711 + ) 1712 + diff --git a/lib/Dialect/Triton/IR/Ops.cpp b/lib/Dialect/Triton/IR/Ops.cpp 1713 + index 3aadbfa0c0..86570359c5 100644 1714 + --- a/lib/Dialect/Triton/IR/Ops.cpp 1715 + +++ b/lib/Dialect/Triton/IR/Ops.cpp 1716 + @@ -1,10 +1,9 @@ 1717 + -#include "triton/Dialect/Triton/IR/Dialect.h" 1718 + -#include "triton/Dialect/Triton/IR/Types.h" 1719 + - 1720 + #include "mlir/IR/Builders.h" 1721 + #include "mlir/IR/BuiltinAttributes.h" 1722 + #include "mlir/IR/BuiltinTypes.h" 1723 + #include "mlir/IR/OperationSupport.h" 1724 + +#include "triton/Dialect/Triton/IR/Dialect.h" 1725 + +#include "triton/Dialect/Triton/IR/Types.h" 1726 + 1727 + namespace mlir { 1728 + namespace triton { 1729 + @@ -38,8 +37,8 @@ static Type getPointerTypeSameShape(Type type) { 1730 + } 1731 + 1732 + // Parser & printer for assembly forms 1733 + -ParseResult parseLoadOp(OpAsmParser &parser, OperationState &result) { 1734 + - SmallVector<OpAsmParser::OperandType, 4> allOperands; 1735 + +ParseResult LoadOp::parse(OpAsmParser &parser, OperationState &result) { 1736 + + SmallVector<OpAsmParser::UnresolvedOperand, 4> allOperands; 1737 + Type resultTypes[1]; 1738 + SMLoc allOperandLoc = parser.getCurrentLocation(); 1739 + if (parser.parseOperandList(allOperands) || 1740 + @@ -73,18 +72,18 @@ ParseResult parseLoadOp(OpAsmParser &parser, OperationState &result) { 1741 + return success(); 1742 + } 1743 + 1744 + -void printLoadOp(OpAsmPrinter &printer, LoadOp loadOp) { 1745 + +void LoadOp::print(OpAsmPrinter &printer) { 1746 + printer << " "; 1747 + - printer << loadOp.getOperation()->getOperands(); 1748 + + printer << getOperation()->getOperands(); 1749 + // "operand_segment_sizes" can be deduced, so we don't print it. 1750 + - printer.printOptionalAttrDict(loadOp->getAttrs(), 1751 + - {loadOp.operand_segment_sizesAttrName()}); 1752 + + printer.printOptionalAttrDict(getOperation()->getAttrs(), 1753 + + {operand_segment_sizesAttrName()}); 1754 + printer << " : "; 1755 + - printer.printStrippedAttrOrType(loadOp.result().getType()); 1756 + + printer.printStrippedAttrOrType(getResult().getType()); 1757 + } 1758 + 1759 + -ParseResult parseStoreOp(OpAsmParser &parser, OperationState &result) { 1760 + - SmallVector<OpAsmParser::OperandType, 4> allOperands; 1761 + +ParseResult StoreOp::parse(OpAsmParser &parser, OperationState &result) { 1762 + + SmallVector<OpAsmParser::UnresolvedOperand, 4> allOperands; 1763 + Type valueType; 1764 + SMLoc allOperandLoc = parser.getCurrentLocation(); 1765 + if (parser.parseOperandList(allOperands) || 1766 + @@ -104,12 +103,12 @@ ParseResult parseStoreOp(OpAsmParser &parser, OperationState &result) { 1767 + return success(); 1768 + } 1769 + 1770 + -void printStoreOp(OpAsmPrinter &printer, StoreOp storeOp) { 1771 + +void StoreOp::print(OpAsmPrinter &printer) { 1772 + printer << " "; 1773 + - printer << storeOp.getOperation()->getOperands(); 1774 + - printer.printOptionalAttrDict(storeOp->getAttrs(), /*elidedAttrs=*/{}); 1775 + + printer << getOperation()->getOperands(); 1776 + + printer.printOptionalAttrDict(getOperation()->getAttrs(), /*elidedAttrs=*/{}); 1777 + printer << " : "; 1778 + - printer.printStrippedAttrOrType(storeOp.value().getType()); 1779 + + printer.printStrippedAttrOrType(value().getType()); 1780 + } 1781 + 1782 + } // namespace triton 1783 + @@ -319,7 +318,8 @@ OpFoldResult SplatOp::fold(ArrayRef<Attribute> operands) { 1784 + if (!constOperand) 1785 + return {}; 1786 + auto shapedType = getType().cast<ShapedType>(); 1787 + - auto ret = SplatElementsAttr::get(shapedType, {constOperand.getValue()}); 1788 + + auto ret = SplatElementsAttr::get( 1789 + + shapedType, ArrayRef<Attribute>(constOperand.getValue())); 1790 + return ret; 1791 + } 1792 + 1793 + diff --git a/lib/Dialect/Triton/Transforms/Combine.cpp b/lib/Dialect/Triton/Transforms/Combine.cpp 1794 + index 2261472170..11570283d6 100644 1795 + --- a/lib/Dialect/Triton/Transforms/Combine.cpp 1796 + +++ b/lib/Dialect/Triton/Transforms/Combine.cpp 1797 + @@ -57,13 +57,13 @@ DenseElementsAttr getConstantValue(Builder &builder, Attribute value, 1798 + class CombineSelectMaskedLoadPattern : public mlir::RewritePattern { 1799 + public: 1800 + CombineSelectMaskedLoadPattern(mlir::MLIRContext *context) 1801 + - : mlir::RewritePattern(mlir::SelectOp::getOperationName(), 3, context, 1802 + - {triton::LoadOp::getOperationName()}) {} 1803 + + : mlir::RewritePattern(mlir::arith::SelectOp::getOperationName(), 3, 1804 + + context, {triton::LoadOp::getOperationName()}) {} 1805 + 1806 + mlir::LogicalResult 1807 + matchAndRewrite(mlir::Operation *op, 1808 + mlir::PatternRewriter &rewriter) const override { 1809 + - auto selectOp = llvm::dyn_cast<mlir::SelectOp>(op); 1810 + + auto selectOp = llvm::dyn_cast<mlir::arith::SelectOp>(op); 1811 + if (!selectOp) 1812 + return mlir::failure(); 1813 + 1814 + diff --git a/lib/Dialect/Triton/Transforms/Combine.td b/lib/Dialect/Triton/Transforms/Combine.td 1815 + index 14f286b26e..ded0e346e6 100644 1816 + --- a/lib/Dialect/Triton/Transforms/Combine.td 1817 + +++ b/lib/Dialect/Triton/Transforms/Combine.td 1818 + @@ -1,9 +1,9 @@ 1819 + #ifndef TRITON_PATTERNS 1820 + #define TRITON_PATTERNS 1821 + 1822 + -include "mlir/Dialect/StandardOps/IR/Ops.td" 1823 + include "mlir/Dialect/Arithmetic/IR/ArithmeticOps.td" 1824 + include "triton/Dialect/Triton/IR/TritonOps.td" 1825 + +include "mlir/IR/PatternBase.td" 1826 + 1827 + 1828 + // AddIOp(DotOp(a, b, c), d) and c==0 => DotOp(a, b, d) 1829 + diff --git a/lib/Dialect/TritonGPU/IR/Dialect.cpp b/lib/Dialect/TritonGPU/IR/Dialect.cpp 1830 + index 1fbc609e88..bfc3f3d3da 100644 1831 + --- a/lib/Dialect/TritonGPU/IR/Dialect.cpp 1832 + +++ b/lib/Dialect/TritonGPU/IR/Dialect.cpp 1833 + @@ -1,14 +1,14 @@ 1834 + +#include "triton/Dialect/Triton/IR/Dialect.h" 1835 + + 1836 + #include <numeric> 1837 + 1838 + #include "mlir/IR/DialectImplementation.h" 1839 + #include "mlir/IR/OpImplementation.h" 1840 + #include "triton/Analysis/Utility.h" 1841 + -#include "triton/Dialect/Triton/IR/Dialect.h" 1842 + +#include "triton/Dialect/TritonGPU/IR/Dialect.cpp.inc" 1843 + #include "triton/Dialect/TritonGPU/IR/Dialect.h" 1844 + #include "llvm/ADT/TypeSwitch.h" 1845 + 1846 + -#include "triton/Dialect/TritonGPU/IR/Dialect.cpp.inc" 1847 + - 1848 + using namespace mlir; 1849 + using namespace mlir::triton::gpu; 1850 + 1851 + @@ -366,7 +366,6 @@ template SmallVector<int64_t> 1852 + SliceEncodingAttr::paddedShape<int64_t>(ArrayRef<int64_t> shape) const; 1853 + 1854 + unsigned SliceEncodingAttr::getElemsPerThread(ArrayRef<int64_t> shape) const { 1855 + - size_t rank = shape.size(); 1856 + auto parent = getParent(); 1857 + return ::getElemsPerThread(parent, paddedShape(shape)); 1858 + } 1859 + @@ -655,9 +654,9 @@ void DotOperandEncodingAttr::print(mlir::AsmPrinter &printer) const { 1860 + // InsertSliceAsyncOp 1861 + //===----------------------------------------------------------------------===// 1862 + 1863 + -ParseResult parseInsertSliceAsyncOp(OpAsmParser &parser, 1864 + - OperationState &result) { 1865 + - SmallVector<OpAsmParser::OperandType, 8> allOperands; 1866 + +ParseResult InsertSliceAsyncOp::parse(OpAsmParser &parser, 1867 + + OperationState &result) { 1868 + + SmallVector<OpAsmParser::UnresolvedOperand, 8> allOperands; 1869 + Type srcType, dstType; 1870 + SMLoc allOperandLoc = parser.getCurrentLocation(); 1871 + if (parser.parseOperandList(allOperands) || 1872 + @@ -696,18 +695,16 @@ ParseResult parseInsertSliceAsyncOp(OpAsmParser &parser, 1873 + return success(); 1874 + } 1875 + 1876 + -void printInsertSliceAsyncOp(OpAsmPrinter &printer, 1877 + - InsertSliceAsyncOp insertSliceAsyncOp) { 1878 + +void InsertSliceAsyncOp::print(OpAsmPrinter &printer) { 1879 + printer << " "; 1880 + - printer << insertSliceAsyncOp.getOperation()->getOperands(); 1881 + + printer << getOperation()->getOperands(); 1882 + // "operand_segment_sizes" can be deduced, so we don't print it. 1883 + - printer.printOptionalAttrDict( 1884 + - insertSliceAsyncOp->getAttrs(), 1885 + - {insertSliceAsyncOp.operand_segment_sizesAttrName()}); 1886 + + printer.printOptionalAttrDict(getOperation()->getAttrs(), 1887 + + {operand_segment_sizesAttrName()}); 1888 + printer << " : "; 1889 + - printer.printStrippedAttrOrType(insertSliceAsyncOp.src().getType()); 1890 + + printer.printStrippedAttrOrType(src().getType()); 1891 + printer << " -> "; 1892 + - printer.printStrippedAttrOrType(insertSliceAsyncOp.result().getType()); 1893 + + printer.printStrippedAttrOrType(result().getType()); 1894 + } 1895 + 1896 + //===----------------------------------------------------------------------===// 1897 + diff --git a/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp b/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp 1898 + index 82407980d3..ee6009f44a 100644 1899 + --- a/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp 1900 + +++ b/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp 1901 + @@ -27,7 +27,11 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> { 1902 + auto origType = ptr.getType().cast<RankedTensorType>(); 1903 + // Get the shape of the tensor. 1904 + size_t rank = origType.getRank(); 1905 + - AxisInfo info = axisInfo.lookupLatticeElement(ptr)->getValue(); 1906 + + dataflow::Lattice<AxisInfo> *latticeElement = 1907 + + axisInfo.getLatticeElement(ptr); 1908 + + AxisInfo info = latticeElement && !latticeElement->isUninitialized() 1909 + + ? latticeElement->getValue() 1910 + + : AxisInfo(); 1911 + // Get the contiguity order of `ptr` 1912 + auto order = argSort(info.getContiguity()); 1913 + // The desired divisibility is the maximum divisibility 1914 + @@ -40,7 +44,7 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> { 1915 + for (Value val : op->getResults()) { 1916 + if (val.getType() != origType) 1917 + continue; 1918 + - auto valInfo = axisInfo.lookupLatticeElement(val); 1919 + + auto valInfo = axisInfo.getLatticeElement(val); 1920 + auto currOrder = argSort(valInfo->getValue().getContiguity()); 1921 + if (order == currOrder) 1922 + withSameOrder.insert(val); 1923 + @@ -55,7 +59,7 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> { 1924 + unsigned elemNumBytes = std::max(elemNumBits / 8, 1u); 1925 + unsigned perThread = 1; 1926 + for (Value val : withSameOrder) { 1927 + - AxisInfo info = axisInfo.lookupLatticeElement(val)->getValue(); 1928 + + AxisInfo info = axisInfo.getLatticeElement(val)->getValue(); 1929 + unsigned maxMultipleBytes = info.getDivisibility(order[0]); 1930 + unsigned maxMultiple = std::max(maxMultipleBytes / elemNumBytes, 1u); 1931 + unsigned maxContig = info.getContiguity(order[0]); 1932 + @@ -123,8 +127,10 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> { 1933 + void runOnOperation() override { 1934 + Operation *op = getOperation(); 1935 + // Run axis info analysis 1936 + - AxisInfoAnalysis axisInfo(&getContext()); 1937 + - axisInfo.run(op); 1938 + + std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver(); 1939 + + AxisInfoAnalysis *axisInfo = solver->load<AxisInfoAnalysis>(); 1940 + + if (failed(solver->initializeAndRun(op))) 1941 + + return signalPassFailure(); 1942 + 1943 + // For each i/o operation, we determine what layout 1944 + // the pointers should have for best memory coalescing 1945 + @@ -146,10 +152,10 @@ struct CoalescePass : public TritonGPUCoalesceBase<CoalescePass> { 1946 + RankedTensorType ty = ptr.getType().template dyn_cast<RankedTensorType>(); 1947 + if (!ty || !ty.getElementType().isa<PointerType>()) 1948 + return; 1949 + - AxisInfo info = axisInfo.lookupLatticeElement(ptr)->getValue(); 1950 + + AxisInfo info = axisInfo->getLatticeElement(ptr)->getValue(); 1951 + auto mod = curr->getParentOfType<ModuleOp>(); 1952 + int numWarps = triton::gpu::TritonGPUDialect::getNumWarps(mod); 1953 + - auto convertType = getTypeConverter(axisInfo, ptr, numWarps); 1954 + + auto convertType = getTypeConverter(*axisInfo, ptr, numWarps); 1955 + layoutMap[ptr] = convertType; 1956 + }); 1957 + 1958 + diff --git a/lib/Dialect/TritonGPU/Transforms/Combine.cpp b/lib/Dialect/TritonGPU/Transforms/Combine.cpp 1959 + index efa37ff2dc..089ce3996c 100644 1960 + --- a/lib/Dialect/TritonGPU/Transforms/Combine.cpp 1961 + +++ b/lib/Dialect/TritonGPU/Transforms/Combine.cpp 1962 + @@ -1,6 +1,6 @@ 1963 + #include "Utility.h" 1964 + #include "mlir/Analysis/SliceAnalysis.h" 1965 + -#include "mlir/Dialect/SCF/SCF.h" 1966 + +#include "mlir/Dialect/SCF/IR/SCF.h" 1967 + #include "mlir/IR/BlockAndValueMapping.h" 1968 + #include "mlir/IR/BuiltinAttributes.h" 1969 + #include "mlir/IR/Matchers.h" 1970 + diff --git a/lib/Dialect/TritonGPU/Transforms/Combine.td b/lib/Dialect/TritonGPU/Transforms/Combine.td 1971 + index 6bf1b14866..6a7b10dbcb 100644 1972 + --- a/lib/Dialect/TritonGPU/Transforms/Combine.td 1973 + +++ b/lib/Dialect/TritonGPU/Transforms/Combine.td 1974 + @@ -3,5 +3,6 @@ 1975 + 1976 + include "triton/Dialect/TritonGPU/IR/TritonGPUOps.td" 1977 + include "triton/Dialect/Triton/IR/TritonOps.td" 1978 + +include "mlir/IR/PatternBase.td" 1979 + 1980 + #endif 1981 + diff --git a/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp b/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp 1982 + index 4bd3bc76bf..b2f8defd81 100644 1983 + --- a/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp 1984 + +++ b/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp 1985 + @@ -1,5 +1,5 @@ 1986 + #include "mlir/Analysis/SliceAnalysis.h" 1987 + -#include "mlir/Dialect/SCF/SCF.h" 1988 + +#include "mlir/Dialect/SCF/IR/SCF.h" 1989 + #include "mlir/IR/BlockAndValueMapping.h" 1990 + #include "mlir/IR/BuiltinAttributes.h" 1991 + #include "mlir/IR/Matchers.h" 1992 + diff --git a/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp b/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp 1993 + index 9b2f42231e..85f746c1dc 100644 1994 + --- a/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp 1995 + +++ b/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp 1996 + @@ -2,6 +2,7 @@ 1997 + #include "mlir/IR/BlockAndValueMapping.h" 1998 + #include "mlir/IR/TypeUtilities.h" 1999 + #include "triton/Analysis/AxisInfo.h" 2000 + +#include "triton/Analysis/Utility.h" 2001 + #include "triton/Dialect/TritonGPU/IR/Dialect.h" 2002 + #include "triton/Dialect/TritonGPU/Transforms/Passes.h" 2003 + 2004 + @@ -160,15 +161,18 @@ ttg::AllocTensorOp LoopPipeliner::allocateEmptyBuffer(Operation *op, 2005 + LogicalResult LoopPipeliner::initialize() { 2006 + Block *loop = forOp.getBody(); 2007 + 2008 + - AxisInfoAnalysis axisInfoAnalysis(forOp.getContext()); 2009 + - axisInfoAnalysis.run(forOp->getParentOfType<ModuleOp>()); 2010 + + std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver(); 2011 + + AxisInfoAnalysis *axisInfoAnalysis = solver->load<AxisInfoAnalysis>(); 2012 + + if (failed(solver->initializeAndRun(forOp->getParentOfType<ModuleOp>()))) { 2013 + + return failure(); 2014 + + } 2015 + 2016 + // can we use forOp.walk(...) here? 2017 + SmallVector<triton::LoadOp, 2> allLoads; 2018 + for (Operation &op : *loop) 2019 + if (auto loadOp = dyn_cast<triton::LoadOp>(&op)) { 2020 + auto ptr = loadOp.ptr(); 2021 + - unsigned vec = axisInfoAnalysis.getPtrContiguity(ptr); 2022 + + unsigned vec = axisInfoAnalysis->getPtrContiguity(ptr); 2023 + auto tensorTy = ptr.getType().dyn_cast<RankedTensorType>(); 2024 + if (!tensorTy) 2025 + continue; 2026 + diff --git a/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp b/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp 2027 + index 0e7dbe5264..b95a4f50a6 100644 2028 + --- a/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp 2029 + +++ b/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp 2030 + @@ -1,5 +1,5 @@ 2031 + #include "mlir/Analysis/SliceAnalysis.h" 2032 + -#include "mlir/Dialect/SCF/SCF.h" 2033 + +#include "mlir/Dialect/SCF/IR/SCF.h" 2034 + #include "mlir/IR/BlockAndValueMapping.h" 2035 + #include "mlir/IR/BuiltinAttributes.h" 2036 + #include "mlir/IR/Matchers.h" 2037 + diff --git a/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp b/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp 2038 + index 37ac710995..762e887f36 100644 2039 + --- a/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp 2040 + +++ b/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp 2041 + @@ -82,12 +82,12 @@ TritonGPUConversionTarget::TritonGPUConversionTarget( 2042 + scf::ReduceReturnOp>(); 2043 + 2044 + addDynamicallyLegalDialect<arith::ArithmeticDialect, math::MathDialect, 2045 + - triton::TritonDialect, StandardOpsDialect, 2046 + - scf::SCFDialect>([&](Operation *op) { 2047 + - if (typeConverter.isLegal(op)) 2048 + - return true; 2049 + - return false; 2050 + - }); 2051 + + triton::TritonDialect, scf::SCFDialect>( 2052 + + [&](Operation *op) { 2053 + + if (typeConverter.isLegal(op)) 2054 + + return true; 2055 + + return false; 2056 + + }); 2057 + 2058 + // We have requirements for the data layouts 2059 + addDynamicallyLegalOp<triton::DotOp>([](triton::DotOp dotOp) -> bool { 2060 + diff --git a/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp b/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp 2061 + index c229104286..c911fd4a5c 100644 2062 + --- a/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp 2063 + +++ b/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp 2064 + @@ -1,5 +1,5 @@ 2065 + #include "Utility.h" 2066 + -#include "mlir/Dialect/SCF/SCF.h" 2067 + +#include "mlir/Dialect/SCF/IR/SCF.h" 2068 + #include "mlir/IR/Matchers.h" 2069 + #include "mlir/IR/PatternMatch.h" 2070 + #include "mlir/Transforms/GreedyPatternRewriteDriver.h" 2071 + @@ -118,8 +118,8 @@ void setOpResultType(Operation *op, ArrayRef<Type> newTypes) { 2072 + .get("value") 2073 + .dyn_cast<mlir::DenseElementsAttr>(); 2074 + if (attr) { 2075 + - auto newAttr = mlir::DenseElementsAttr::getFromRawBuffer( 2076 + - newType, attr.getRawData(), true); 2077 + + auto newAttr = 2078 + + mlir::DenseElementsAttr::getFromRawBuffer(newType, attr.getRawData()); 2079 + op->setAttr("value", newAttr); 2080 + } 2081 + } 2082 + diff --git a/lib/Dialect/TritonGPU/Transforms/Utility.cpp b/lib/Dialect/TritonGPU/Transforms/Utility.cpp 2083 + index ed15f02f67..6400f1633a 100644 2084 + --- a/lib/Dialect/TritonGPU/Transforms/Utility.cpp 2085 + +++ b/lib/Dialect/TritonGPU/Transforms/Utility.cpp 2086 + @@ -1,5 +1,5 @@ 2087 + #include "Utility.h" 2088 + -#include "mlir/Dialect/SCF/SCF.h" 2089 + +#include "mlir/Dialect/SCF/IR/SCF.h" 2090 + #include "mlir/IR/BlockAndValueMapping.h" 2091 + #include "mlir/Transforms/GreedyPatternRewriteDriver.h" 2092 + 2093 + diff --git a/lib/Target/LLVMIR/CMakeLists.txt b/lib/Target/LLVMIR/CMakeLists.txt 2094 + index f1bbd0bf4e..ac8973ad19 100644 2095 + --- a/lib/Target/LLVMIR/CMakeLists.txt 2096 + +++ b/lib/Target/LLVMIR/CMakeLists.txt 2097 + @@ -6,8 +6,7 @@ add_mlir_translation_library(TritonLLVMIR 2098 + 2099 + LINK_LIBS PUBLIC 2100 + MLIRIR 2101 + - MLIRLLVMIR 2102 + - MLIRSCFToStandard 2103 + + MLIRLLVMDialect 2104 + MLIRSupport 2105 + MLIRTargetLLVMIRExport 2106 + ) 2107 + diff --git a/lib/Target/PTX/PTXTranslation.cpp b/lib/Target/PTX/PTXTranslation.cpp 2108 + index 4cb0d8193c..6a5453a6e7 100644 2109 + --- a/lib/Target/PTX/PTXTranslation.cpp 2110 + +++ b/lib/Target/PTX/PTXTranslation.cpp 2111 + @@ -1,11 +1,14 @@ 2112 + #include "triton/Target/PTX/PTXTranslation.h" 2113 + #include "triton/Target/LLVMIR/LLVMIRTranslation.h" 2114 + +#include <optional> 2115 + 2116 + #include "llvm/IR/IRBuilder.h" 2117 + #include "llvm/IR/LegacyPassManager.h" 2118 + #include "llvm/IR/Module.h" 2119 + #include "llvm/IR/Verifier.h" 2120 + #include "llvm/MC/TargetRegistry.h" 2121 + +#include "llvm/Pass.h" 2122 + +#include "llvm/Support/CommandLine.h" 2123 + #include "llvm/Support/TargetSelect.h" 2124 + #include "llvm/Target/TargetMachine.h" 2125 + 2126 + diff --git a/python/setup.py b/python/setup.py 2127 + index 2ac3accd25..4530b36714 100644 2128 + --- a/python/setup.py 2129 + +++ b/python/setup.py 2130 + @@ -57,19 +57,10 @@ def get_pybind11_package_info(): 2131 + def get_llvm_package_info(): 2132 + # download if nothing is installed 2133 + system = platform.system() 2134 + - if system == "Darwin": 2135 + - system_suffix = "apple-darwin" 2136 + - elif system == "Linux": 2137 + - vglibc = tuple(map(int, platform.libc_ver()[1].split('.'))) 2138 + - vglibc = vglibc[0] * 100 + vglibc[1] 2139 + - linux_suffix = 'ubuntu-18.04' if vglibc > 217 else 'centos-7' 2140 + - system_suffix = f"linux-gnu-{linux_suffix}" 2141 + - else: 2142 + - raise RuntimeError(f"unsupported system: {system}") 2143 + + system_suffix = {"Linux": "linux-gnu-ubuntu-18.04", "Darwin": "apple-darwin"}[system] 2144 + use_assert_enabled_llvm = check_env_flag("TRITON_USE_ASSERT_ENABLED_LLVM", "False") 2145 + - release_suffix = "assert" if use_assert_enabled_llvm else "release" 2146 + - name = f'llvm+mlir-14.0.6-x86_64-{system_suffix}-{release_suffix}' 2147 + - url = f"https://github.com/ptillet/triton-llvm-releases/releases/download/llvm-14.0.6-f28c006a5895/{name}.tar.xz" 2148 + + name = 'llvm+mlir-15.0.7-x86_64-{}-{}'.format(system_suffix, "assert" if use_assert_enabled_llvm else "release") 2149 + + url = "https://github.com/ptillet/triton-llvm-releases/releases/download/llvm-15.0.7-8dfdcc7b7bf6/{}.tar.xz".format(name) 2150 + return Package("llvm", name, url, "lib", "LLVM_INCLUDE_DIRS", "LLVM_LIBRARY_DIR", "LLVM_SYSPATH") 2151 + 2152 + 2153 + diff --git a/python/src/triton.cc b/python/src/triton.cc 2154 + index c40b117a55..f190eacc34 100644 2155 + --- a/python/src/triton.cc 2156 + +++ b/python/src/triton.cc 2157 + @@ -8,9 +8,10 @@ 2158 + #include "mlir/Pass/PassManager.h" 2159 + #include "mlir/Transforms/Passes.h" 2160 + 2161 + -#include "mlir/Parser.h" 2162 + +#include "mlir/Parser/Parser.h" 2163 + #include "mlir/Support/FileUtilities.h" 2164 + 2165 + +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" 2166 + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" 2167 + #include "triton/Analysis/Allocation.h" 2168 + #include "triton/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.h" 2169 + @@ -195,7 +196,7 @@ void init_triton_ir(py::module &&m) { 2170 + std::string attrName = name + "_arg" + std::to_string(id); 2171 + mlir::Block *owner = arg.getOwner(); 2172 + if (owner->isEntryBlock() && 2173 + - !mlir::isa<mlir::FuncOp>(owner->getParentOp())) { 2174 + + !mlir::isa<mlir::func::FuncOp>(owner->getParentOp())) { 2175 + owner->getParentOp()->setAttr(attrName, attr); 2176 + } 2177 + } 2178 + @@ -348,7 +349,7 @@ void init_triton_ir(py::module &&m) { 2179 + return str; 2180 + }) 2181 + .def("push_back", 2182 + - [](mlir::ModuleOp &self, mlir::FuncOp &funcOp) -> void { 2183 + + [](mlir::ModuleOp &self, mlir::func::FuncOp &funcOp) -> void { 2184 + self.push_back(funcOp); 2185 + }) 2186 + .def("has_function", 2187 + @@ -358,16 +359,18 @@ void init_triton_ir(py::module &&m) { 2188 + return false; 2189 + }) 2190 + .def("get_function", 2191 + - [](mlir::ModuleOp &self, std::string &funcName) -> mlir::FuncOp { 2192 + - return self.lookupSymbol<mlir::FuncOp>(funcName); 2193 + - }) 2194 + - .def("get_single_function", [](mlir::ModuleOp &self) -> mlir::FuncOp { 2195 + - llvm::SmallVector<mlir::FuncOp> funcs; 2196 + - self.walk([&](mlir::FuncOp func) { funcs.push_back(func); }); 2197 + - if (funcs.size() != 1) 2198 + - throw std::runtime_error("Expected a single function"); 2199 + - return funcs[0]; 2200 + - }); 2201 + + [](mlir::ModuleOp &self, 2202 + + std::string &funcName) -> mlir::func::FuncOp { 2203 + + return self.lookupSymbol<mlir::func::FuncOp>(funcName); 2204 + + }) 2205 + + .def("get_single_function", 2206 + + [](mlir::ModuleOp &self) -> mlir::func::FuncOp { 2207 + + llvm::SmallVector<mlir::func::FuncOp> funcs; 2208 + + self.walk([&](mlir::func::FuncOp func) { funcs.push_back(func); }); 2209 + + if (funcs.size() != 1) 2210 + + throw std::runtime_error("Expected a single function"); 2211 + + return funcs[0]; 2212 + + }); 2213 + 2214 + m.def("make_attr", 2215 + [](const std::vector<int> &values, mlir::MLIRContext &context) { 2216 + @@ -388,47 +391,48 @@ void init_triton_ir(py::module &&m) { 2217 + registry.insert<mlir::triton::TritonDialect, 2218 + mlir::triton::gpu::TritonGPUDialect, 2219 + mlir::math::MathDialect, mlir::arith::ArithmeticDialect, 2220 + - mlir::StandardOpsDialect, mlir::scf::SCFDialect>(); 2221 + + mlir::func::FuncDialect, mlir::scf::SCFDialect>(); 2222 + context.appendDialectRegistry(registry); 2223 + context.loadAllAvailableDialects(); 2224 + 2225 + // parse module 2226 + - mlir::OwningOpRef<mlir::ModuleOp> module( 2227 + - mlir::parseSourceFile(inputFilename, &context)); 2228 + + mlir::OwningOpRef<mlir::ModuleOp> module = 2229 + + mlir::parseSourceFile<mlir::ModuleOp>(inputFilename, &context); 2230 + + if (!module) 2231 + + throw std::runtime_error("Parse MLIR file failed."); 2232 + // locations are incompatible with ptx < 7.5 ! 2233 + module->walk([](mlir::Operation *op) { 2234 + op->setLoc(mlir::UnknownLoc::get(op->getContext())); 2235 + }); 2236 + - if (!module) 2237 + - throw std::runtime_error("Parse MLIR file failed."); 2238 + 2239 + return module->clone(); 2240 + }, 2241 + ret::take_ownership); 2242 + 2243 + - py::class_<mlir::FuncOp, mlir::OpState>(m, "function") 2244 + + py::class_<mlir::func::FuncOp, mlir::OpState>(m, "function") 2245 + // .def_property_readonly("attrs", &ir::function::attrs) 2246 + // .def("add_attr", &ir::function::add_attr); 2247 + .def("args", 2248 + - [](mlir::FuncOp &self, unsigned idx) -> mlir::BlockArgument { 2249 + + [](mlir::func::FuncOp &self, unsigned idx) -> mlir::BlockArgument { 2250 + return self.getArgument(idx); 2251 + }) 2252 + .def( 2253 + "add_entry_block", 2254 + - [](mlir::FuncOp &self) -> mlir::Block * { 2255 + + [](mlir::func::FuncOp &self) -> mlir::Block * { 2256 + return self.addEntryBlock(); 2257 + }, 2258 + ret::reference) 2259 + .def( 2260 + "set_arg_attr", 2261 + - [](mlir::FuncOp &self, int arg_no, const std::string &name, int val) { 2262 + + [](mlir::func::FuncOp &self, int arg_no, const std::string &name, 2263 + + int val) { 2264 + // set arg attributes "name" to value "val" 2265 + auto attrTy = mlir::IntegerType::get(self.getContext(), 32); 2266 + self.setArgAttr(arg_no, name, mlir::IntegerAttr::get(attrTy, val)); 2267 + }, 2268 + ret::reference) 2269 + - .def_property_readonly("type", &mlir::FuncOp::getType) 2270 + - .def("reset_type", &mlir::FuncOp::setType); 2271 + + .def_property_readonly("type", &mlir::func::FuncOp::getFunctionType) 2272 + + .def("reset_type", &mlir::func::FuncOp::setType); 2273 + 2274 + py::class_<mlir::OpBuilder::InsertPoint>(m, "InsertPoint"); 2275 + 2276 + @@ -445,13 +449,13 @@ void init_triton_ir(py::module &&m) { 2277 + .def("ret", 2278 + [](mlir::OpBuilder &self, std::vector<mlir::Value> &vals) -> void { 2279 + auto loc = self.getUnknownLoc(); 2280 + - self.create<mlir::ReturnOp>(loc, vals); 2281 + + self.create<mlir::func::ReturnOp>(loc, vals); 2282 + }) 2283 + .def("call", 2284 + - [](mlir::OpBuilder &self, mlir::FuncOp &func, 2285 + + [](mlir::OpBuilder &self, mlir::func::FuncOp &func, 2286 + std::vector<mlir::Value> &args) -> mlir::OpState { 2287 + auto loc = self.getUnknownLoc(); 2288 + - return self.create<mlir::CallOp>(loc, func, args); 2289 + + return self.create<mlir::func::CallOp>(loc, func, args); 2290 + }) 2291 + // insertion block/point 2292 + .def("set_insertion_point_to_start", 2293 + @@ -618,15 +622,16 @@ void init_triton_ir(py::module &&m) { 2294 + .def("get_or_insert_function", 2295 + [](mlir::OpBuilder &self, mlir::ModuleOp &module, 2296 + std::string &funcName, mlir::Type &funcType, 2297 + - std::string &visibility) -> mlir::FuncOp { 2298 + + std::string &visibility) -> mlir::func::FuncOp { 2299 + if (mlir::Operation *funcOperation = module.lookupSymbol(funcName)) 2300 + - return llvm::dyn_cast<mlir::FuncOp>(funcOperation); 2301 + + return llvm::dyn_cast<mlir::func::FuncOp>(funcOperation); 2302 + auto loc = self.getUnknownLoc(); 2303 + if (auto funcTy = funcType.dyn_cast<mlir::FunctionType>()) { 2304 + llvm::SmallVector<mlir::NamedAttribute> attrs = { 2305 + mlir::NamedAttribute(self.getStringAttr("sym_visibility"), 2306 + self.getStringAttr(visibility))}; 2307 + - return self.create<mlir::FuncOp>(loc, funcName, funcTy, attrs); 2308 + + return self.create<mlir::func::FuncOp>(loc, funcName, funcTy, 2309 + + attrs); 2310 + } 2311 + throw std::runtime_error("invalid function type"); 2312 + }) 2313 + @@ -658,15 +663,15 @@ void init_triton_ir(py::module &&m) { 2314 + [](mlir::OpBuilder &self, mlir::Value condition, 2315 + mlir::Block *trueDest, mlir::Block *falseDest) { 2316 + auto loc = self.getUnknownLoc(); 2317 + - self.create<mlir::CondBranchOp>(loc, condition, trueDest, 2318 + - falseDest); 2319 + + self.create<mlir::cf::CondBranchOp>(loc, condition, trueDest, 2320 + + falseDest); 2321 + return; 2322 + }) 2323 + .def("create_branch", 2324 + [](mlir::OpBuilder &self, mlir::Block *dest, 2325 + std::vector<mlir::Value> &args) { 2326 + auto loc = self.getUnknownLoc(); 2327 + - self.create<mlir::BranchOp>(loc, dest, args); 2328 + + self.create<mlir::cf::BranchOp>(loc, dest, args); 2329 + return; 2330 + }) 2331 + // Structured control flow 2332 + @@ -792,14 +797,14 @@ void init_triton_ir(py::module &&m) { 2333 + .def("create_to_index", 2334 + [](mlir::OpBuilder &self, mlir::Value &input) -> mlir::Value { 2335 + auto loc = self.getUnknownLoc(); 2336 + - return self.create<mlir::arith::IndexCastOp>(loc, input, 2337 + - self.getIndexType()); 2338 + + return self.create<mlir::arith::IndexCastOp>( 2339 + + loc, self.getIndexType(), input); 2340 + }) 2341 + .def("create_index_to_si", 2342 + [](mlir::OpBuilder &self, mlir::Value &input) -> mlir::Value { 2343 + auto loc = self.getUnknownLoc(); 2344 + - return self.create<mlir::arith::IndexCastOp>(loc, input, 2345 + - self.getI32Type()); 2346 + + return self.create<mlir::arith::IndexCastOp>( 2347 + + loc, self.getI32Type(), input); 2348 + }) 2349 + .def("create_fmul", 2350 + [](mlir::OpBuilder &self, mlir::Value &lhs, 2351 + @@ -1316,8 +1321,8 @@ void init_triton_ir(py::module &&m) { 2352 + [](mlir::OpBuilder &self, mlir::Value &condition, 2353 + mlir::Value &trueValue, mlir::Value &falseValue) -> mlir::Value { 2354 + auto loc = self.getUnknownLoc(); 2355 + - return self.create<mlir::SelectOp>(loc, condition, trueValue, 2356 + - falseValue); 2357 + + return self.create<mlir::arith::SelectOp>(loc, condition, 2358 + + trueValue, falseValue); 2359 + }) 2360 + .def("create_printf", 2361 + [](mlir::OpBuilder &self, const std::string &prefix, 2362 + @@ -1429,7 +1434,7 @@ void init_triton_ir(py::module &&m) { 2363 + self.addPass(mlir::triton::createConvertTritonGPUToLLVMPass()); 2364 + }) 2365 + .def("add_scf_to_cfg", [](mlir::PassManager &self) { 2366 + - self.addPass(mlir::createLowerToCFGPass()); 2367 + + self.addPass(mlir::createConvertSCFToCFPass()); 2368 + }); 2369 + } 2370 + 2371 + diff --git a/python/test/unit/language/test_core.py b/python/test/unit/language/test_core.py 2372 + index 432544a8a4..018f544714 100644 2373 + --- a/python/test/unit/language/test_core.py 2374 + +++ b/python/test/unit/language/test_core.py 2375 + @@ -1918,7 +1918,7 @@ def test_convert2d(dtype, shape, src_layout, dst_layout, device='cuda'): 2376 + #dst = {dst_layout} 2377 + """ + """ 2378 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 2379 + - func public @kernel_0d1d(%arg0: !tt.ptr<f16> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f16> {tt.divisibility = 16 : i32}) { 2380 + + func.func public @kernel_0d1d(%arg0: !tt.ptr<f16> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f16> {tt.divisibility = 16 : i32}) { 2381 + %cst = arith.constant dense<128> : tensor<128x1xi32, #src> 2382 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #triton_gpu.slice<{dim = 1, parent = #src}>> 2383 + %1 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #triton_gpu.slice<{dim = 0, parent = #src}>> 2384 + diff --git a/python/triton/compiler.py b/python/triton/compiler.py 2385 + index 5d167634df..c36589037c 100644 2386 + --- a/python/triton/compiler.py 2387 + +++ b/python/triton/compiler.py 2388 + @@ -1514,14 +1514,14 @@ def make_hash(fn, **kwargs): 2389 + return hashlib.md5((Path(fn).read_text() + triton.runtime.jit.version_key()).encode("utf-8")).hexdigest() 2390 + 2391 + 2392 + -# - ^\s*func\s+ : match the start of the string, any leading whitespace, the keyword func, 2393 + +# - ^\s*func\.func\s+ : match the start of the string, any leading whitespace, the keyword func, 2394 + # and any following whitespace 2395 + # - (public\s+)? : optionally match the keyword public and any following whitespace 2396 + # - (@\w+) : match an @ symbol followed by one or more word characters 2397 + # (letters, digits, or underscores), and capture it as group 1 (the function name) 2398 + # - (\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\)) : match a pair of parentheses enclosing 2399 + # zero or more arguments separated by commas, and capture it as group 2 (the argument list) 2400 + -mlir_prototype_pattern = r'^\s*func\s+(?:public\s+)?(@\w+)(\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\))\s*\{\s*$' 2401 + +mlir_prototype_pattern = r'^\s*func\.func\s+(?:public\s+)?(@\w+)(\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\))\s*\{\s*$' 2402 + ptx_prototype_pattern = r"\.(?:visible|extern)\s+\.(?:entry|func)\s+(\w+)\s*\(([^)]*)\)" 2403 + prototype_pattern = { 2404 + "ttir": mlir_prototype_pattern, 2405 + diff --git a/test/Analysis/test-alias.mlir b/test/Analysis/test-alias.mlir 2406 + index b3d5673f85..bb21615e68 100644 2407 + --- a/test/Analysis/test-alias.mlir 2408 + +++ b/test/Analysis/test-alias.mlir 2409 + @@ -11,7 +11,7 @@ 2410 + 2411 + // CHECK-LABEL: matmul_loop 2412 + // There shouldn't be any aliasing with the dot op encoding. 2413 + -func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 2414 + +func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 2415 + %a_ptr_init = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL> 2416 + %b_ptr_init = tt.broadcast %B : (!tt.ptr<f16>) -> tensor<32x128x!tt.ptr<f16>, #BL> 2417 + %a_mask = arith.constant dense<true> : tensor<128x32xi1, #AL> 2418 + @@ -36,7 +36,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B 2419 + } 2420 + 2421 + // CHECK-LABEL: alloc 2422 + -func @alloc(%A : !tt.ptr<f16>) { 2423 + +func.func @alloc(%A : !tt.ptr<f16>) { 2424 + // CHECK: %cst -> %cst 2425 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 2426 + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #AL> 2427 + @@ -46,7 +46,7 @@ func @alloc(%A : !tt.ptr<f16>) { 2428 + } 2429 + 2430 + // CHECK-LABEL: convert 2431 + -func @convert(%A : !tt.ptr<f16>) { 2432 + +func.func @convert(%A : !tt.ptr<f16>) { 2433 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> 2434 + // CHECK: %0 -> %0 2435 + %cst1 = triton_gpu.convert_layout %cst0 : (tensor<16x16xf16, #AL>) -> tensor<16x16xf16, #A_SHARED> 2436 + @@ -54,7 +54,7 @@ func @convert(%A : !tt.ptr<f16>) { 2437 + } 2438 + 2439 + // CHECK-LABEL: trans 2440 + -func @trans(%A : !tt.ptr<f16>) { 2441 + +func.func @trans(%A : !tt.ptr<f16>) { 2442 + // CHECK: %cst -> %cst 2443 + %tensor = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #A_SHARED> 2444 + // CHECK: %0 -> %cst 2445 + @@ -63,7 +63,7 @@ func @trans(%A : !tt.ptr<f16>) { 2446 + } 2447 + 2448 + // CHECK-LABEL: insert_slice_async 2449 + -func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) { 2450 + +func.func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) { 2451 + %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL> 2452 + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> 2453 + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> 2454 + @@ -76,7 +76,7 @@ func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) { 2455 + } 2456 + 2457 + // CHECK-LABEL: insert_slice 2458 + -func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) { 2459 + +func.func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) { 2460 + %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL> 2461 + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> 2462 + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> 2463 + @@ -90,7 +90,7 @@ func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) { 2464 + } 2465 + 2466 + // CHECK-LABEL: extract_slice 2467 + -func @extract_slice(%A : !tt.ptr<f16>) { 2468 + +func.func @extract_slice(%A : !tt.ptr<f16>) { 2469 + // CHECK: %cst -> %cst 2470 + %cst0 = arith.constant dense<0.000000e+00> : tensor<1x16x16xf16, #A_SHARED> 2471 + %index = arith.constant 0 : index 2472 + @@ -100,7 +100,7 @@ func @extract_slice(%A : !tt.ptr<f16>) { 2473 + } 2474 + 2475 + // CHECK-LABEL: if_cat 2476 + -func @if_cat(%i1 : i1) { 2477 + +func.func @if_cat(%i1 : i1) { 2478 + // CHECK: %cst -> %cst 2479 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 2480 + // CHECK: %cst_0 -> %cst_0 2481 + @@ -119,7 +119,7 @@ func @if_cat(%i1 : i1) { 2482 + } 2483 + 2484 + // CHECK-LABEL: if_alias 2485 + -func @if_alias(%i1 : i1) { 2486 + +func.func @if_alias(%i1 : i1) { 2487 + // CHECK: %cst -> %cst 2488 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 2489 + // CHECK-NEXT: %cst_0 -> %cst_0 2490 + @@ -134,7 +134,7 @@ func @if_alias(%i1 : i1) { 2491 + } 2492 + 2493 + // CHECK-LABEL: for 2494 + -func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 2495 + +func.func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 2496 + // CHECK: %cst -> %cst 2497 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 2498 + // CHECK-NEXT: %cst_0 -> %cst_0 2499 + @@ -154,7 +154,7 @@ func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.p 2500 + } 2501 + 2502 + // CHECK-LABEL: for_if 2503 + -func @for_if(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 2504 + +func.func @for_if(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 2505 + // CHECK: %cst -> %cst 2506 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 2507 + // CHECK-NEXT: %cst_0 -> %cst_0 2508 + @@ -180,7 +180,7 @@ func @for_if(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !t 2509 + } 2510 + 2511 + // CHECK-LABEL: for_if_for 2512 + -func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 2513 + +func.func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 2514 + // CHECK: %cst -> %cst 2515 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 2516 + // CHECK-NEXT: %cst_0 -> %cst_0 2517 + diff --git a/test/Analysis/test-alignment.mlir b/test/Analysis/test-alignment.mlir 2518 + index 0ab34c7a78..af8ea6f856 100644 2519 + --- a/test/Analysis/test-alignment.mlir 2520 + +++ b/test/Analysis/test-alignment.mlir 2521 + @@ -1,288 +1,288 @@ 2522 + -// RUN: triton-opt %s -test-print-alignment -split-input-file 2>&1 | FileCheck %s 2523 + +// RUN: triton-opt %s -test-print-alignment -split-input-file -o %t 2>&1 | FileCheck %s 2524 + 2525 + -// CHECK-LABEL: cast 2526 + -func @cast() { 2527 + - // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [1] 2528 + +// CHECK-LABEL: @cast 2529 + +func.func @cast() { 2530 + + // CHECK: contiguity = [1], divisibility = [1], constancy = [1], constant_value = 1 2531 + %cst = arith.constant 1 : i32 2532 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [1] 2533 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = 1 2534 + %0 = arith.extsi %cst : i32 to i64 2535 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] 2536 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 2537 + %cst_tensor = arith.constant dense<1> : tensor<128xi32> 2538 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] 2539 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 2540 + %1 = tt.bitcast %cst_tensor : tensor<128xi32> -> tensor<128xi64> 2541 + return 2542 + } 2543 + 2544 + // ----- 2545 + 2546 + -// CHECK-LABEL: add 2547 + -func @add() { 2548 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2549 + +// CHECK-LABEL: @add 2550 + +func.func @add() { 2551 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2552 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2553 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] 2554 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 2555 + %1 = arith.constant dense<1> : tensor<128xi32> 2556 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2557 + + // CHECK-NEXT: contiguity = [128], divisibility = [1], constancy = [1], constant_value = <none> 2558 + %2 = arith.addi %0, %1 : tensor<128xi32> 2559 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [127] 2560 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 127 2561 + %3 = arith.constant dense<127> : tensor<128xi32> 2562 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] 2563 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 2564 + %4 = arith.addi %1, %3 : tensor<128xi32> 2565 + return 2566 + } 2567 + 2568 + // ----- 2569 + 2570 + -// CHECK-LABEL: sub 2571 + -func @sub() { 2572 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2573 + +// CHECK-LABEL: @sub 2574 + +func.func @sub() { 2575 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2576 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2577 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] 2578 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 2579 + %1 = arith.constant dense<1> : tensor<128xi32> 2580 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2581 + + // CHECK-NEXT: contiguity = [128], divisibility = [1], constancy = [1], constant_value = <none> 2582 + %2 = arith.subi %0, %1 : tensor<128xi32> 2583 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [129] 2584 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 129 2585 + %3 = arith.constant dense<129> : tensor<128xi32> 2586 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] 2587 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 2588 + %4 = arith.subi %3, %1 : tensor<128xi32> 2589 + return 2590 + } 2591 + 2592 + // ----- 2593 + 2594 + -// CHECK-LABEL: mul 2595 + -func @mul() { 2596 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2597 + +// CHECK-LABEL: @mul 2598 + +func.func @mul() { 2599 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2600 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2601 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] 2602 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 2603 + %1 = arith.constant dense<1> : tensor<128xi32> 2604 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2605 + + // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2606 + %2 = arith.muli %0, %1 : tensor<128xi32> 2607 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] 2608 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 2609 + %3 = arith.constant dense<128> : tensor<128xi32> 2610 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] 2611 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 2612 + %4 = arith.muli %3, %1 : tensor<128xi32> 2613 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [2] ; Constancy: [128] ; ConstantValue: [2] 2614 + + // CHECK-NEXT: contiguity = [1], divisibility = [2], constancy = [128], constant_value = 2 2615 + %5 = arith.constant dense<2> : tensor<128xi32> 2616 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [256] ; Constancy: [128] ; ConstantValue: [256] 2617 + + // CHECK-NEXT: contiguity = [1], divisibility = [256], constancy = [128], constant_value = 256 2618 + %6 = arith.muli %4, %5 : tensor<128xi32> 2619 + return 2620 + } 2621 + 2622 + // ----- 2623 + 2624 + -// CHECK-LABEL: div 2625 + -func @div() { 2626 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2627 + +// CHECK-LABEL: @div 2628 + +func.func @div() { 2629 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2630 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2631 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] 2632 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 2633 + %1 = arith.constant dense<1> : tensor<128xi32> 2634 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2635 + + // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2636 + %2 = arith.divsi %0, %1 : tensor<128xi32> 2637 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2638 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 2639 + %3 = arith.divui %1, %0 : tensor<128xi32> 2640 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] 2641 + + // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 2642 + %4 = arith.constant dense<64> : tensor<128xi32> 2643 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16777216] ; Constancy: [64] ; ConstantValue: [None] 2644 + + // CHECK-NEXT: contiguity = [1], divisibility = [16777216], constancy = [64], constant_value = <none> 2645 + %5 = arith.divsi %0, %4 : tensor<128xi32> 2646 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2647 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 2648 + %6 = arith.divsi %4, %0 : tensor<128xi32> 2649 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] 2650 + + // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 2651 + %7 = arith.divsi %4, %1 : tensor<128xi32> 2652 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [2] ; Constancy: [128] ; ConstantValue: [66] 2653 + + // CHECK-NEXT: contiguity = [1], divisibility = [2], constancy = [128], constant_value = 66 2654 + %8 = arith.constant dense<66> : tensor<128xi32> 2655 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [2] ; ConstantValue: [None] 2656 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [2], constant_value = <none> 2657 + %9 = arith.divui %0, %8 : tensor<128xi32> 2658 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [8192] ; Constancy: [1] ; ConstantValue: [None] 2659 + + // CHECK-NEXT: contiguity = [128], divisibility = [8192], constancy = [1], constant_value = <none> 2660 + %10 = tt.make_range {end = 8320 : i32, start = 8192 : i32} : tensor<128xi32> 2661 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [64] ; ConstantValue: [None] 2662 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [64], constant_value = <none> 2663 + %11 = arith.divsi %10, %4 : tensor<128xi32> 2664 + - return 2665 + + return 2666 + } 2667 + 2668 + // ----- 2669 + 2670 + -// CHECK-LABEL: rem 2671 + -func @rem() { 2672 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2673 + +// CHECK-LABEL: @rem 2674 + +func.func @rem() { 2675 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2676 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2677 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] 2678 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 2679 + %1 = arith.constant dense<1> : tensor<128xi32> 2680 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] 2681 + + // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 2682 + %2 = arith.remsi %0, %1 : tensor<128xi32> 2683 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2684 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 2685 + %3 = arith.remui %1, %0 : tensor<128xi32> 2686 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] 2687 + + // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 2688 + %4 = arith.constant dense<64> : tensor<128xi32> 2689 + - // CHECK-NEXT: Contiguity: [64] ; Divisibility: [64] ; Constancy: [1] ; ConstantValue: [None] 2690 + + // CHECK-NEXT: contiguity = [64], divisibility = [64], constancy = [1], constant_value = <none> 2691 + %5 = arith.remsi %0, %4 : tensor<128xi32> 2692 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [1] ; ConstantValue: [None] 2693 + + // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [1], constant_value = <none> 2694 + %6 = arith.remsi %4, %0 : tensor<128xi32> 2695 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [2] ; Constancy: [128] ; ConstantValue: [66] 2696 + + // CHECK-NEXT: contiguity = [1], divisibility = [2], constancy = [128], constant_value = 66 2697 + %7 = arith.constant dense<66> : tensor<128xi32> 2698 + - // CHECK-NEXT: Contiguity: [2] ; Divisibility: [2] ; Constancy: [1] ; ConstantValue: [None] 2699 + + // CHECK-NEXT: contiguity = [2], divisibility = [2], constancy = [1], constant_value = <none> 2700 + %8 = arith.remui %0, %7 : tensor<128xi32> 2701 + - return 2702 + + return 2703 + } 2704 + 2705 + // ----- 2706 + 2707 + -// CHECK-LABEL: broadcast 2708 + -func @broadcast() { 2709 + - // CHECK: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] 2710 + +// CHECK-LABEL: @broadcast 2711 + +func.func @broadcast() { 2712 + + // CHECK: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 2713 + %0 = arith.constant dense<64> : tensor<128xi32> 2714 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [64, 1] ; Constancy: [128, 1] ; ConstantValue: [64] 2715 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [64, 1], constancy = [128, 1], constant_value = 64 2716 + %1 = tt.expand_dims %0 {axis = 1 : i32} : (tensor<128xi32>) -> tensor<128x1xi32> 2717 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [64, 1] ; Constancy: [128, 128] ; ConstantValue: [64] 2718 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [64, 1], constancy = [128, 128], constant_value = 64 2719 + %2 = tt.broadcast %1 : (tensor<128x1xi32>) -> tensor<128x128xi32> 2720 + return 2721 + } 2722 + 2723 + // ----- 2724 + 2725 + -// CHECK-LABEL: splat 2726 + -func @splat(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}) { 2727 + - // CHECK: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 128] ; ConstantValue: [None] 2728 + +// CHECK-LABEL: @splat 2729 + +func.func @splat(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}) { 2730 + + // CHECK: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 128], constant_value = <none> 2731 + %0 = tt.splat %arg0 : (!tt.ptr<f32>) -> tensor<128x128x!tt.ptr<f32>> 2732 + return 2733 + } 2734 + 2735 + // ----- 2736 + 2737 + -// CHECK-LABEL: cmp 2738 + -func @cmp() { 2739 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2740 + +// CHECK-LABEL: @cmp 2741 + +func.func @cmp() { 2742 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2743 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2744 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] 2745 + + // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 2746 + %1 = arith.constant dense<0> : tensor<128xi32> 2747 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] 2748 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none> 2749 + %2 = arith.cmpi eq, %0, %1 : tensor<128xi32> 2750 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] 2751 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none> 2752 + %3 = arith.cmpi slt, %0, %1 : tensor<128xi32> 2753 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2754 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 2755 + %4 = arith.cmpi sle, %0, %1 : tensor<128xi32> 2756 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] 2757 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none> 2758 + %5 = arith.cmpi sge, %0, %1 : tensor<128xi32> 2759 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8] 2760 + + // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8 2761 + %6 = arith.constant dense<8> : tensor<128xi32> 2762 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None] 2763 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = <none> 2764 + %7 = arith.cmpi sgt, %0, %6 : tensor<128xi32> 2765 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [0] 2766 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 0 2767 + %8 = arith.cmpi sgt, %1, %6 : tensor<128xi32> 2768 + return 2769 + } 2770 + 2771 + // ----- 2772 + 2773 + -// CHECK-LABEL: logic 2774 + -func @logic() { 2775 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2776 + +// CHECK-LABEL: @logic 2777 + +func.func @logic() { 2778 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2779 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2780 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] 2781 + + // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 2782 + %1 = arith.constant dense<64> : tensor<128xi32> 2783 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16777216] ; Constancy: [64] ; ConstantValue: [None] 2784 + + // CHECK-NEXT: contiguity = [1], divisibility = [16777216], constancy = [64], constant_value = <none> 2785 + %2 = arith.divsi %0, %1 : tensor<128xi32> 2786 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8] 2787 + + // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8 2788 + %3 = arith.constant dense<8> : tensor<128xi32> 2789 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [134217728] ; Constancy: [8] ; ConstantValue: [None] 2790 + + // CHECK-NEXT: contiguity = [1], divisibility = [134217728], constancy = [8], constant_value = <none> 2791 + %4 = arith.divsi %0, %3 : tensor<128xi32> 2792 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2793 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 2794 + %5 = arith.andi %0, %1 : tensor<128xi32> 2795 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2796 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 2797 + %6 = arith.ori %0, %1 : tensor<128xi32> 2798 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2799 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 2800 + %7 = arith.xori %0, %1 : tensor<128xi32> 2801 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None] 2802 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = <none> 2803 + %8 = arith.andi %2, %4 : tensor<128xi32> 2804 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None] 2805 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = <none> 2806 + %9 = arith.ori %2, %4 : tensor<128xi32> 2807 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None] 2808 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = <none> 2809 + %10 = arith.xori %2, %4 : tensor<128xi32> 2810 + return 2811 + } 2812 + 2813 + // ----- 2814 + 2815 + -// CHECK-LABEL: select 2816 + -func @select() { 2817 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2818 + +// CHECK-LABEL: @select 2819 + +func.func @select() { 2820 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2821 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2822 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] 2823 + + // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 2824 + %1 = arith.constant dense<0> : tensor<128xi32> 2825 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] 2826 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none> 2827 + %2 = arith.cmpi eq, %0, %1 : tensor<128xi32> 2828 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] 2829 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none> 2830 + %3 = arith.cmpi slt, %0, %1 : tensor<128xi32> 2831 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [1] ; ConstantValue: [0] 2832 + + // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [1], constant_value = 0 2833 + %4 = arith.constant 0 : i1 2834 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] 2835 + + // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 2836 + %7 = tt.splat %4 : (i1) -> tensor<128xi1> 2837 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] 2838 + - %5 = select %4, %3, %7 : tensor<128xi1> 2839 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] 2840 + + // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 2841 + + %5 = arith.select %4, %3, %7 : tensor<128xi1> 2842 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = <none> 2843 + %8 = "triton_gpu.select"(%7, %3, %2) : (tensor<128xi1>, tensor<128xi1>, tensor<128xi1>) -> tensor<128xi1> 2844 + return 2845 + } 2846 + 2847 + // ----- 2848 + 2849 + -func @shift() { 2850 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2851 + +func.func @shift() { 2852 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2853 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2854 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8] 2855 + + // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8 2856 + %1 = arith.constant dense<8> : tensor<128xi32> 2857 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4] ; Constancy: [128] ; ConstantValue: [4] 2858 + + // CHECK-NEXT: contiguity = [1], divisibility = [4], constancy = [128], constant_value = 4 2859 + %2 = arith.constant dense<4> : tensor<128xi32> 2860 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [274877906944] ; Constancy: [1] ; ConstantValue: [None] 2861 + + // CHECK-NEXT: contiguity = [1], divisibility = [274877906944], constancy = [1], constant_value = <none> 2862 + %3 = arith.shli %0, %1 : tensor<128xi32> 2863 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [67108864] ; Constancy: [1] ; ConstantValue: [None] 2864 + + // CHECK-NEXT: contiguity = [1], divisibility = [67108864], constancy = [1], constant_value = <none> 2865 + %4 = arith.shrsi %0, %2 : tensor<128xi32> 2866 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] 2867 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 2868 + %5 = arith.shli %1, %2 : tensor<128xi32> 2869 + return 2870 + } 2871 + 2872 + // ----- 2873 + 2874 + -func @max_min() { 2875 + - // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2876 + +func.func @max_min() { 2877 + + // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2878 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2879 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [64] ; Constancy: [1] ; ConstantValue: [None] 2880 + + // CHECK-NEXT: contiguity = [128], divisibility = [64], constancy = [1], constant_value = <none> 2881 + %1 = tt.make_range {end = 192 : i32, start = 64 : i32} : tensor<128xi32> 2882 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2883 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 2884 + %2 = arith.maxsi %0, %1 : tensor<128xi32> 2885 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 2886 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 2887 + %3 = arith.minsi %0, %1 : tensor<128xi32> 2888 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8] 2889 + + // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8 2890 + %4 = arith.constant dense<8> : tensor<128xi32> 2891 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4] ; Constancy: [128] ; ConstantValue: [4] 2892 + + // CHECK-NEXT: contiguity = [1], divisibility = [4], constancy = [128], constant_value = 4 2893 + %5 = arith.constant dense<4> : tensor<128xi32> 2894 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [8] 2895 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = 8 2896 + %6 = arith.maxsi %4, %5 : tensor<128xi32> 2897 + return 2898 + } 2899 + 2900 + // ----- 2901 + 2902 + -// CHECK-LABEL: for 2903 + -func @for() { 2904 + - // CHECK: Contiguity: [1, 1] ; Divisibility: [4611686018427387904, 4611686018427387904] ; Constancy: [128, 32] ; ConstantValue: [0] 2905 + +// CHECK-LABEL: @for 2906 + +func.func @for() { 2907 + + // CHECK: contiguity = [1, 1], divisibility = [4611686018427387904, 4611686018427387904], constancy = [128, 32], constant_value = 0 2908 + %a_init = arith.constant dense<0> : tensor<128x32xi32> 2909 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 32] ; ConstantValue: [1] 2910 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 32], constant_value = 1 2911 + %b_init = arith.constant dense<1> : tensor<128x32xi32> 2912 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [4, 4] ; Constancy: [128, 32] ; ConstantValue: [4] 2913 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [4, 4], constancy = [128, 32], constant_value = 4 2914 + %c_init = arith.constant dense<4> : tensor<128x32xi32> 2915 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [128] 2916 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [1], constant_value = 128 2917 + %ub = arith.constant 128 : index 2918 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [1] ; ConstantValue: [0] 2919 + + // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [1], constant_value = 0 2920 + %lb = arith.constant 0 : index 2921 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [16] 2922 + + // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [1], constant_value = 16 2923 + %step = arith.constant 16 : index 2924 + %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>) { 2925 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [None] 2926 + + // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [1], constant_value = <none> 2927 + %t = arith.index_cast %iv : index to i32 2928 + - // CHECK: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 32] ; ConstantValue: [None] 2929 + - // CHECK: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 32] ; ConstantValue: [None] 2930 + - // CHECK: Contiguity: [1, 1] ; Divisibility: [4, 4] ; Constancy: [128, 32] ; ConstantValue: [4] 2931 + + // CHECK: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 32], constant_value = <none> 2932 + + // CHECK: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 32], constant_value = <none> 2933 + + // CHECK: contiguity = [1, 1], divisibility = [4, 4], constancy = [128, 32], constant_value = 4 2934 + scf.yield %b, %a, %c : tensor<128x32xi32>, tensor<128x32xi32>, tensor<128x32xi32> 2935 + } 2936 + return 2937 + @@ -290,53 +290,53 @@ func @for() { 2938 + 2939 + // ----- 2940 + 2941 + -// CHECK-LABEL: permute_2d 2942 + -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}) { 2943 + - // CHECK: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 128] ; ConstantValue: [1] 2944 + +// CHECK-LABEL: @permute_2d 2945 + +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}) { 2946 + + // CHECK: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 128], constant_value = 1 2947 + %cst = arith.constant dense<true> : tensor<128x128xi1> 2948 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [1, 1] ; ConstantValue: [None] 2949 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [1, 1], constancy = [1, 1], constant_value = <none> 2950 + %cst_0 = arith.constant dense<0.000000e+00> : tensor<128x128xf32> 2951 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2952 + + // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2953 + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2954 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 2955 + + // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 2956 + %1 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 2957 + - // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [1073741824, 1] ; Constancy: [1, 1] ; ConstantValue: [None] 2958 + + // CHECK-NEXT: contiguity = [128, 1], divisibility = [1073741824, 1], constancy = [1, 1], constant_value = <none> 2959 + %2 = tt.expand_dims %0 {axis = 1 : i32} : (tensor<128xi32>) -> tensor<128x1xi32> 2960 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 1] ; ConstantValue: [None] 2961 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 1], constant_value = <none> 2962 + %3 = tt.splat %arg1 : (i32) -> tensor<128x1xi32> 2963 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [17179869184, 16] ; Constancy: [1, 1] ; ConstantValue: [None] 2964 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [17179869184, 16], constancy = [1, 1], constant_value = <none> 2965 + %4 = arith.muli %2, %3 : tensor<128x1xi32> 2966 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 1] ; ConstantValue: [None] 2967 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 1], constant_value = <none> 2968 + %5 = tt.splat %arg0 : (!tt.ptr<f32>) -> tensor<128x1x!tt.ptr<f32>> 2969 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [1, 1] ; ConstantValue: [None] 2970 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [1, 1], constant_value = <none> 2971 + %6 = tt.addptr %5, %4 : tensor<128x1x!tt.ptr<f32>>, tensor<128x1xi32> 2972 + - // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 1073741824] ; Constancy: [1, 1] ; ConstantValue: [None] 2973 + + // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 1073741824], constancy = [1, 1], constant_value = <none> 2974 + %7 = tt.expand_dims %1 {axis = 0 : i32}: (tensor<128xi32>) -> tensor<1x128xi32> 2975 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [1, 128] ; ConstantValue: [None] 2976 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [1, 128], constant_value = <none> 2977 + %8 = tt.broadcast %6 : (tensor<128x1x!tt.ptr<f32>>) -> tensor<128x128x!tt.ptr<f32>> 2978 + - // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 1073741824] ; Constancy: [128, 1] ; ConstantValue: [None] 2979 + + // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 1073741824], constancy = [128, 1], constant_value = <none> 2980 + %9 = tt.broadcast %7 : (tensor<1x128xi32>) -> tensor<128x128xi32> 2981 + - // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 16] ; Constancy: [1, 1] ; ConstantValue: [None] 2982 + + // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 16], constancy = [1, 1], constant_value = <none> 2983 + %10 = tt.addptr %8, %9 : tensor<128x128x!tt.ptr<f32>>, tensor<128x128xi32> 2984 + - // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [1073741824, 1] ; Constancy: [1, 1] ; ConstantValue: [None] 2985 + + // CHECK-NEXT: contiguity = [128, 1], divisibility = [1073741824, 1], constancy = [1, 1], constant_value = <none> 2986 + %11 = tt.expand_dims %0 {axis = 1 : i32}: (tensor<128xi32>) -> tensor<128x1xi32> 2987 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 1] ; ConstantValue: [None] 2988 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 1], constant_value = <none> 2989 + %12 = tt.splat %arg2 : (!tt.ptr<f32>) -> tensor<128x1x!tt.ptr<f32>> 2990 + - // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [16, 1] ; Constancy: [1, 1] ; ConstantValue: [None] 2991 + + // CHECK-NEXT: contiguity = [128, 1], divisibility = [16, 1], constancy = [1, 1], constant_value = <none> 2992 + %13 = tt.addptr %12, %11 : tensor<128x1x!tt.ptr<f32>>, tensor<128x1xi32> 2993 + - // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 1073741824] ; Constancy: [1, 1] ; ConstantValue: [None] 2994 + + // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 1073741824], constancy = [1, 1], constant_value = <none> 2995 + %14 = tt.expand_dims %1 {axis = 0 : i32} : (tensor<128xi32>) -> tensor<1x128xi32> 2996 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [1, 128] ; ConstantValue: [None] 2997 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [1, 128], constant_value = <none> 2998 + %15 = tt.splat %arg3 : (i32) -> tensor<1x128xi32> 2999 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 17179869184] ; Constancy: [1, 1] ; ConstantValue: [None] 3000 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 17179869184], constancy = [1, 1], constant_value = <none> 3001 + %16 = arith.muli %14, %15 : tensor<1x128xi32> 3002 + - // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [16, 1] ; Constancy: [1, 128] ; ConstantValue: [None] 3003 + + // CHECK-NEXT: contiguity = [128, 1], divisibility = [16, 1], constancy = [1, 128], constant_value = <none> 3004 + %17 = tt.broadcast %13 : (tensor<128x1x!tt.ptr<f32>>) -> tensor<128x128x!tt.ptr<f32>> 3005 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 17179869184] ; Constancy: [128, 1] ; ConstantValue: [None] 3006 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 17179869184], constancy = [128, 1], constant_value = <none> 3007 + %18 = tt.broadcast %16 : (tensor<1x128xi32>) -> tensor<128x128xi32> 3008 + - // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [16, 1] ; Constancy: [1, 1] ; ConstantValue: [None] 3009 + + // CHECK-NEXT: contiguity = [128, 1], divisibility = [16, 1], constancy = [1, 1], constant_value = <none> 3010 + %19 = tt.addptr %17, %18 : tensor<128x128x!tt.ptr<f32>>, tensor<128x128xi32> 3011 + - // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [1, 1] ; ConstantValue: [None] 3012 + + // CHECK-NEXT: contiguity = [1, 1], divisibility = [1, 1], constancy = [1, 1], constant_value = <none> 3013 + %20 = tt.load %10, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<128x128xf32> 3014 + tt.store %19, %20, %cst : tensor<128x128xf32> 3015 + return 3016 + @@ -347,29 +347,29 @@ func @permute_2d(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: i32 {t 3017 + module { 3018 + 3019 + // This is a tiny test for verifying StoreOp-related alignment, It simply store a constant to a buffer. 3020 + -// CHECK-LABEL: store_constant_align 3021 + -func @store_constant_align(%addr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n: i32 {tt.divisibility = 16 : i32}) { 3022 + - // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 3023 + +// CHECK-LABEL: @store_constant_align 3024 + +func.func @store_constant_align(%addr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n: i32 {tt.divisibility = 16 : i32}) { 3025 + + // CHECK: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 3026 + %pid = tt.get_program_id {axis = 0 : i32} : i32 3027 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [128] 3028 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [1], constant_value = 128 3029 + %c128_i32 = arith.constant 128 : i32 3030 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [None] 3031 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [1], constant_value = <none> 3032 + %1 = arith.muli %pid, %c128_i32 : i32 3033 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] 3034 + + // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = <none> 3035 + %2 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> 3036 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [None] 3037 + + // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = <none> 3038 + %3 = tt.splat %1 : (i32) -> tensor<128xi32> 3039 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [None] 3040 + + // CHECK-NEXT: contiguity = [128], divisibility = [128], constancy = [1], constant_value = <none> 3041 + %4 = arith.addi %3, %2 : tensor<128xi32> 3042 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [128] ; ConstantValue: [None] 3043 + + // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [128], constant_value = <none> 3044 + %5 = tt.splat %addr : (!tt.ptr<f32>) -> tensor<128x!tt.ptr<f32>> 3045 + - // CHECK-NEXT: Contiguity: [128] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [None] 3046 + + // CHECK-NEXT: contiguity = [128], divisibility = [16], constancy = [1], constant_value = <none> 3047 + %6 = tt.addptr %5, %4 : tensor<128x!tt.ptr<f32>>, tensor<128xi32> 3048 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [128] ; ConstantValue: [None] 3049 + + // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [128], constant_value = <none> 3050 + %9 = tt.splat %n : (i32) -> tensor<128xi32> 3051 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [16] ; ConstantValue: [None] 3052 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [16], constant_value = <none> 3053 + %mask = arith.cmpi slt, %4, %9 : tensor<128xi32> 3054 + - // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] 3055 + + // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 3056 + %cst = arith.constant dense<0.0> : tensor<128xf32> 3057 + tt.store %5, %cst, %mask : tensor<128xf32> 3058 + return 3059 + @@ -381,8 +381,8 @@ func @store_constant_align(%addr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %n: 3060 + 3061 + // This IR is dumped from vecadd test. 3062 + // Note, the hint {tt.divisibility = 16 : i32} for %n_elements affects the alignment of mask. 3063 + -// CHECK-LABEL: vecadd_mask_align_16 3064 + -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}) { 3065 + +// CHECK-LABEL: @vecadd_mask_align_16 3066 + +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}) { 3067 + %c64_i32 = arith.constant 64 : i32 3068 + %0 = tt.get_program_id {axis = 0 : i32} : i32 3069 + %1 = arith.muli %0, %c64_i32 : i32 3070 + @@ -394,13 +394,13 @@ func @vecadd_mask_align_16(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %ar 3071 + %7 = tt.splat %arg1 : (!tt.ptr<f32>) -> tensor<64x!tt.ptr<f32>> 3072 + %8 = tt.addptr %7, %4 : tensor<64x!tt.ptr<f32>>, tensor<64xi32> 3073 + %9 = tt.splat %n_elements : (i32) -> tensor<64xi32> 3074 + - // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [16] ; ConstantValue: [None] ( %{{.*}} = arith.cmpi slt, %{{.*}}, %{{.*}} : tensor<64xi32> ) 3075 + + // CHECK: arith.cmpi slt, %{{.*}} => contiguity = [1], divisibility = [1], constancy = [16], constant_value = <none> 3076 + %mask = arith.cmpi slt, %4, %9 : tensor<64xi32> 3077 + %11 = tt.load %6, %mask {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32> 3078 + %12 = tt.load %8, %mask {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32> 3079 + %13 = arith.addf %11, %12 : tensor<64xf32> 3080 + %14 = tt.splat %arg2 : (!tt.ptr<f32>) -> tensor<64x!tt.ptr<f32>> 3081 + - // CHECK: Contiguity: [64] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [None] ( %{{.*}} = tt.addptr %{{.*}}, %{{.*}} : tensor<64x!tt.ptr<f32>>, tensor<64xi32> ) 3082 + + // CHECK: tt.addptr %{{.*}} => contiguity = [64], divisibility = [16], constancy = [1], constant_value = <none> 3083 + %15 = tt.addptr %14, %4 : tensor<64x!tt.ptr<f32>>, tensor<64xi32> 3084 + tt.store %15, %13, %mask : tensor<64xf32> 3085 + return 3086 + @@ -410,8 +410,8 @@ func @vecadd_mask_align_16(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %ar 3087 + 3088 + // This IR is dumped from vecadd test. 3089 + // Note, there is no divisibility hint for %n_elements, Triton should assume its divisibility to be 1 by default. 3090 + -// CHECK-LABEL: vecadd_mask_align_1 3091 + -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) { 3092 + +// CHECK-LABEL: @vecadd_mask_align_1 3093 + +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) { 3094 + %c64_i32 = arith.constant 64 : i32 3095 + %0 = tt.get_program_id {axis = 0 : i32} : i32 3096 + %1 = arith.muli %0, %c64_i32 : i32 3097 + @@ -423,7 +423,7 @@ func @vecadd_mask_align_1(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg 3098 + %7 = tt.splat %arg1 : (!tt.ptr<f32>) -> tensor<64x!tt.ptr<f32>> 3099 + %8 = tt.addptr %7, %4 : tensor<64x!tt.ptr<f32>>, tensor<64xi32> 3100 + %9 = tt.splat %n_elements : (i32) -> tensor<64xi32> 3101 + - // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ( %{{.*}} = arith.cmpi slt, %{{.*}}, %{{.*}} : tensor<64xi32> ) 3102 + + // CHECK: arith.cmpi slt, %{{.*}} => contiguity = [1], divisibility = [1], constancy = [1], constant_value = <none> 3103 + %10 = arith.cmpi slt, %4, %9 : tensor<64xi32> 3104 + %11 = tt.load %6, %10 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32> 3105 + %12 = tt.load %8, %10 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32> 3106 + diff --git a/test/Analysis/test-allocation.mlir b/test/Analysis/test-allocation.mlir 3107 + index efb00c404d..f79222aa7b 100644 3108 + --- a/test/Analysis/test-allocation.mlir 3109 + +++ b/test/Analysis/test-allocation.mlir 3110 + @@ -13,7 +13,7 @@ 3111 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3112 + 3113 + // CHECK-LABEL: matmul_loop 3114 + -func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3115 + +func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3116 + %a_ptr_init = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL> 3117 + %b_ptr_init = tt.broadcast %B : (!tt.ptr<f16>) -> tensor<32x128x!tt.ptr<f16>, #BL> 3118 + 3119 + @@ -46,7 +46,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B 3120 + 3121 + // Shared memory is available after a tensor's liveness range ends 3122 + // CHECK-LABEL: reusable 3123 + -func @reusable(%A : !tt.ptr<f16>) { 3124 + +func.func @reusable(%A : !tt.ptr<f16>) { 3125 + %cst1 = arith.constant dense<true> : tensor<128x32xi1, #AL> 3126 + %cst2 = arith.constant dense<0.000000e+00> : tensor<128x32xf16, #AL> 3127 + %cst3 = arith.constant dense<true> : tensor<32x128xi1, #AL> 3128 + @@ -78,7 +78,7 @@ func @reusable(%A : !tt.ptr<f16>) { 3129 + // %cst1->%cst4 3130 + // %cst3->%g->%h->%i 3131 + // CHECK-LABEL: preallocate 3132 + -func @preallocate(%A : !tt.ptr<f16>) { 3133 + +func.func @preallocate(%A : !tt.ptr<f16>) { 3134 + // CHECK: offset = 0, size = 512 3135 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3136 + // CHECK-NEXT: offset = 1024, size = 512 3137 + @@ -113,7 +113,7 @@ func @preallocate(%A : !tt.ptr<f16>) { 3138 + 3139 + // Unused tensors are immediately released 3140 + // CHECK-LABEL: unused 3141 + -func @unused(%A : !tt.ptr<f16>) { 3142 + +func.func @unused(%A : !tt.ptr<f16>) { 3143 + // CHECK: offset = 0, size = 1024 3144 + %cst0 = arith.constant dense<0.000000e+00> : tensor<32x16xf16, #A_SHARED> 3145 + // CHECK-NEXT: offset = 0, size = 512 3146 + @@ -128,7 +128,7 @@ func @unused(%A : !tt.ptr<f16>) { 3147 + 3148 + // cst0 is alive through the entire function, it cannot be released before the end of the function 3149 + // CHECK-LABEL: longlive 3150 + -func @longlive(%A : !tt.ptr<f16>) { 3151 + +func.func @longlive(%A : !tt.ptr<f16>) { 3152 + // CHECK: offset = 0, size = 512 3153 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3154 + // CHECK-NEXT: offset = 512, size = 512 3155 + @@ -156,7 +156,7 @@ func @longlive(%A : !tt.ptr<f16>) { 3156 + } 3157 + 3158 + // CHECK-LABEL: alloc 3159 + -func @alloc(%A : !tt.ptr<f16>) { 3160 + +func.func @alloc(%A : !tt.ptr<f16>) { 3161 + // CHECK: offset = 0, size = 512 3162 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3163 + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #AL> 3164 + @@ -167,7 +167,7 @@ func @alloc(%A : !tt.ptr<f16>) { 3165 + } 3166 + 3167 + // CHECK-LABEL: scratch 3168 + -func @scratch() { 3169 + +func.func @scratch() { 3170 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> 3171 + // CHECK: scratch offset = 0, size = 512 3172 + %b = tt.reduce %cst0 {redOp = 1 : i32, axis = 0 : i32} : tensor<16x16xf16, #AL> -> tensor<16xf16, #sliceAd0> 3173 + @@ -176,7 +176,7 @@ func @scratch() { 3174 + } 3175 + 3176 + // CHECK-LABEL: trans 3177 + -func @trans(%A : !tt.ptr<f16>) { 3178 + +func.func @trans(%A : !tt.ptr<f16>) { 3179 + // CHECK: offset = 0, size = 1024 3180 + %tensor = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #A_SHARED> 3181 + %b = tt.trans %tensor : (tensor<16x32xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED_T> 3182 + @@ -184,7 +184,7 @@ func @trans(%A : !tt.ptr<f16>) { 3183 + } 3184 + 3185 + // CHECK-LABEL: insert_slice_async 3186 + -func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) { 3187 + +func.func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) { 3188 + %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL> 3189 + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> 3190 + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> 3191 + @@ -197,7 +197,7 @@ func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) { 3192 + } 3193 + 3194 + // CHECK-LABEL: extract_slice 3195 + -func @extract_slice(%A : !tt.ptr<f16>) { 3196 + +func.func @extract_slice(%A : !tt.ptr<f16>) { 3197 + // CHECK: offset = 0, size = 512 3198 + %cst0 = arith.constant dense<0.000000e+00> : tensor<1x16x16xf16, #A_SHARED> 3199 + %index = arith.constant 0 : index 3200 + @@ -209,7 +209,7 @@ func @extract_slice(%A : !tt.ptr<f16>) { 3201 + // B0 -> (B1) -> B0 3202 + // Memory used by B1 can be reused by B0. 3203 + // CHECK-LABEL: if 3204 + -func @if(%i1 : i1) { 3205 + +func.func @if(%i1 : i1) { 3206 + // CHECK: offset = 0, size = 512 3207 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3208 + // CHECK-NEXT: offset = 512, size = 512 3209 + @@ -233,7 +233,7 @@ func @if(%i1 : i1) { 3210 + // B0 -> (B1) -> (B2) -> B0 3211 + // Memory used by B0 cannot be reused by B1 or B2. 3212 + // CHECK-LABEL: if_else 3213 + -func @if_else(%i1 : i1) { 3214 + +func.func @if_else(%i1 : i1) { 3215 + // CHECK: offset = 0, size = 512 3216 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3217 + // CHECK-NEXT: offset = 512, size = 512 3218 + @@ -260,7 +260,7 @@ func @if_else(%i1 : i1) { 3219 + // Block arguments and yields are memory aliases that do not trigger a new 3220 + // allocation. 3221 + // CHECK-LABEL: for 3222 + -func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3223 + +func.func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3224 + // CHECK: offset = 0, size = 8192 3225 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3226 + // CHECK-NEXT: offset = 8192, size = 8192 3227 + @@ -275,7 +275,7 @@ func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.p 3228 + } 3229 + 3230 + // CHECK-LABEL: for_if_slice 3231 + -func @for_if_slice(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 3232 + +func.func @for_if_slice(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 3233 + // CHECK: offset = 0, size = 8192 3234 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3235 + // CHECK-NEXT: offset = 8192, size = 8192 3236 + @@ -296,7 +296,7 @@ func @for_if_slice(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, % 3237 + 3238 + // c0 cannot be released in the loop 3239 + // CHECK-LABEL: for_use_ancestor 3240 + -func @for_use_ancestor(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 3241 + +func.func @for_use_ancestor(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 3242 + // CHECK: offset = 0, size = 8192 3243 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3244 + // CHECK-NEXT: offset = 8192, size = 8192 3245 + @@ -316,7 +316,7 @@ func @for_use_ancestor(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16 3246 + // a_shared_init, b_shared_init, and c_shared_init's liveness ranges are span over the entire function before cst2. 3247 + // So they cannot be reused by cst0 and cst1, but can be reused by cst2. 3248 + // CHECK-LABEL: for_if_for 3249 + -func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 3250 + +func.func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>, %i1 : i1) { 3251 + // CHECK: offset = 0, size = 8192 3252 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3253 + // CHECK-NEXT: offset = 8192, size = 8192 3254 + diff --git a/test/Analysis/test-membar.mlir b/test/Analysis/test-membar.mlir 3255 + index 7199e5f53d..17880b2094 100644 3256 + --- a/test/Analysis/test-membar.mlir 3257 + +++ b/test/Analysis/test-membar.mlir 3258 + @@ -14,7 +14,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3259 + 3260 + // CHECK-LABEL: matmul_loop 3261 + // There shouldn't be any membar with the dot op encoding. 3262 + -func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3263 + +func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3264 + %a_ptr_init = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL> 3265 + %b_ptr_init = tt.broadcast %B : (!tt.ptr<f16>) -> tensor<32x128x!tt.ptr<f16>, #BL> 3266 + 3267 + @@ -42,7 +42,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B 3268 + } 3269 + 3270 + // CHECK-LABEL: raw_single_block 3271 + -func @raw_single_block(%A : !tt.ptr<f16>) { 3272 + +func.func @raw_single_block(%A : !tt.ptr<f16>) { 3273 + %cst1 = arith.constant dense<true> : tensor<128x32xi1, #AL> 3274 + %cst2 = arith.constant dense<0.000000e+00> : tensor<128x32xf16, #AL> 3275 + %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL> 3276 + @@ -54,7 +54,7 @@ func @raw_single_block(%A : !tt.ptr<f16>) { 3277 + } 3278 + 3279 + // CHECK-LABEL: war_single_block 3280 + -func @war_single_block(%A : !tt.ptr<f16>) { 3281 + +func.func @war_single_block(%A : !tt.ptr<f16>) { 3282 + %cst1 = arith.constant dense<true> : tensor<128x32xi1, #AL> 3283 + %cst2 = arith.constant dense<0.000000e+00> : tensor<128x32xf16, #AL> 3284 + %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL> 3285 + @@ -70,7 +70,7 @@ func @war_single_block(%A : !tt.ptr<f16>) { 3286 + } 3287 + 3288 + // CHECK-LABEL: scratch 3289 + -func @scratch() { 3290 + +func.func @scratch() { 3291 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3292 + // CHECK: Membar 1 3293 + %a = tt.cat %cst0, %cst0 {axis = 0} : (tensor<16x16xf16, #A_SHARED>, tensor<16x16xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED> 3294 + @@ -81,7 +81,7 @@ func @scratch() { 3295 + } 3296 + 3297 + // CHECK-LABEL: async_wait 3298 + -func @async_wait() { 3299 + +func.func @async_wait() { 3300 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3301 + // CHECK: Membar 1 3302 + %a = tt.cat %cst0, %cst0 {axis = 0} : (tensor<16x16xf16, #A_SHARED>, tensor<16x16xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED> 3303 + @@ -92,7 +92,7 @@ func @async_wait() { 3304 + } 3305 + 3306 + // CHECK-LABEL: alloc 3307 + -func @alloc() { 3308 + +func.func @alloc() { 3309 + %cst0 = triton_gpu.alloc_tensor : tensor<16x16xf16, #A_SHARED> 3310 + %a = tt.cat %cst0, %cst0 {axis = 0} : (tensor<16x16xf16, #A_SHARED>, tensor<16x16xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED> 3311 + // CHECK: Membar 2 3312 + @@ -101,7 +101,7 @@ func @alloc() { 3313 + } 3314 + 3315 + // CHECK-LABEL: extract_slice 3316 + -func @extract_slice() { 3317 + +func.func @extract_slice() { 3318 + %cst0 = arith.constant dense<0.000000e+00> : tensor<1x16x16xf16, #A_SHARED> 3319 + %index = arith.constant 0 : index 3320 + %cst1 = tensor.extract_slice %cst0[%index, 0, 0][1, 16, 16][1, 1, 1] : tensor<1x16x16xf16, #A_SHARED> to tensor<16x16xf16, #A_SHARED> 3321 + @@ -113,14 +113,14 @@ func @extract_slice() { 3322 + } 3323 + 3324 + // CHECK-LABEL: trans 3325 + -func @trans() { 3326 + +func.func @trans() { 3327 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #A_SHARED> 3328 + %b = tt.trans %cst0 : (tensor<16x32xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED_T> 3329 + return 3330 + } 3331 + 3332 + // CHECK-LABEL: insert_slice_async 3333 + -func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) { 3334 + +func.func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) { 3335 + %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL> 3336 + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> 3337 + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> 3338 + @@ -135,7 +135,7 @@ func @insert_slice_async(%A : !tt.ptr<f16>, %i1 : i1) { 3339 + } 3340 + 3341 + // CHECK-LABEL: insert_slice 3342 + -func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) { 3343 + +func.func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) { 3344 + %a_ptr = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<16x16x!tt.ptr<f16>, #AL> 3345 + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> 3346 + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> 3347 + @@ -153,7 +153,7 @@ func @insert_slice(%A : !tt.ptr<f16>, %i1 : i1) { 3348 + 3349 + // If branch inserted a barrier for %cst0 and %cst1, but else didn't, then the barrier should be inserted in the parent region 3350 + // CHECK-LABEL: multi_blocks 3351 + -func @multi_blocks(%i1 : i1) { 3352 + +func.func @multi_blocks(%i1 : i1) { 3353 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3354 + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3355 + scf.if %i1 { 3356 + @@ -174,7 +174,7 @@ func @multi_blocks(%i1 : i1) { 3357 + 3358 + // Both branches inserted a barrier for %cst0 and %cst1, then the barrier doesn't need to be inserted in the parent region 3359 + // CHECK-LABEL: multi_blocks_join_barrier 3360 + -func @multi_blocks_join_barrier(%i1 : i1) { 3361 + +func.func @multi_blocks_join_barrier(%i1 : i1) { 3362 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3363 + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3364 + scf.if %i1 { 3365 + @@ -192,7 +192,7 @@ func @multi_blocks_join_barrier(%i1 : i1) { 3366 + 3367 + // Read yielded tensor requires a barrier 3368 + // CHECK-LABEL: multi_blocks_yield 3369 + -func @multi_blocks_yield(%i1 : i1) { 3370 + +func.func @multi_blocks_yield(%i1 : i1) { 3371 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3372 + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3373 + %a = scf.if %i1 -> (tensor<32x16xf16, #A_SHARED>) { 3374 + @@ -212,7 +212,7 @@ func @multi_blocks_yield(%i1 : i1) { 3375 + 3376 + // Conservatively add a barrier as if the branch (%i1) is never taken 3377 + // CHECK-LABEL: multi_blocks_noelse 3378 + -func @multi_blocks_noelse(%i1 : i1) { 3379 + +func.func @multi_blocks_noelse(%i1 : i1) { 3380 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3381 + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3382 + scf.if %i1 { 3383 + @@ -226,7 +226,7 @@ func @multi_blocks_noelse(%i1 : i1) { 3384 + 3385 + // Conservatively add a barrier as if the branch (%i2) is never taken 3386 + // CHECK-LABEL: multi_blocks_nested_scf 3387 + -func @multi_blocks_nested_scf(%i1 : i1, %i2 : i1) { 3388 + +func.func @multi_blocks_nested_scf(%i1 : i1, %i2 : i1) { 3389 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3390 + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> 3391 + scf.if %i1 { 3392 + @@ -247,7 +247,7 @@ func @multi_blocks_nested_scf(%i1 : i1, %i2 : i1) { 3393 + } 3394 + 3395 + // CHECK-LABEL: for 3396 + -func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3397 + +func.func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3398 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3399 + %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3400 + %c_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3401 + @@ -262,7 +262,7 @@ func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.p 3402 + // Although a_shared and b_shared are synced before entering the loop, 3403 + // they are reassociated with aliases (c_shared) and thus require a barrier. 3404 + // CHECK-LABEL: for_alias 3405 + -func @for_alias(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3406 + +func.func @for_alias(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3407 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3408 + %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3409 + // CHECK-NEXT: Membar 2 3410 + @@ -282,7 +282,7 @@ func @for_alias(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : 3411 + // Although cst2 is not an argument of scf.yield, its memory is reused by cst1. 3412 + // So we need a barrier both before and after cst1 3413 + // CHECK-LABEL: for_reuse 3414 + -func @for_reuse(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3415 + +func.func @for_reuse(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3416 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3417 + %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3418 + // CHECK-NEXT: Membar 2 3419 + @@ -302,7 +302,7 @@ func @for_reuse(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : 3420 + 3421 + 3422 + // CHECK-LABEL: for_reuse_nested 3423 + -func @for_reuse_nested(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3424 + +func.func @for_reuse_nested(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 3425 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3426 + %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> 3427 + // CHECK-NEXT: Membar 2 3428 + diff --git a/test/Conversion/triton_ops.mlir b/test/Conversion/triton_ops.mlir 3429 + index e9ee502435..0e979b148d 100644 3430 + --- a/test/Conversion/triton_ops.mlir 3431 + +++ b/test/Conversion/triton_ops.mlir 3432 + @@ -1,6 +1,6 @@ 3433 + // RUN: triton-opt %s | FileCheck %s 3434 + 3435 + -func @cast_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_f32: f32, %scalar_i64: i64) { 3436 + +func.func @cast_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_f32: f32, %scalar_i64: i64) { 3437 + // scalar -> scalar 3438 + // CHECK: i64 -> !tt.ptr<f32> 3439 + %0 = tt.int_to_ptr %scalar_i64 : i64 -> !tt.ptr<f32> 3440 + @@ -35,7 +35,7 @@ func @cast_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_f32: f32, %scalar_i64: i64) { 3441 + return 3442 + } 3443 + 3444 + -func @addptr_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_i32: i32) { 3445 + +func.func @addptr_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_i32: i32) { 3446 + // scalar -> scalar 3447 + // CHECK: !tt.ptr<f32> 3448 + %0 = tt.addptr %scalar_ptr, %scalar_i32 : !tt.ptr<f32>, i32 3449 + @@ -54,7 +54,7 @@ func @addptr_ops(%scalar_ptr: !tt.ptr<f32>, %scalar_i32: i32) { 3450 + return 3451 + } 3452 + 3453 + -func @load_store_ops_scalar(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %mask : i1) { 3454 + +func.func @load_store_ops_scalar(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %mask : i1) { 3455 + // Test if Load/Store ops can handle scalar values 3456 + %other = arith.constant 0.0e+0 : f32 3457 + 3458 + @@ -76,7 +76,7 @@ func @load_store_ops_scalar(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %ma 3459 + return 3460 + } 3461 + 3462 + -func @reduce_ops_infer(%ptr: !tt.ptr<f32>, %v : tensor<1x2x4xf32>) { 3463 + +func.func @reduce_ops_infer(%ptr: !tt.ptr<f32>, %v : tensor<1x2x4xf32>) { 3464 + // Test if reduce ops infer types correctly 3465 + 3466 + // CHECK: %{{.*}} = tt.reduce %{{.*}} -> tensor<2x4xf32> 3467 + @@ -101,7 +101,7 @@ func @reduce_ops_infer(%ptr: !tt.ptr<f32>, %v : tensor<1x2x4xf32>) { 3468 + return 3469 + } 3470 + 3471 + -func @dot_ops_infer(%ptr: !tt.ptr<f32>, %v : f32) { 3472 + +func.func @dot_ops_infer(%ptr: !tt.ptr<f32>, %v : f32) { 3473 + // Test if reduce ops infer types correctly 3474 + %v128x32 = tt.splat %v : (f32) -> tensor<128x32xf32> 3475 + %v32x128 = tt.splat %v : (f32) -> tensor<32x128xf32> 3476 + diff --git a/test/Conversion/triton_to_tritongpu.mlir b/test/Conversion/triton_to_tritongpu.mlir 3477 + index a160bc8815..b461ca542f 100644 3478 + --- a/test/Conversion/triton_to_tritongpu.mlir 3479 + +++ b/test/Conversion/triton_to_tritongpu.mlir 3480 + @@ -1,6 +1,6 @@ 3481 + // RUN: triton-opt %s -split-input-file -convert-triton-to-tritongpu=num-warps=2 | FileCheck %s 3482 + 3483 + -func @ops() { 3484 + +func.func @ops() { 3485 + // CHECK: module attributes {"triton_gpu.num-warps" = 2 : i32} {{.*}} 3486 + %a = arith.constant dense<1.00e+00> : tensor<128x32xf16> 3487 + %b = arith.constant dense<2.00e+00> : tensor<32x128xf16> 3488 + @@ -11,7 +11,7 @@ func @ops() { 3489 + 3490 + // ----- 3491 + 3492 + -func @load_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) { 3493 + +func.func @load_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) { 3494 + // Test if LoadOp is lowered properly (see #771) 3495 + %ptrs = tt.splat %ptr : (!tt.ptr<f32>) -> tensor<128x!tt.ptr<f32>> 3496 + %mask = arith.constant dense<true> : tensor<128xi1> 3497 + @@ -30,7 +30,7 @@ func @load_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) { 3498 + 3499 + // ----- 3500 + 3501 + -func @reduce_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) { 3502 + +func.func @reduce_ops(%ptr: !tt.ptr<f32> {tt.divisibility = 16 : i32}) { 3503 + // Test if the total number of threadsPerWarp is 32 3504 + // Test if the total number of warps is 2 3505 + // CHECK: #blocked0 = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [4, 8], warpsPerCTA = [1, 2], order = [0, 1]}> 3506 + diff --git a/test/Conversion/tritongpu_to_llvm.mlir b/test/Conversion/tritongpu_to_llvm.mlir 3507 + index e9e7d5a340..507b362c99 100644 3508 + --- a/test/Conversion/tritongpu_to_llvm.mlir 3509 + +++ b/test/Conversion/tritongpu_to_llvm.mlir 3510 + @@ -4,7 +4,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3511 + // CHECK: llvm.func @test_empty_kernel(%arg0: i32, %arg1: !llvm.ptr<f16, 1>) 3512 + // Here the 128 comes from the 4 in module attribute multiples 32 3513 + // CHECK: attributes {nvvm.kernel = 1 : ui1, nvvm.maxntid = 128 : i32} {{.*}} 3514 + - func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) { 3515 + + func.func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) { 3516 + // CHECK: llvm.return 3517 + return 3518 + } 3519 + @@ -15,7 +15,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3520 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3521 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3522 + // CHECK-LABEL: basic_load 3523 + - func @basic_load(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) { 3524 + + func.func @basic_load(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) { 3525 + // CHECK: llvm.inline_asm 3526 + // CHECK: llvm.inline_asm 3527 + %1 = tt.load %a_ptr_init, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked0> 3528 + @@ -28,7 +28,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3529 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3530 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3531 + // CHECK-LABEL: vectorized_load 3532 + - func @vectorized_load(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) { 3533 + + func.func @vectorized_load(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) { 3534 + // CHECK: llvm.inline_asm 3535 + // CHECK-SAME: ld.global.b32 3536 + // CHECK: llvm.inline_asm 3537 + @@ -43,7 +43,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3538 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> 3539 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3540 + // CHECK-LABEL: vectorized_load_f16 3541 + - func @vectorized_load_f16(%a_ptr_init: tensor<256x!tt.ptr<f16>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf16, #blocked0>) { 3542 + + func.func @vectorized_load_f16(%a_ptr_init: tensor<256x!tt.ptr<f16>, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf16, #blocked0>) { 3543 + // CHECK: llvm.inline_asm 3544 + // CHECK-SAME: ld.global.b16 3545 + // CHECK: llvm.inline_asm 3546 + @@ -59,7 +59,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3547 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [8], order = [0]}> 3548 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3549 + // CHECK-LABEL: masked_load_const_other 3550 + - func @masked_load_const_other(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>) { 3551 + + func.func @masked_load_const_other(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>) { 3552 + %cst_0 = arith.constant dense<0.000000e+00> : tensor<256xf32, #blocked0> 3553 + %1 = tt.load %a_ptr_init, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked0> 3554 + return 3555 + @@ -72,7 +72,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3556 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [8], order = [0]}> 3557 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3558 + // CHECK-LABEL: masked_load_const_other_vec 3559 + - func @masked_load_const_other_vec(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>) { 3560 + + func.func @masked_load_const_other_vec(%a_ptr_init : tensor<256x!tt.ptr<f32>, #blocked0>, %cst : tensor<256xi1, #blocked0>) { 3561 + %cst_0 = arith.constant dense<0.000000e+00> : tensor<256xf32, #blocked0> 3562 + %1 = tt.load %a_ptr_init, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked0> 3563 + return 3564 + @@ -84,7 +84,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3565 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [2], order = [0]}> 3566 + module attributes {"triton_gpu.num-warps" = 2 : i32} { 3567 + // CHECK-LABEL: global_load_store_no_vec 3568 + - 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) { 3569 + + 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) { 3570 + %c256_i32 = arith.constant 256 : i32 3571 + %0 = tt.get_program_id {axis = 0 : i32} : i32 3572 + %1 = arith.muli %0, %c256_i32 : i32 3573 + @@ -128,7 +128,7 @@ module attributes {"triton_gpu.num-warps" = 2 : i32} { 3574 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [2], order = [0]}> 3575 + module attributes {"triton_gpu.num-warps" = 2 : i32} { 3576 + // CHECK-LABEL: global_load_store_vec4 3577 + - 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) { 3578 + + 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) { 3579 + %c256_i32 = arith.constant 256 : i32 3580 + %0 = tt.get_program_id {axis = 0 : i32} : i32 3581 + %1 = arith.muli %0, %c256_i32 : i32 3582 + @@ -165,7 +165,7 @@ module attributes {"triton_gpu.num-warps" = 2 : i32} { 3583 + #blocked = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [2], order = [0]}> 3584 + // 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. 3585 + module attributes {"triton_gpu.num-warps" = 2 : i32} { 3586 + - 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) { 3587 + + 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) { 3588 + %c64_i32 = arith.constant 64 : i32 3589 + %0 = tt.get_program_id {axis = 0 : i32} : i32 3590 + %1 = arith.muli %0, %c64_i32 : i32 3591 + @@ -195,7 +195,7 @@ module attributes {"triton_gpu.num-warps" = 2 : i32} { 3592 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> 3593 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3594 + // CHECK-LABEL: global_load_store_vec2 3595 + - 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) { 3596 + + 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) { 3597 + %c256_i32 = arith.constant 256 : i32 3598 + %0 = tt.get_program_id {axis = 0 : i32} : i32 3599 + %1 = arith.muli %0, %c256_i32 : i32 3600 + @@ -240,7 +240,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3601 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> 3602 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3603 + // CHECK-LABEL: global_load_store_vec8 3604 + - 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) { 3605 + + 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) { 3606 + %c256_i32 = arith.constant 256 : i32 3607 + %0 = tt.get_program_id {axis = 0 : i32} : i32 3608 + %1 = arith.muli %0, %c256_i32 : i32 3609 + @@ -283,7 +283,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3610 + #blocked2 = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [4, 1], order = [0, 1]}> 3611 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3612 + // CHECK-LABEL: basic_view_broadcast 3613 + - func @basic_view_broadcast(%arg : tensor<256xf32,#blocked0>) { 3614 + + func.func @basic_view_broadcast(%arg : tensor<256xf32,#blocked0>) { 3615 + // CHECK: llvm.mlir.undef 3616 + // CHECK: %[[T0:.*]] = llvm.extractvalue 3617 + // CHECK: %[[T1:.*]] = llvm.extractvalue 3618 + @@ -307,7 +307,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3619 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> 3620 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3621 + // CHECK-LABEL: basic_make_range 3622 + - func @basic_make_range() { 3623 + + func.func @basic_make_range() { 3624 + // CHECK: nvvm.read.ptx.sreg.tid.x 3625 + // CHECK: llvm.mlir.undef 3626 + // CHECK: llvm.insertvalue 3627 + @@ -322,7 +322,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3628 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3629 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3630 + // CHECK-LABEL: basic_addf 3631 + - func @basic_addf(%arg0 : tensor<256xf32,#blocked0>, %arg1 : tensor<256xf32,#blocked0>) { 3632 + + func.func @basic_addf(%arg0 : tensor<256xf32,#blocked0>, %arg1 : tensor<256xf32,#blocked0>) { 3633 + // CHECK: llvm.fadd 3634 + // CHECK: llvm.fadd 3635 + %1 = arith.addf %arg0, %arg1 : tensor<256xf32,#blocked0> 3636 + @@ -335,7 +335,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3637 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3638 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3639 + // CHECK-LABEL: basic_addi 3640 + - func @basic_addi(%arg0 : tensor<256xi32,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) { 3641 + + func.func @basic_addi(%arg0 : tensor<256xi32,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) { 3642 + // CHECK: llvm.add 3643 + // CHECK: llvm.add 3644 + %1 = arith.addi %arg0, %arg1 : tensor<256xi32,#blocked0> 3645 + @@ -347,7 +347,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3646 + 3647 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3648 + // CHECK-LABEL: basic_program_id 3649 + - func @basic_program_id() { 3650 + + func.func @basic_program_id() { 3651 + // CHECK: nvvm.read.ptx.sreg.ctaid.x : i32 3652 + %0 = tt.get_program_id {axis = 0 : i32} : i32 3653 + return 3654 + @@ -359,7 +359,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3655 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3656 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3657 + // CHECK-LABEL: basic_addptr 3658 + - func @basic_addptr(%arg0 : tensor<256x!tt.ptr<f32>,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) { 3659 + + func.func @basic_addptr(%arg0 : tensor<256x!tt.ptr<f32>,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) { 3660 + // CHECK: llvm.getelementptr 3661 + // CHECK: llvm.getelementptr 3662 + %0 = tt.addptr %arg0, %arg1 : tensor<256x!tt.ptr<f32>, #blocked0>, tensor<256xi32, #blocked0> 3663 + @@ -373,7 +373,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3664 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3665 + // CHECK: llvm.mlir.global external @global_smem 3666 + // CHECK-LABEL: basic_alloc_tensor 3667 + - func @basic_alloc_tensor() { 3668 + + func.func @basic_alloc_tensor() { 3669 + // CHECK: llvm.mlir.addressof @global_smem 3670 + // CHECK-NEXT: llvm.bitcast 3671 + // CHECK-NEXT: llvm.mlir.constant 3672 + @@ -390,7 +390,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3673 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3674 + // CHECK: llvm.mlir.global external @global_smem 3675 + // CHECK-LABEL: basic_extract_slice 3676 + - func @basic_extract_slice() { 3677 + + func.func @basic_extract_slice() { 3678 + // CHECK: llvm.mlir.addressof @global_smem 3679 + // CHECK: llvm.extractvalue 3680 + // CHECK-NEXT: llvm.extractvalue 3681 + @@ -423,7 +423,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3682 + 3683 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3684 + // CHECK-LABEL: basic_async_wait 3685 + - func @basic_async_wait() { 3686 + + func.func @basic_async_wait() { 3687 + // CHECK: cp.async.wait_group 0x4 3688 + triton_gpu.async_wait {num = 4: i32} 3689 + return 3690 + @@ -442,7 +442,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3691 + #A = #triton_gpu.shared<{vec = 8, perPhase = 1, maxPhase = 4, order = [1, 0]}> 3692 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3693 + // CHECK-LABEL: basic_insert_slice_async_fallback 3694 + - func @basic_insert_slice_async_fallback(%arg0: !tt.ptr<f16> {tt.divisibility = 1 : i32}) { 3695 + + func.func @basic_insert_slice_async_fallback(%arg0: !tt.ptr<f16> {tt.divisibility = 1 : i32}) { 3696 + %off0_ = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #slice2d1> 3697 + %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<64xi32, #slice3d0> 3698 + %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<16xi32, #slice2d1>) -> tensor<16x1xi32, #block2> 3699 + @@ -481,7 +481,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3700 + #A = #triton_gpu.shared<{vec = 8, perPhase = 1, maxPhase = 4, order = [1, 0]}> 3701 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3702 + // CHECK-LABEL: basic_insert_slice_async_v4 3703 + - func @basic_insert_slice_async_v4(%arg0: !tt.ptr<f32> {tt.divisibility = 32 : i32}) { 3704 + + func.func @basic_insert_slice_async_v4(%arg0: !tt.ptr<f32> {tt.divisibility = 32 : i32}) { 3705 + %off0_ = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #slice2d1> 3706 + %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<64xi32, #slice3d0> 3707 + %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<16xi32, #slice2d1>) -> tensor<16x1xi32, #block2> 3708 + @@ -523,7 +523,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3709 + #A = #triton_gpu.shared<{vec = 1, perPhase = 1, maxPhase = 4, order = [1, 0]}> 3710 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3711 + // CHECK-LABEL: basic_insert_slice_async_v1 3712 + - func @basic_insert_slice_async_v1(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}) { 3713 + + func.func @basic_insert_slice_async_v1(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}) { 3714 + %off0_ = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #slice2d1> 3715 + %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #slice3d0> 3716 + %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<16xi32, #slice2d1>) -> tensor<16x1xi32, #block2> 3717 + @@ -568,7 +568,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3718 + #A = #triton_gpu.shared<{vec = 1, perPhase = 1, maxPhase = 4, order = [1, 0]}> 3719 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3720 + // CHECK-LABEL: basic_insert_slice_async_v1_multictas 3721 + - func @basic_insert_slice_async_v1_multictas(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}) { 3722 + + func.func @basic_insert_slice_async_v1_multictas(%arg0: !tt.ptr<f32> {tt.divisibility = 4 : i32}) { 3723 + %off0_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #slice2d1> 3724 + %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #slice3d0> 3725 + %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<32xi32, #slice2d1>) -> tensor<32x1xi32, #block2> 3726 + @@ -619,7 +619,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3727 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3728 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3729 + // CHECK: basic_splat 3730 + - func @basic_splat(%ptr: !tt.ptr<f32>) { 3731 + + func.func @basic_splat(%ptr: !tt.ptr<f32>) { 3732 + // CHECK: llvm.mlir.undef 3733 + // CHECK: llvm.insertvalue 3734 + // CHECK: llvm.insertvalue 3735 + @@ -633,7 +633,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3736 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3737 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3738 + // CHECK-LABEL: basic_store 3739 + - func @basic_store(%ptrs: tensor<256x!tt.ptr<f32>, #blocked0>, %vals: tensor<256xf32, #blocked0>, %mask: tensor<256xi1, #blocked0>) { 3740 + + func.func @basic_store(%ptrs: tensor<256x!tt.ptr<f32>, #blocked0>, %vals: tensor<256xf32, #blocked0>, %mask: tensor<256xi1, #blocked0>) { 3741 + // CHECK: llvm.inline_asm 3742 + // CHECK-SAME: st.global.b32 [ ${{.*}} + 0 ], { ${{.*}} }; 3743 + // CHECK: llvm.inline_asm 3744 + @@ -650,7 +650,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3745 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3746 + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> 3747 + // CHECK-LABEL: convert_layout_blocked_blocked 3748 + - func @convert_layout_blocked_blocked(%arg0: tensor<16x16xf32, #blocked0>) { 3749 + + func.func @convert_layout_blocked_blocked(%arg0: tensor<16x16xf32, #blocked0>) { 3750 + // CHECK: llvm.mlir.addressof @global_smem 3751 + // CHECK: llvm.store 3752 + // CHECK-SAME: !llvm.ptr<vector<1xf32>, 3> 3753 + @@ -697,7 +697,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3754 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3755 + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> 3756 + // CHECK-LABEL: convert_layout_blocked_blocked_vec 3757 + - func @convert_layout_blocked_blocked_vec(%arg0: tensor<16x16xf32, #blocked0>) { 3758 + + func.func @convert_layout_blocked_blocked_vec(%arg0: tensor<16x16xf32, #blocked0>) { 3759 + // CHECK: llvm.mlir.addressof @global_smem 3760 + // CHECK: llvm.store 3761 + // CHECK-SAME: !llvm.ptr<vector<4xf32>, 3> 3762 + @@ -720,7 +720,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3763 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3764 + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> 3765 + // CHECK-LABEL: convert_layout_blocked_blocked_multi_rep 3766 + - func @convert_layout_blocked_blocked_multi_rep(%arg0: tensor<16x16xf32, #blocked0>) { 3767 + + func.func @convert_layout_blocked_blocked_multi_rep(%arg0: tensor<16x16xf32, #blocked0>) { 3768 + // CHECK: llvm.mlir.addressof @global_smem 3769 + // CHECK: llvm.store 3770 + // CHECK-SAME: !llvm.ptr<vector<4xf32>, 3> 3771 + @@ -751,7 +751,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3772 + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma0}> 3773 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3774 + // CHECK-LABEL: convert_dot 3775 + - func @convert_dot(%A: tensor<16x16xf16, #blocked0>, %B: tensor<16x16xf16, #blocked0>) { 3776 + + func.func @convert_dot(%A: tensor<16x16xf16, #blocked0>, %B: tensor<16x16xf16, #blocked0>) { 3777 + %AA = triton_gpu.convert_layout %A : (tensor<16x16xf16, #blocked0>) -> tensor<16x16xf16, #shared0> 3778 + %BB = triton_gpu.convert_layout %B : (tensor<16x16xf16, #blocked0>) -> tensor<16x16xf16, #shared0> 3779 + // CHECK: llvm.inline_asm 3780 + @@ -775,7 +775,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3781 + // TODO: problems in MLIR's parser on slice layout 3782 + // #blocked0 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [8, 4], warpsPerCTA = [1, 1], order = [1, 0]}> 3783 + // module attributes {"triton_gpu.num-warps" = 1 : i32} { 3784 + -// func @make_range_sliced_layout() { 3785 + +// func.func @make_range_sliced_layout() { 3786 + // %0 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #triton_gpu.slice<{dim = 0, parent = #blocked0}>> 3787 + // return 3788 + // } 3789 + @@ -788,7 +788,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3790 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3791 + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> 3792 + // CHECK-LABEL: convert_layout_mmav2_block 3793 + - func @convert_layout_mmav2_blocked(%arg0: tensor<32x16xf32, #mma>) { 3794 + + func.func @convert_layout_mmav2_blocked(%arg0: tensor<32x16xf32, #mma>) { 3795 + // CHECK: llvm.store 3796 + // CHECK-SAME: !llvm.ptr<vector<2xf32>, 3> 3797 + // CHECK: llvm.store 3798 + @@ -808,7 +808,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3799 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3800 + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> 3801 + // CHECK-LABEL: convert_layout_mmav1_block 3802 + - func @convert_layout_mmav1_blocked(%arg0: tensor<32x64xf32, #mma>) { 3803 + + func.func @convert_layout_mmav1_blocked(%arg0: tensor<32x64xf32, #mma>) { 3804 + // CHECK: llvm.store 3805 + // CHECK-SAME: !llvm.ptr<vector<2xf32>, 3> 3806 + // CHECK: llvm.store 3807 + @@ -831,7 +831,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3808 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3809 + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> 3810 + // CHECK-LABEL: convert_layout_blocked_shared 3811 + - func @convert_layout_blocked_shared(%arg0: tensor<128x32xf32, #blocked0>) { 3812 + + func.func @convert_layout_blocked_shared(%arg0: tensor<128x32xf32, #blocked0>) { 3813 + // CHECK: llvm.store 3814 + // CHECK-SAME: !llvm.ptr<vector<8xf32>, 3> 3815 + // CHECK: llvm.store 3816 + @@ -847,7 +847,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3817 + #blocked1 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [4, 8], warpsPerCTA = [1, 1], order = [1, 0]}> 3818 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3819 + // CHECK-LABEL: convert_blocked1d_to_slice0 3820 + - func @convert_blocked1d_to_slice0(%src:tensor<32xi32, #blocked0>) { 3821 + + func.func @convert_blocked1d_to_slice0(%src:tensor<32xi32, #blocked0>) { 3822 + // CHECK-COUNT-4: llvm.load {{.*}} : !llvm.ptr<vector<1xi32>, 3> 3823 + %cvt = triton_gpu.convert_layout %src : (tensor<32xi32, #blocked0>) -> tensor<32xi32, #triton_gpu.slice<{dim = 0, parent = #blocked1}>> 3824 + return 3825 + @@ -860,7 +860,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3826 + #blocked1 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [4, 8], warpsPerCTA = [1, 1], order = [1, 0]}> 3827 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3828 + // CHECK-LABEL: convert_blocked1d_to_slice1 3829 + - func @convert_blocked1d_to_slice1(%src:tensor<32xi32, #blocked0>) { 3830 + + func.func @convert_blocked1d_to_slice1(%src:tensor<32xi32, #blocked0>) { 3831 + // CHECK-COUNT-32: llvm.load {{.*}} : !llvm.ptr<vector<1xi32>, 3> 3832 + %cvt = triton_gpu.convert_layout %src : (tensor<32xi32, #blocked0>) -> tensor<32xi32, #triton_gpu.slice<{dim = 1, parent = #blocked1}>> 3833 + return 3834 + @@ -873,7 +873,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3835 + #blocked1 = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> 3836 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3837 + // CHECK-LABEL: convert_blocked_to_blocked_ptr 3838 + - func @convert_blocked_to_blocked_ptr(%src:tensor<32x!tt.ptr<f32>, #blocked0>) { 3839 + + func.func @convert_blocked_to_blocked_ptr(%src:tensor<32x!tt.ptr<f32>, #blocked0>) { 3840 + // CHECK: llvm.ptrtoint 3841 + // CHECK: llvm.store 3842 + // CHECK: nvvm.barrier0 3843 + @@ -892,7 +892,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3844 + #dot_operand_a = #triton_gpu.dot_op<{opIdx=0, parent=#mma}> 3845 + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma}> 3846 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3847 + - func @matmul_kernel_dot_operand_layout(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32}, 3848 + + func.func @matmul_kernel_dot_operand_layout(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32}, 3849 + %a:tensor<128x32xf16, #shared>, %b:tensor<32x256xf16, #shared>) { 3850 + %cst = arith.constant dense<0.000000e+00> : tensor<128x256xf32, #mma> 3851 + // CHECK: ldmatrix.sync.aligned.m8n8.x4.shared.b16 3852 + @@ -918,7 +918,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3853 + #dot_operand_a = #triton_gpu.dot_op<{opIdx=0, parent=#mma, isMMAv1Row=true}> 3854 + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma, isMMAv1Row=true}> 3855 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3856 + - func @matmul884_kernel_dot_operand_layout(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32}, 3857 + + func.func @matmul884_kernel_dot_operand_layout(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32}, 3858 + %a:tensor<32x64xf16, #shared0>, %b:tensor<64x64xf16, #shared1>) { 3859 + %cst = arith.constant dense<0.000000e+00> : tensor<32x64xf32, #mma> 3860 + // CHECK: ldmatrix.sync.aligned.m8n8.x4.shared.b16 3861 + @@ -941,7 +941,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3862 + #dot_operand_a = #triton_gpu.dot_op<{opIdx=0, parent=#blocked}> 3863 + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#blocked}> 3864 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3865 + - func @matmul_fmadot(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32}, 3866 + + func.func @matmul_fmadot(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32}, 3867 + %a:tensor<32x16xf32, #shared>, %b:tensor<16x32xf32, #shared>) { 3868 + %cst = arith.constant dense<0.000000e+00> : tensor<32x32xf32, #blocked> 3869 + // CHECK: llvm.intr.fmuladd 3870 + @@ -965,7 +965,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3871 + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma}> 3872 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3873 + // CHECK-LABEL: matmul_tf32dot 3874 + - func @matmul_tf32dot(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32}, 3875 + + func.func @matmul_tf32dot(%ptr:!tt.ptr<f32> {tt.divisibility = 16 : i32}, 3876 + %a:tensor<32x16xf32, #shared>, %b:tensor<16x32xf32, #shared>) { 3877 + %cst = arith.constant dense<0.000000e+00> : tensor<32x32xf32, #mma> 3878 + // CHECK: llvm.inline_asm 3879 + @@ -1000,7 +1000,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3880 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3881 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3882 + // CHECK-LABEL: atomic_add_f32 3883 + - func @atomic_add_f32(%arg0 : tensor<256x!tt.ptr<f32>, #blocked0>, %arg1 : tensor<256xi1, #blocked0>, %arg2 : tensor<256xf32, #blocked0>) { 3884 + + func.func @atomic_add_f32(%arg0 : tensor<256x!tt.ptr<f32>, #blocked0>, %arg1 : tensor<256xi1, #blocked0>, %arg2 : tensor<256xf32, #blocked0>) { 3885 + // CHECK: llvm.inline_asm 3886 + // CHECK-SAME: atom.global.gpu.add.f32 3887 + %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> 3888 + @@ -1012,7 +1012,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3889 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3890 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3891 + 3892 + -func @test_get_program_id(%a: tensor<32x!tt.ptr<i32>, #blocked0>) { 3893 + +func.func @test_get_program_id(%a: tensor<32x!tt.ptr<i32>, #blocked0>) { 3894 + %blockidx = tt.get_program_id {axis=0:i32} : i32 3895 + %blockidy = tt.get_program_id {axis=1:i32} : i32 3896 + %blockidz = tt.get_program_id {axis=2:i32} : i32 3897 + @@ -1032,7 +1032,7 @@ func @test_get_program_id(%a: tensor<32x!tt.ptr<i32>, #blocked0>) { 3898 + // ----- 3899 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> 3900 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3901 + - func @test_get_num_program(%a: tensor<32x!tt.ptr<i32>, #blocked0>) { 3902 + + func.func @test_get_num_program(%a: tensor<32x!tt.ptr<i32>, #blocked0>) { 3903 + // CHECK: nvvm.read.ptx.sreg.nctaid.x 3904 + // CHECK: nvvm.read.ptx.sreg.nctaid.y 3905 + // CHECK: nvvm.read.ptx.sreg.nctaid.z 3906 + @@ -1052,7 +1052,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3907 + #blocked0 = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> 3908 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3909 + // CHECK-LABEL: test_index_cache 3910 + - func @test_index_cache() { 3911 + + func.func @test_index_cache() { 3912 + // CHECK: nvvm.read.ptx.sreg.tid.x 3913 + %0 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32, #blocked0> 3914 + // CHECK-NOT: nvvm.read.ptx.sreg.tid.x 3915 + @@ -1066,7 +1066,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 3916 + #shared0 = #triton_gpu.shared<{vec = 8, perPhase = 2, maxPhase = 4, order = [1, 0]}> 3917 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3918 + // CHECK-LABEL: test_base_index_cache 3919 + - func @test_base_index_cache(%arg0: tensor<128x32xf32, #blocked0>) { 3920 + + func.func @test_base_index_cache(%arg0: tensor<128x32xf32, #blocked0>) { 3921 + // CHECK: nvvm.read.ptx.sreg.tid.x 3922 + %0 = triton_gpu.convert_layout %arg0 : (tensor<128x32xf32, #blocked0>) -> tensor<128x32xf32, #shared0> 3923 + // CHECK-NOT: nvvm.read.ptx.sreg.tid.x 3924 + @@ -1080,7 +1080,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { 3925 + #shared0 = #triton_gpu.shared<{vec = 8, perPhase = 2, maxPhase = 4, order = [1, 0]}> 3926 + module attributes {"triton_gpu.num-warps" = 1 : i32} { 3927 + // CHECK-LABEL: test_index_cache_different_block 3928 + - func @test_index_cache_different_block(%arg0: tensor<128x32xf32, #blocked0>, %arg1: i1) { 3929 + + func.func @test_index_cache_different_block(%arg0: tensor<128x32xf32, #blocked0>, %arg1: i1) { 3930 + // CHECK: nvvm.read.ptx.sreg.tid.x 3931 + %0 = triton_gpu.convert_layout %arg0 : (tensor<128x32xf32, #blocked0>) -> tensor<128x32xf32, #shared0> 3932 + scf.if %arg1 { 3933 + diff --git a/test/Target/tritongpu_to_llvmir.mlir b/test/Target/tritongpu_to_llvmir.mlir 3934 + index cafff3ca60..114d3a9eb2 100644 3935 + --- a/test/Target/tritongpu_to_llvmir.mlir 3936 + +++ b/test/Target/tritongpu_to_llvmir.mlir 3937 + @@ -4,11 +4,11 @@ 3938 + // CHECK-LABEL: ; ModuleID = 'LLVMDialectModule' 3939 + // CHECK: define void @test_empty_kernel 3940 + // CHECK: !nvvm.annotations 3941 + -// CHECK: !{void (i32, half addrspace(1)*)* @test_empty_kernel, !"maxntidx", i32 128} 3942 + +// CHECK: !{ptr @test_empty_kernel, !"maxntidx", i32 128} 3943 + 3944 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3945 + 3946 + -func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) { 3947 + +func.func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) { 3948 + 3949 + return 3950 + } 3951 + diff --git a/test/Target/tritongpu_to_ptx.mlir b/test/Target/tritongpu_to_ptx.mlir 3952 + index 404e970a29..12742ad9e2 100644 3953 + --- a/test/Target/tritongpu_to_ptx.mlir 3954 + +++ b/test/Target/tritongpu_to_ptx.mlir 3955 + @@ -6,7 +6,7 @@ 3956 + 3957 + module attributes {"triton_gpu.num-warps" = 4 : i32} { 3958 + 3959 + -func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) { 3960 + +func.func @test_empty_kernel(%lb : index, %A : !tt.ptr<f16>) { 3961 + 3962 + return 3963 + } 3964 + diff --git a/test/Triton/combine.mlir b/test/Triton/combine.mlir 3965 + index 050a3f7565..5ef6790e69 100644 3966 + --- a/test/Triton/combine.mlir 3967 + +++ b/test/Triton/combine.mlir 3968 + @@ -2,10 +2,10 @@ 3969 + // RUN: triton-opt %s -split-input-file -canonicalize -triton-combine | FileCheck %s 3970 + 3971 + // CHECK-LABEL: @test_combine_dot_add_pattern 3972 + -func @test_combine_dot_add_pattern() -> (tensor<128x128xf32>, tensor<128x128xf32>) { 3973 + - // CHECK: %[[d:.*]] = arith.constant dense<3.000000e+00> : tensor<128x128xf32> 3974 + - // CHECK: %[[b:.*]] = arith.constant dense<2.000000e+00> : tensor<128x128xf32> 3975 + - // CHECK: %[[a:.*]] = arith.constant dense<1.000000e+00> : tensor<128x128xf32> 3976 + +func.func @test_combine_dot_add_pattern() -> (tensor<128x128xf32>, tensor<128x128xf32>) { 3977 + + // CHECK-DAG: %[[d:.*]] = arith.constant dense<3.000000e+00> : tensor<128x128xf32> 3978 + + // CHECK-DAG: %[[b:.*]] = arith.constant dense<2.000000e+00> : tensor<128x128xf32> 3979 + + // CHECK-DAG: %[[a:.*]] = arith.constant dense<1.000000e+00> : tensor<128x128xf32> 3980 + %a = arith.constant dense<1.0> : tensor<128x128xf32> 3981 + %b = arith.constant dense<2.0> : tensor<128x128xf32> 3982 + %zero = arith.constant dense<0.0> : tensor<128x128xf32> 3983 + @@ -24,7 +24,7 @@ func @test_combine_dot_add_pattern() -> (tensor<128x128xf32>, tensor<128x128xf32 3984 + 3985 + 3986 + // COM: CHECK-LABEL: @test_combine_addptr_pattern 3987 + -func @test_combine_addptr_pattern(%base: !tt.ptr<f32>) -> tensor<8x!tt.ptr<f32>> { 3988 + +func.func @test_combine_addptr_pattern(%base: !tt.ptr<f32>) -> tensor<8x!tt.ptr<f32>> { 3989 + %off0 = arith.constant 10 : i32 3990 + %off1 = arith.constant 15 : i32 3991 + 3992 + @@ -47,46 +47,46 @@ func @test_combine_addptr_pattern(%base: !tt.ptr<f32>) -> tensor<8x!tt.ptr<f32>> 3993 + 3994 + 3995 + // CHECK-LABEL: @test_combine_select_masked_load_pattern 3996 + -func @test_combine_select_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %cond: i1) -> (tensor<8xf32>, tensor<8xf32>) { 3997 + +func.func @test_combine_select_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %cond: i1) -> (tensor<8xf32>, tensor<8xf32>) { 3998 + %mask = tt.broadcast %cond : (i1) -> tensor<8xi1> 3999 + %false_val = arith.constant dense<0.0> : tensor<8xf32> 4000 + 4001 + // CHECK: %[[res1:.*]] = tt.load %{{.*}}, %{{.*}}, %{{.*}} {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> 4002 + %x = tt.load %ptr, %mask, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> 4003 + - %0 = select %cond, %x, %false_val : tensor<8xf32> 4004 + + %0 = arith.select %cond, %x, %false_val : tensor<8xf32> 4005 + 4006 + // CHECK: %[[res2:.*]] = tt.load %{{.*}}, %{{.*}}, %{{.*}} {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> 4007 + %y = tt.load %ptr, %mask, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> 4008 + - %1 = select %cond, %y, %false_val : tensor<8xf32> 4009 + + %1 = arith.select %cond, %y, %false_val : tensor<8xf32> 4010 + 4011 + // CHECK: return %[[res1]], %[[res2]] : tensor<8xf32>, tensor<8xf32> 4012 + return %0, %1 : tensor<8xf32>, tensor<8xf32> 4013 + } 4014 + 4015 + // CHECK-LABEL: @test_combine_select_masked_load_fail_pattern 4016 + -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>) { 4017 + +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>) { 4018 + %false_val = arith.constant dense<0.0> : tensor<8xf32> 4019 + 4020 + // Case 1: value at the "load" position is not an "op". Select should not be canonicalized. 4021 + - // CHECK: %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> 4022 + - %0 = select %cond0, %dummy_load, %false_val : tensor<8xf32> 4023 + + // CHECK: %{{.*}} = arith.select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> 4024 + + %0 = arith.select %cond0, %dummy_load, %false_val : tensor<8xf32> 4025 + 4026 + // Case 2: value at the "broadcast" position is not an "op". Select should not be canonicalized. 4027 + %real_load0 = tt.load %ptr, %dummy_broadcast, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> 4028 + - // CHECK: %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> 4029 + - %1 = select %cond0, %real_load0, %false_val : tensor<8xf32> 4030 + + // CHECK: %{{.*}} = arith.select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> 4031 + + %1 = arith.select %cond0, %real_load0, %false_val : tensor<8xf32> 4032 + 4033 + // Case 3: condition of "broadcast" is not the same as the condition of "select". Select should not be canonicalized. 4034 + %cond0_ = tt.broadcast %cond0 : (i1) -> tensor<8xi1> 4035 + %real_load1 = tt.load %ptr, %cond0_, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> 4036 + - // CHECK: %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> 4037 + - %2 = select %cond1, %real_load1, %false_val : tensor<8xf32> 4038 + + // CHECK: %{{.*}} = arith.select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> 4039 + + %2 = arith.select %cond1, %real_load1, %false_val : tensor<8xf32> 4040 + 4041 + return %0, %1, %2 : tensor<8xf32>, tensor<8xf32>, tensor<8xf32> 4042 + } 4043 + 4044 + // CHECK-LABEL: @test_combine_broadcast_constant_pattern 4045 + -func @test_combine_broadcast_constant_pattern(%cst : f32) -> tensor<8x2xf32> { 4046 + +func.func @test_combine_broadcast_constant_pattern(%cst : f32) -> tensor<8x2xf32> { 4047 + // CHECK: %[[cst:.*]] = arith.constant dense<1.000000e+00> : tensor<8x2xf32> 4048 + %const = arith.constant dense<1.0> : tensor<8xf32> 4049 + %bst_out = tt.broadcast %const : (tensor<8xf32>) -> tensor<8x2xf32> 4050 + @@ -96,7 +96,7 @@ func @test_combine_broadcast_constant_pattern(%cst : f32) -> tensor<8x2xf32> { 4051 + } 4052 + 4053 + // CHECK-LABEL: @test_canonicalize_masked_load_pattern 4054 + -func @test_canonicalize_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) { 4055 + +func.func @test_canonicalize_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) { 4056 + %true_mask = arith.constant dense<true> : tensor<8xi1> 4057 + %false_mask = arith.constant dense<false> : tensor<8xi1> 4058 + %other_val = arith.constant dense<0.0> : tensor<8xf32> 4059 + @@ -117,7 +117,7 @@ func @test_canonicalize_masked_load_pattern(%ptr: tensor<8x!tt.ptr<f32>>) -> (te 4060 + } 4061 + 4062 + // CHECK-LABEL: @test_canonicalize_masked_load_fail_pattern 4063 + -func @test_canonicalize_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %mask: tensor<8xi1>) -> (tensor<8xf32>, tensor<8xf32>) { 4064 + +func.func @test_canonicalize_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %mask: tensor<8xi1>) -> (tensor<8xf32>, tensor<8xf32>) { 4065 + %other_val = arith.constant dense<0.0> : tensor<8xf32> 4066 + 4067 + // Case: value at the "mask" position is not an "op". Load should not be canonicalized. 4068 + @@ -130,7 +130,7 @@ func @test_canonicalize_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, % 4069 + } 4070 + 4071 + // CHECK-LABEL: @test_canonicalize_masked_store_pattern 4072 + -func @test_canonicalize_masked_store_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val: tensor<8xf32>) { 4073 + +func.func @test_canonicalize_masked_store_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val: tensor<8xf32>) { 4074 + %true_mask = arith.constant dense<true> : tensor<8xi1> 4075 + %false_mask = arith.constant dense<false> : tensor<8xi1> 4076 + 4077 + @@ -144,7 +144,7 @@ func @test_canonicalize_masked_store_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val: 4078 + } 4079 + 4080 + // CHECK-LABEL: @test_canonicalize_masked_store_fail_pattern 4081 + -func @test_canonicalize_masked_store_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val: tensor<8xf32>, %mask: tensor<8xi1>) { 4082 + +func.func @test_canonicalize_masked_store_fail_pattern(%ptr: tensor<8x!tt.ptr<f32>>, %val: tensor<8xf32>, %mask: tensor<8xi1>) { 4083 + // Case: value at the "mask" position is not an "op". Store should not be canonicalized. 4084 + // CHECK: tt.store %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> 4085 + tt.store %ptr, %val, %mask : tensor<8xf32> 4086 + diff --git a/test/Triton/vecadd.mlir b/test/Triton/vecadd.mlir 4087 + index 0b69ef3054..f5019b1cdd 100644 4088 + --- a/test/Triton/vecadd.mlir 4089 + +++ b/test/Triton/vecadd.mlir 4090 + @@ -1,7 +1,7 @@ 4091 + // RUN: triton-opt %s -verify-diagnostics 4092 + 4093 + module { 4094 + - 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) { 4095 + + 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) { 4096 + %0 = tt.get_program_id {axis = 0 : i32} : i32 4097 + %c256_i32 = arith.constant 256 : i32 4098 + %1 = arith.muli %0, %c256_i32 : i32 4099 + @@ -43,7 +43,7 @@ module { 4100 + } 4101 + } 4102 + // module { 4103 + -// 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) { 4104 + +// 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) { 4105 + // %c64 = arith.constant 64 : index 4106 + // %c32 = arith.constant 32 : index 4107 + // %c0 = arith.constant 0 : index 4108 + diff --git a/test/TritonGPU/coalesce.mlir b/test/TritonGPU/coalesce.mlir 4109 + index 60e359f527..51cccccfbd 100644 4110 + --- a/test/TritonGPU/coalesce.mlir 4111 + +++ b/test/TritonGPU/coalesce.mlir 4112 + @@ -19,7 +19,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { 4113 + // CHECK: [[store_val:%.*]] = triton_gpu.convert_layout {{.*}} -> tensor<64x64xf32, [[col_layout]]> 4114 + // CHECK: [[store_mask:%.*]] = triton_gpu.convert_layout {{.*}} -> tensor<64x64xi1, [[col_layout]]> 4115 + // CHECK: tt.store [[store_ptr]], [[store_val]], [[store_mask]] 4116 + -func @transpose(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, 4117 + +func.func @transpose(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, 4118 + %arg1: i32 {tt.divisibility = 16 : i32}, 4119 + %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, 4120 + %arg3: i32 {tt.divisibility = 16 : i32}) { 4121 + diff --git a/test/TritonGPU/combine.mlir b/test/TritonGPU/combine.mlir 4122 + index 2c009ffa48..7e9cb9d504 100644 4123 + --- a/test/TritonGPU/combine.mlir 4124 + +++ b/test/TritonGPU/combine.mlir 4125 + @@ -9,7 +9,7 @@ 4126 + // CHECK: [[col_layout:#.*]] = #triton_gpu.blocked<{sizePerThread = [4, 1], threadsPerWarp = [16, 2], warpsPerCTA = [4, 1], order = [0, 1]}> 4127 + // CHECK: [[col_layout_novec:#.*]] = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [4, 1], order = [0, 1]}> 4128 + // CHECK-LABEL: cst 4129 + -func @cst() -> tensor<1024xi32, #layout1> { 4130 + +func.func @cst() -> tensor<1024xi32, #layout1> { 4131 + %cst = arith.constant dense<0> : tensor<1024xi32, #layout0> 4132 + %1 = triton_gpu.convert_layout %cst : (tensor<1024xi32, #layout0>) -> tensor<1024xi32, #layout1> 4133 + // CHECK-NOT: triton_gpu.convert_layout 4134 + @@ -18,7 +18,7 @@ func @cst() -> tensor<1024xi32, #layout1> { 4135 + } 4136 + 4137 + // CHECK-LABEL: range 4138 + -func @range() -> tensor<1024xi32, #layout1> { 4139 + +func.func @range() -> tensor<1024xi32, #layout1> { 4140 + %0 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #layout0> 4141 + %1 = triton_gpu.convert_layout %0 : (tensor<1024xi32, #layout0>) -> tensor<1024xi32, #layout1> 4142 + // CHECK-NOT: triton_gpu.convert_layout 4143 + @@ -27,7 +27,7 @@ func @range() -> tensor<1024xi32, #layout1> { 4144 + } 4145 + 4146 + // CHECK-LABEL: splat 4147 + -func @splat(%arg0: i32) -> tensor<1024xi32, #layout1> { 4148 + +func.func @splat(%arg0: i32) -> tensor<1024xi32, #layout1> { 4149 + %0 = tt.splat %arg0 : (i32) -> tensor<1024xi32, #layout0> 4150 + %1 = triton_gpu.convert_layout %0 : (tensor<1024xi32, #layout0>) -> tensor<1024xi32, #layout1> 4151 + // CHECK-NOT: triton_gpu.convert_layout 4152 + @@ -36,7 +36,7 @@ func @splat(%arg0: i32) -> tensor<1024xi32, #layout1> { 4153 + } 4154 + 4155 + // CHECK-LABEL: remat 4156 + -func @remat(%arg0: i32) -> tensor<1024xi32, #layout1> { 4157 + +func.func @remat(%arg0: i32) -> tensor<1024xi32, #layout1> { 4158 + %0 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #layout0> 4159 + %1 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #layout0> 4160 + %2 = arith.muli %0, %1 : tensor<1024xi32, #layout0> 4161 + @@ -56,7 +56,7 @@ func @remat(%arg0: i32) -> tensor<1024xi32, #layout1> { 4162 + } 4163 + 4164 + // CHECK-LABEL: remat_load_store 4165 + -func @remat_load_store(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4166 + +func.func @remat_load_store(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4167 + %0 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #layout0> 4168 + %1 = tt.splat %arg : (!tt.ptr<i32>) -> tensor<64x!tt.ptr<i32>, #layout0> 4169 + %2 = tt.addptr %1, %0 : tensor<64x!tt.ptr<i32>, #layout0>, tensor<64xi32, #layout0> 4170 + @@ -70,7 +70,7 @@ func @remat_load_store(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4171 + 4172 + // Don't rematerialize vectorized loads 4173 + // CHECK-LABEL: remat_expensive 4174 + -func @remat_expensive(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4175 + +func.func @remat_expensive(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4176 + %0 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #layout1> 4177 + %1 = tt.splat %arg : (!tt.ptr<i32>) -> tensor<64x!tt.ptr<i32>, #layout1> 4178 + %2 = tt.addptr %1, %0 : tensor<64x!tt.ptr<i32>, #layout1>, tensor<64xi32, #layout1> 4179 + @@ -85,7 +85,7 @@ func @remat_expensive(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4180 + 4181 + // Don't rematerialize loads when original and target layouts are different 4182 + // CHECK-LABEL: remat_multi_layout 4183 + -func @remat_multi_layout(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4184 + +func.func @remat_multi_layout(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4185 + %0 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #layout0> 4186 + %1 = tt.splat %arg : (!tt.ptr<i32>) -> tensor<64x!tt.ptr<i32>, #layout0> 4187 + %2 = tt.addptr %1, %0 : tensor<64x!tt.ptr<i32>, #layout0>, tensor<64xi32, #layout0> 4188 + @@ -100,7 +100,7 @@ func @remat_multi_layout(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4189 + 4190 + // Always rematerialize single value loads 4191 + // CHECK-LABEL: remat_single_value 4192 + -func @remat_single_value(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4193 + +func.func @remat_single_value(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4194 + %0 = tt.splat %arg : (!tt.ptr<i32>) -> tensor<1x!tt.ptr<i32>, #layout1> 4195 + %1 = tt.load %0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<1xi32, #layout1> 4196 + // CHECK-NOT: triton_gpu.convert_layout 4197 + @@ -111,7 +111,7 @@ func @remat_single_value(%arg: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4198 + } 4199 + 4200 + // CHECK-LABEL: if 4201 + -func @if(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4202 + +func.func @if(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4203 + // CHECK-NOT: triton_gpu.convert_layout 4204 + %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout1> 4205 + %0 = tt.get_program_id {axis = 0 : i32} : i32 4206 + @@ -128,7 +128,7 @@ func @if(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4207 + } 4208 + 4209 + // CHECK-LABEL: if_convert_else_not 4210 + -func @if_convert_else_not(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4211 + +func.func @if_convert_else_not(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4212 + %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout0> 4213 + %0 = tt.get_program_id {axis = 0 : i32} : i32 4214 + %1 = tt.splat %0 : (i32) -> tensor<1024xi32, #layout0> 4215 + @@ -149,7 +149,7 @@ func @if_convert_else_not(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 4216 + } 4217 + 4218 + // CHECK-LABEL: if_not_else_convert 4219 + -func @if_not_else_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4220 + +func.func @if_not_else_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4221 + %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout0> 4222 + %0 = tt.get_program_id {axis = 0 : i32} : i32 4223 + %1 = tt.splat %0 : (i32) -> tensor<1024xi32, #layout0> 4224 + @@ -170,7 +170,7 @@ func @if_not_else_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 4225 + } 4226 + 4227 + // CHECK-LABEL: if_else_both_convert 4228 + -func @if_else_both_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4229 + +func.func @if_else_both_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 : i32}) { 4230 + %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout0> 4231 + %0 = tt.get_program_id {axis = 0 : i32} : i32 4232 + %1 = tt.splat %0 : (i32) -> tensor<1024xi32, #layout0> 4233 + @@ -200,7 +200,7 @@ func @if_else_both_convert(%arg0: i32, %arg1: !tt.ptr<i32> {tt.divisibility = 16 4234 + #blocked4 = #triton_gpu.blocked<{sizePerThread = [4, 1], threadsPerWarp = [16, 2], warpsPerCTA = [4, 1], order = [0, 1]}> 4235 + 4236 + // CHECK-LABEL: transpose 4237 + -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}) { 4238 + +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}) { 4239 + // CHECK-NOT: triton_gpu.convert_layout 4240 + // CHECK: [[loaded_val:%.*]] = tt.load {{.*}}, {{%cst.*}}, {{%cst.*}} {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64x64xf32, [[row_layout]]> 4241 + // CHECK: [[cvt_val:%.*]] = triton_gpu.convert_layout [[loaded_val]] : (tensor<64x64xf32, [[row_layout]]>) -> tensor<64x64xf32, [[col_layout]]> 4242 + @@ -241,7 +241,7 @@ func @transpose(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: i32 {tt 4243 + } 4244 + 4245 + // CHECK-LABEL: loop 4246 + -func @loop(%arg0: !tt.ptr<f32>, %arg1: i32, %arg2: !tt.ptr<f32>, %arg3: i32, %arg4: i32) { 4247 + +func.func @loop(%arg0: !tt.ptr<f32>, %arg1: i32, %arg2: !tt.ptr<f32>, %arg3: i32, %arg4: i32) { 4248 + // CHECK-NOT: triton_gpu.convert_layout 4249 + // CHECK: [[loop_ret:%.*]]:2 = scf.for {{.*}} -> (tensor<64x64xf32, [[row_layout]]>, tensor<64x64x!tt.ptr<f32>, [[row_layout]]>) 4250 + // CHECK-NEXT: {{.*}} = tt.load {{.*}} : tensor<64x64xf32, [[row_layout]]> 4251 + @@ -295,7 +295,7 @@ func @loop(%arg0: !tt.ptr<f32>, %arg1: i32, %arg2: !tt.ptr<f32>, %arg3: i32, %ar 4252 + } 4253 + 4254 + // CHECK-LABEL: vecadd 4255 + -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) { 4256 + +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) { 4257 + // CHECK-NOT: triton_gpu.convert_layout 4258 + %c256_i32 = arith.constant 256 : i32 4259 + %0 = tt.get_program_id {axis = 0 : i32} : i32 4260 + @@ -327,7 +327,7 @@ func @vecadd(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f3 4261 + 4262 + // Select has args with different element types 4263 + // CHECK-LABEL: select 4264 + -func @select(%arg0: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}) { 4265 + +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}) { 4266 + // CHECK-NOT: triton_gpu.convert_layout 4267 + %cst = arith.constant dense<30000> : tensor<1x1xi32, #blocked2> 4268 + %cst_0 = arith.constant dense<30000> : tensor<1x512xi32, #blocked2> 4269 + @@ -378,7 +378,7 @@ func @select(%arg0: !tt.ptr<f64> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f6 4270 + 4271 + // Make sure the following IR doesn't hang the compiler. 4272 + // CHECK-LABEL: long_func 4273 + -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}) { 4274 + +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}) { 4275 + %cst = arith.constant dense<1.000000e+00> : tensor<1024xf32, #blocked0> 4276 + %cst_0 = arith.constant dense<5.000000e-04> : tensor<1024xf32, #blocked0> 4277 + %cst_1 = arith.constant dense<0.999499976> : tensor<1024xf32, #blocked0> 4278 + @@ -775,7 +775,7 @@ func public @long_func(%arg0: !tt.ptr<i64> {tt.divisibility = 16 : i32}, %arg1: 4279 + // A mnist model from torch inductor. 4280 + // Check if topological sort is working correct and there's no unnecessary convert 4281 + // CHECK-LABEL: mnist 4282 + -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) { 4283 + +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) { 4284 + // CHECK-NOT: triton_gpu.convert_layout 4285 + %cst = arith.constant dense<10> : tensor<16x1xi32, #blocked2> 4286 + %cst_0 = arith.constant dense<10> : tensor<1x16xi32, #blocked3> 4287 + @@ -862,7 +862,7 @@ func public @mnist(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg1: !tt. 4288 + #blocked5 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [2, 16], warpsPerCTA = [8, 1], order = [1, 0]}> 4289 + // cmpf and cmpi have different operands and result types 4290 + // CHECK-LABEL: cmp 4291 + -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}) { 4292 + +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}) { 4293 + %c64 = arith.constant 64 : index 4294 + %c2048 = arith.constant 2048 : index 4295 + %c0 = arith.constant 0 : index 4296 + diff --git a/test/TritonGPU/loop-pipeline.mlir b/test/TritonGPU/loop-pipeline.mlir 4297 + index 6ee3b15fbc..663f2da7b0 100644 4298 + --- a/test/TritonGPU/loop-pipeline.mlir 4299 + +++ b/test/TritonGPU/loop-pipeline.mlir 4300 + @@ -10,7 +10,7 @@ 4301 + #A = #triton_gpu.dot_op<{opIdx = 0, parent = #C}> 4302 + #B = #triton_gpu.dot_op<{opIdx = 1, parent = #C}> 4303 + 4304 + -// CHECK: func @matmul_loop 4305 + +// CHECK: func.func @matmul_loop 4306 + // CHECK-DAG: %[[CONSTANT_0:.*]] = arith.constant 0 : i32 4307 + // CHECK-DAG: %[[CONSTANT_1:.*]] = arith.constant 1 : i32 4308 + // CHECK-DAG: %[[CONSTANT_2:.*]] = arith.constant 2 : i32 4309 + @@ -46,8 +46,8 @@ 4310 + // CHECK-DAG: %[[NEXT_PIPELINE_IDX:.*]] = arith.addi %[[PIPELINE_IDX]], %[[CONSTANT_1]] 4311 + // CHECK-DAG: %[[NEXT_LOOP_IDX:.*]] = arith.addi %[[LOOP_IDX]], %[[CONSTANT_1]] 4312 + // CHECK: scf.yield {{.*}}, {{.*}}, {{.*}}, %[[NEXT_A_BUFFER]], %[[NEXT_B_BUFFER]], %[[NEXT_A]], %[[NEXT_B]], {{.*}}, {{.*}}, {{.*}}, %[[NEXT_PIPELINE_IDX]], %[[NEXT_LOOP_IDX]] 4313 + -func @matmul_loop(%lb : index, %ub : index, %step : index, 4314 + - %A : !tt.ptr<f16> {tt.divisibility = 16 : i32}, 4315 + +func.func @matmul_loop(%lb : index, %ub : index, %step : index, 4316 + + %A : !tt.ptr<f16> {tt.divisibility = 16 : i32}, 4317 + %B : !tt.ptr<f16> {tt.divisibility = 16 : i32}) { 4318 + // A ptrs 4319 + %a_ptr_splat = tt.splat %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL> 4320 + @@ -61,7 +61,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, 4321 + %b_tmp1 = tt.expand_dims %b_tmp0 {axis = 0 : i32} : (tensor<128xi32, #BLs0>) -> tensor<1x128xi32, #BL> 4322 + %b_offs = tt.broadcast %b_tmp1 : (tensor<1x128xi32, #BL>) -> tensor<32x128xi32, #BL> 4323 + %b_ptr_init = tt.addptr %b_ptr_splat, %b_offs : tensor<32x128x!tt.ptr<f16>, #BL>, tensor<32x128xi32, #BL> 4324 + - 4325 + + 4326 + 4327 + %a_mask = arith.constant dense<true> : tensor<128x32xi1, #AL> 4328 + %a_other = arith.constant dense<0.00e+00> : tensor<128x32xf16, #AL> 4329 + @@ -88,7 +88,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, 4330 + } 4331 + 4332 + 4333 + -// CHECK: func @matmul_loop_nested 4334 + +// CHECK: func.func @matmul_loop_nested 4335 + // CHECK-DAG: %[[CONSTANT_0:.*]] = arith.constant 0 : i32 4336 + // CHECK-DAG: %[[CONSTANT_1:.*]] = arith.constant 1 : i32 4337 + // CHECK-DAG: %[[CONSTANT_2:.*]] = arith.constant 2 : i32 4338 + @@ -118,8 +118,8 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, 4339 + // CHECK-DAG: %[[NEXT_PIPELINE_IDX:.*]] = arith.addi %[[PIPELINE_IDX]], %[[CONSTANT_1]] 4340 + // CHECK-DAG: %[[NEXT_LOOP_IDX:.*]] = arith.addi %[[LOOP_IDX]], %[[CONSTANT_1]] 4341 + // CHECK: scf.yield {{.*}}, {{.*}}, {{.*}}, %[[NEXT_A_BUFFER]], %[[NEXT_B_BUFFER]], %[[NEXT_A]], %[[NEXT_B]], {{.*}}, {{.*}}, {{.*}}, %[[NEXT_PIPELINE_IDX]], %[[NEXT_LOOP_IDX]] 4342 + -func @matmul_loop_nested(%lb : index, %ub : index, %step : index, 4343 + - %A : !tt.ptr<f16> {tt.divisibility = 16 : i32}, 4344 + +func.func @matmul_loop_nested(%lb : index, %ub : index, %step : index, 4345 + + %A : !tt.ptr<f16> {tt.divisibility = 16 : i32}, 4346 + %B : !tt.ptr<f16> {tt.divisibility = 16 : i32}) { 4347 + scf.for %iv0 = %lb to %ub step %step { 4348 + // A ptrs 4349 + @@ -134,7 +134,7 @@ func @matmul_loop_nested(%lb : index, %ub : index, %step : index, 4350 + %b_tmp1 = tt.expand_dims %b_tmp0 {axis = 0 : i32} : (tensor<128xi32, #BLs0>) -> tensor<1x128xi32, #BL> 4351 + %b_offs = tt.broadcast %b_tmp1 : (tensor<1x128xi32, #BL>) -> tensor<32x128xi32, #BL> 4352 + %b_ptr_init = tt.addptr %b_ptr_splat, %b_offs : tensor<32x128x!tt.ptr<f16>, #BL>, tensor<32x128xi32, #BL> 4353 + - 4354 + + 4355 + %a_mask = arith.constant dense<true> : tensor<128x32xi1, #AL> 4356 + %a_other = arith.constant dense<0.00e+00> : tensor<128x32xf16, #AL> 4357 + %b_mask = arith.constant dense<true> : tensor<32x128xi1, #BL> 4358 + @@ -161,7 +161,7 @@ func @matmul_loop_nested(%lb : index, %ub : index, %step : index, 4359 + } 4360 + 4361 + 4362 + -// CHECK: func @matmul_loop_single_pipeline 4363 + +// CHECK: func.func @matmul_loop_single_pipeline 4364 + // CHECK-DAG: %[[CONSTANT_0:.*]] = arith.constant 0 : i32 4365 + // CHECK-DAG: %[[CONSTANT_1:.*]] = arith.constant 1 : i32 4366 + // CHECK-DAG: %[[CONSTANT_2:.*]] = arith.constant 2 : i32 4367 + @@ -183,8 +183,8 @@ func @matmul_loop_nested(%lb : index, %ub : index, %step : index, 4368 + // CHECK-DAG: %[[NEXT_PIPELINE_IDX:.*]] = arith.addi %[[PIPELINE_IDX]], %[[CONSTANT_1]] 4369 + // CHECK-DAG: %[[NEXT_LOOP_IDX:.*]] = arith.addi %[[LOOP_IDX]], %[[CONSTANT_1]] 4370 + // CHECK: scf.yield {{.*}}, {{.*}}, %[[NEXT_B_BUFFER]], %[[NEXT_B]], {{.*}}, {{.*}}, %[[NEXT_PIPELINE_IDX]], %[[NEXT_LOOP_IDX]] 4371 + -func @matmul_loop_single_pipeline(%lb : index, %ub : index, %step : index, 4372 + - %A : !tt.ptr<f16> {tt.divisibility = 16 : i32}, 4373 + +func.func @matmul_loop_single_pipeline(%lb : index, %ub : index, %step : index, 4374 + + %A : !tt.ptr<f16> {tt.divisibility = 16 : i32}, 4375 + %B : !tt.ptr<f16> {tt.divisibility = 16 : i32}) { 4376 + // A ptrs 4377 + %a_ptr_splat = tt.splat %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL> 4378 + diff --git a/test/TritonGPU/matmul.mlir b/test/TritonGPU/matmul.mlir 4379 + index 9bd5318e1e..01dc3f0ab1 100644 4380 + --- a/test/TritonGPU/matmul.mlir 4381 + +++ b/test/TritonGPU/matmul.mlir 4382 + @@ -4,7 +4,7 @@ 4383 + // CHECK: offset = 49152, size = 49152 4384 + // CHECK: size = 98304 4385 + module { 4386 + -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) { 4387 + +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) { 4388 + %cst = arith.constant dense<true> : tensor<64x64xi1> 4389 + %c64 = arith.constant 64 : index 4390 + %c0 = arith.constant 0 : index 4391 + @@ -22,7 +22,7 @@ func @matmul_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32_i32_i32_i32_i32_i32_i32__12c6 4392 + %7 = arith.muli %6, %c8_i32 : i32 4393 + %8 = arith.subi %2, %7 : i32 4394 + %9 = arith.cmpi slt, %8, %c8_i32 : i32 4395 + - %10 = select %9, %8, %c8_i32 : i32 4396 + + %10 = arith.select %9, %8, %c8_i32 : i32 4397 + %11 = arith.remsi %0, %10 : i32 4398 + %12 = arith.addi %7, %11 : i32 4399 + %13 = arith.remsi %0, %5 : i32 4400 + diff --git a/test/TritonGPU/prefetch.mlir b/test/TritonGPU/prefetch.mlir 4401 + index 52b4dddec1..b427547890 100644 4402 + --- a/test/TritonGPU/prefetch.mlir 4403 + +++ b/test/TritonGPU/prefetch.mlir 4404 + @@ -11,7 +11,7 @@ 4405 + #B_OP = #triton_gpu.dot_op<{opIdx = 1, parent = #C}> 4406 + 4407 + 4408 + -// CHECK: func @matmul_loop 4409 + +// CHECK: func.func @matmul_loop 4410 + // CHECK-DAG: %[[A0_PREFETCH_SMEM:.*]] = tensor.extract_slice %[[A0:.*]][0, 0] [128, 16] 4411 + // CHECK-DAG: %[[A0_PREFETCH:.*]] = triton_gpu.convert_layout %[[A0_PREFETCH_SMEM]] 4412 + // CHECK-DAG: %[[B0_PREFETCH_SMEM:.*]] = tensor.extract_slice %[[B0:.*]][0, 0] [16, 128] 4413 + @@ -28,7 +28,7 @@ 4414 + // CHECK-DAG: %[[NEXT_B_PREFETCH_SMEM:.*]] = tensor.extract_slice {{.*}}[0, 0] [16, 128] 4415 + // CHECK-DAG: %[[NEXT_B_PREFETCH:.*]] = triton_gpu.convert_layout %[[NEXT_B_PREFETCH_SMEM]] 4416 + // CHECK: scf.yield {{.*}}, {{.*}}, {{.*}}, {{.*}}, {{.*}}, %[[NEXT_A_PREFETCH]], %[[NEXT_B_PREFETCH]] 4417 + -func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 4418 + +func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr<f16>, %B : !tt.ptr<f16>) { 4419 + %a_ptr_init = tt.broadcast %A : (!tt.ptr<f16>) -> tensor<128x32x!tt.ptr<f16>, #AL> 4420 + %b_ptr_init = tt.broadcast %B : (!tt.ptr<f16>) -> tensor<32x128x!tt.ptr<f16>, #BL> 4421 + 4422 + diff --git a/test/TritonGPU/update-mma-for-volta.mlir b/test/TritonGPU/update-mma-for-volta.mlir 4423 + index d587fffcca..7571ec6185 100644 4424 + --- a/test/TritonGPU/update-mma-for-volta.mlir 4425 + +++ b/test/TritonGPU/update-mma-for-volta.mlir 4426 + @@ -15,7 +15,7 @@ 4427 + // CHECK: [[new_mma:#mma.*]] = #triton_gpu.mma<{versionMajor = 1, versionMinor = 3, warpsPerCTA = [4, 2]}> 4428 + module attributes {"triton_gpu.num-warps" = 16 : i32} { 4429 + // CHECK-LABEL: dot_mmav1 4430 + - func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> { 4431 + + func.func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> { 4432 + %C = arith.constant dense<0.000000e+00> : tensor<64x64xf32, #blocked0> 4433 + %AA = triton_gpu.convert_layout %A : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_a> 4434 + %BB = triton_gpu.convert_layout %B : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_b> 4435 + @@ -50,7 +50,7 @@ module attributes {"triton_gpu.num-warps" = 16 : i32} { 4436 + 4437 + module attributes {"triton_gpu.num-warps" = 16 : i32} { 4438 + // CHECK-LABEL: dot_mmav1 4439 + - func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> { 4440 + + func.func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> { 4441 + %C = arith.constant dense<0.000000e+00> : tensor<64x64xf32, #blocked0> 4442 + %AA = triton_gpu.convert_layout %A : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_a> 4443 + %BB = triton_gpu.convert_layout %B : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_b> 4444 + diff --git a/test/lib/Analysis/TestAlias.cpp b/test/lib/Analysis/TestAlias.cpp 4445 + index 88a4118fe9..3fd0cfd0d3 100644 4446 + --- a/test/lib/Analysis/TestAlias.cpp 4447 + +++ b/test/lib/Analysis/TestAlias.cpp 4448 + @@ -9,10 +9,10 @@ using namespace mlir; 4449 + namespace { 4450 + 4451 + struct TestAliasPass 4452 + - : public PassWrapper<TestAliasPass, OperationPass<FuncOp>> { 4453 + + : public PassWrapper<TestAliasPass, OperationPass<func::FuncOp>> { 4454 + + 4455 + + MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAliasPass); 4456 + 4457 + - // LLVM15+ 4458 + - // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAliasPass); 4459 + static void print(StringRef name, SmallVector<std::string, 4> &vals, 4460 + raw_ostream &os) { 4461 + if (vals.empty()) 4462 + @@ -39,23 +39,24 @@ struct TestAliasPass 4463 + auto opName = SymbolTable::getSymbolName(operation).getValue().str(); 4464 + os << opName << "\n"; 4465 + 4466 + - SharedMemoryAliasAnalysis analysis(&getContext()); 4467 + - analysis.run(operation); 4468 + + std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver(); 4469 + + SharedMemoryAliasAnalysis *analysis = 4470 + + solver->load<SharedMemoryAliasAnalysis>(); 4471 + + if (failed(solver->initializeAndRun(operation))) 4472 + + return signalPassFailure(); 4473 + 4474 + AsmState state(operation->getParentOfType<ModuleOp>()); 4475 + // Get operation ids of value's aliases 4476 + auto getAllocOpNames = [&](Value value) { 4477 + - LatticeElement<AliasInfo> *latticeElement = 4478 + - analysis.lookupLatticeElement(value); 4479 + + dataflow::Lattice<AliasInfo> *latticeElement = 4480 + + analysis->getLatticeElement(value); 4481 + SmallVector<std::string, 4> opNames; 4482 + - if (latticeElement) { 4483 + + if (latticeElement && !latticeElement->isUninitialized()) { 4484 + auto &info = latticeElement->getValue(); 4485 + - if (!info.getAllocs().empty()) { 4486 + - for (auto &alias : info.getAllocs()) { 4487 + - auto opName = 4488 + - getValueOperandName(alias.getDefiningOp()->getResult(0), state); 4489 + - opNames.push_back(std::move(opName)); 4490 + - } 4491 + + for (auto &alias : info.getAllocs()) { 4492 + + auto opName = 4493 + + getValueOperandName(alias.getDefiningOp()->getResult(0), state); 4494 + + opNames.push_back(std::move(opName)); 4495 + } 4496 + } 4497 + // Ensure deterministic output 4498 + diff --git a/test/lib/Analysis/TestAllocation.cpp b/test/lib/Analysis/TestAllocation.cpp 4499 + index 84108c4d36..35e42242bd 100644 4500 + --- a/test/lib/Analysis/TestAllocation.cpp 4501 + +++ b/test/lib/Analysis/TestAllocation.cpp 4502 + @@ -6,10 +6,9 @@ using namespace mlir; 4503 + namespace { 4504 + 4505 + struct TestAllocationPass 4506 + - : public PassWrapper<TestAllocationPass, OperationPass<FuncOp>> { 4507 + + : public PassWrapper<TestAllocationPass, OperationPass<func::FuncOp>> { 4508 + 4509 + - // LLVM15+ 4510 + - // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAllocationPass); 4511 + + MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAllocationPass); 4512 + 4513 + StringRef getArgument() const final { return "test-print-allocation"; } 4514 + StringRef getDescription() const final { 4515 + diff --git a/test/lib/Analysis/TestAxisInfo.cpp b/test/lib/Analysis/TestAxisInfo.cpp 4516 + index a5205bb0a0..22347c32f0 100644 4517 + --- a/test/lib/Analysis/TestAxisInfo.cpp 4518 + +++ b/test/lib/Analysis/TestAxisInfo.cpp 4519 + @@ -1,25 +1,15 @@ 4520 + #include "mlir/Pass/Pass.h" 4521 + #include "triton/Analysis/AxisInfo.h" 4522 + +#include "triton/Analysis/Utility.h" 4523 + 4524 + using namespace mlir; 4525 + 4526 + namespace { 4527 + 4528 + struct TestAxisInfoPass 4529 + - : public PassWrapper<TestAxisInfoPass, OperationPass<FuncOp>> { 4530 + + : public PassWrapper<TestAxisInfoPass, OperationPass<func::FuncOp>> { 4531 + 4532 + - // LLVM15+ 4533 + - // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAlignmentPass); 4534 + - 4535 + - void print(const std::string &name, raw_ostream &os, ArrayRef<int64_t> vals) { 4536 + - os << name << ": ["; 4537 + - for (size_t d = 0; d < vals.size(); d++) { 4538 + - if (d != 0) 4539 + - os << ", "; 4540 + - os << vals[d]; 4541 + - } 4542 + - os << "]"; 4543 + - } 4544 + + MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAxisInfoPass); 4545 + 4546 + StringRef getArgument() const final { return "test-print-alignment"; } 4547 + StringRef getDescription() const final { 4548 + @@ -30,38 +20,19 @@ struct TestAxisInfoPass 4549 + Operation *operation = getOperation(); 4550 + auto &os = llvm::errs(); 4551 + auto opName = SymbolTable::getSymbolName(operation).getValue().str(); 4552 + - os << opName << "\n"; 4553 + - AxisInfoAnalysis analysis(&getContext()); 4554 + - analysis.run(operation); 4555 + + os << "@" << opName << "\n"; 4556 + + 4557 + + std::unique_ptr<DataFlowSolver> solver = createDataFlowSolver(); 4558 + + AxisInfoAnalysis *analysis = solver->load<AxisInfoAnalysis>(); 4559 + + if (failed(solver->initializeAndRun(operation))) 4560 + + return signalPassFailure(); 4561 + operation->walk([&](Operation *op) { 4562 + if (op->getNumResults() < 1) 4563 + return; 4564 + for (Value result : op->getResults()) { 4565 + - // std::ostringstream oss; 4566 + - // result.print(oss); 4567 + - // os << " => "; 4568 + - LatticeElement<AxisInfo> *latticeElement = 4569 + - analysis.lookupLatticeElement(result); 4570 + - if (!latticeElement) { 4571 + - os << "None\n"; 4572 + - return; 4573 + - } 4574 + - AxisInfo &info = latticeElement->getValue(); 4575 + - print("Contiguity", os, info.getContiguity()); 4576 + - os << " ; "; 4577 + - print("Divisibility", os, info.getDivisibility()); 4578 + - os << " ; "; 4579 + - print("Constancy", os, info.getConstancy()); 4580 + - os << " ; "; 4581 + - auto constantValue = info.getConstantValue(); 4582 + - os << "ConstantValue: ["; 4583 + - if (constantValue.has_value()) 4584 + - os << constantValue.value(); 4585 + - else 4586 + - os << "None"; 4587 + - os << "] ( "; 4588 + result.print(os); 4589 + - os << " ) "; 4590 + + os << " => "; 4591 + + analysis->getLatticeElement(result)->getValue().print(os); 4592 + os << "\n"; 4593 + } 4594 + }); 4595 + diff --git a/test/lib/Analysis/TestMembar.cpp b/test/lib/Analysis/TestMembar.cpp 4596 + index df4279fe24..ab9b9f3fb7 100644 4597 + --- a/test/lib/Analysis/TestMembar.cpp 4598 + +++ b/test/lib/Analysis/TestMembar.cpp 4599 + @@ -1,4 +1,4 @@ 4600 + -#include "mlir/Dialect/GPU/GPUDialect.h" 4601 + +#include "mlir/Dialect/GPU/IR/GPUDialect.h" 4602 + #include "mlir/IR/Dialect.h" 4603 + #include "mlir/Pass/Pass.h" 4604 + #include "triton/Analysis/Allocation.h" 4605 + @@ -9,10 +9,9 @@ using namespace mlir; 4606 + namespace { 4607 + 4608 + struct TestMembarPass 4609 + - : public PassWrapper<TestMembarPass, OperationPass<FuncOp>> { 4610 + + : public PassWrapper<TestMembarPass, OperationPass<func::FuncOp>> { 4611 + 4612 + - // LLVM15+ 4613 + - // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestMembarPass); 4614 + + MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestMembarPass); 4615 + 4616 + StringRef getArgument() const final { return "test-print-membar"; } 4617 + StringRef getDescription() const final {
+44 -18
pkgs/development/python-modules/torch/default.nix
··· 6 6 7 7 # Native build inputs 8 8 cmake, util-linux, linkFarm, symlinkJoin, which, pybind11, removeReferencesTo, 9 + pythonRelaxDepsHook, 9 10 10 11 # Build inputs 11 12 numactl, 12 13 Accelerate, CoreServices, libobjc, 13 14 14 15 # Propagated build inputs 16 + filelock, 17 + jinja2, 18 + networkx, 19 + openai-triton, 20 + sympy, 15 21 numpy, pyyaml, cffi, click, typing-extensions, 16 22 17 23 # Unit tests ··· 49 55 inherit (cudaPackages) cudatoolkit cudaFlags cudnn nccl; 50 56 in 51 57 52 - # assert that everything needed for cuda is present and that the correct cuda versions are used 53 - assert !cudaSupport || (let majorIs = lib.versions.major cudatoolkit.version; 54 - in majorIs == "9" || majorIs == "10" || majorIs == "11"); 58 + assert cudaSupport -> (cudaPackages.cudaMajorVersion == "11"); 55 59 56 60 # confirm that cudatoolkits are sync'd across dependencies 57 61 assert !(MPISupport && cudaSupport) || mpi.cudatoolkit == cudatoolkit; ··· 129 133 in buildPythonPackage rec { 130 134 pname = "torch"; 131 135 # Don't forget to update torch-bin to the same version. 132 - version = "1.13.1"; 136 + version = "2.0.0"; 133 137 format = "setuptools"; 134 138 135 - disabled = pythonOlder "3.7.0"; 139 + disabled = pythonOlder "3.8.0"; 136 140 137 141 outputs = [ 138 142 "out" # output standard python package ··· 145 149 repo = "pytorch"; 146 150 rev = "refs/tags/v${version}"; 147 151 fetchSubmodules = true; 148 - hash = "sha256-yQz+xHPw9ODRBkV9hv1th38ZmUr/fXa+K+d+cvmX3Z8="; 152 + hash = "sha256-cSw7+AYBUcZLz3UyK/+JWWjQxKwVBXcFvBq0XAcL3tE="; 149 153 }; 150 154 151 155 patches = lib.optionals (stdenv.isDarwin && stdenv.isx86_64) [ ··· 155 159 # base is 10.12. Until we upgrade, we can fall back on the older 156 160 # pthread support. 157 161 ./pthreadpool-disable-gcd.diff 158 - ] ++ [ 159 - # PyTorch fails to build on gcc 12 due to gloo 160 - # https://github.com/pytorch/pytorch/issues/77614 161 - (fetchpatch { 162 - url = "https://github.com/facebookincubator/gloo/commit/4a5e339b764261d20fc409071dc7a8b8989aa195.patch"; 163 - stripLen = 1; 164 - extraPrefix = "third_party/gloo/"; 165 - hash = "sha256-UxR1r7F6g76BWj3GBIrSy5t+YZDCWy6mMddwx+hon5w="; 166 - }) 167 162 ]; 168 163 169 164 postPatch = lib.optionalString rocmSupport '' ··· 261 256 # Suppress gcc regression: avx512 math function raises uninitialized variable warning 262 257 # https://gcc.gnu.org/bugzilla/show_bug.cgi?id=105593 263 258 # See also: Fails to compile with GCC 12.1.0 https://github.com/pytorch/pytorch/issues/77939 264 - ++ lib.optionals stdenv.cc.isGNU [ "-Wno-error=maybe-uninitialized" "-Wno-error=uninitialized" ])); 259 + ++ lib.optionals (stdenv.cc.isGNU && lib.versionAtLeast stdenv.cc.version "12.0.0") [ 260 + "-Wno-error=maybe-uninitialized" 261 + "-Wno-error=uninitialized" 262 + ] 263 + # Since pytorch 2.0: 264 + # gcc-12.2.0/include/c++/12.2.0/bits/new_allocator.h:158:33: error: ‘void operator delete(void*, std::size_t)’ 265 + # ... called on pointer ‘<unknown>’ with nonzero offset [1, 9223372036854775800] [-Werror=free-nonheap-object] 266 + ++ lib.optionals (stdenv.cc.isGNU && lib.versions.major stdenv.cc.version == "12" ) [ 267 + "-Wno-error=free-nonheap-object" 268 + ])); 265 269 266 270 nativeBuildInputs = [ 267 271 cmake ··· 269 273 which 270 274 ninja 271 275 pybind11 276 + pythonRelaxDepsHook 272 277 removeReferencesTo 273 278 ] ++ lib.optionals cudaSupport [ cudatoolkit_joined ] 274 279 ++ lib.optionals rocmSupport [ rocmtoolkit_joined ]; ··· 286 291 click 287 292 numpy 288 293 pyyaml 294 + 295 + # From install_requires: 296 + filelock 289 297 typing-extensions 298 + sympy 299 + networkx 300 + jinja2 301 + 290 302 # the following are required for tensorboard support 291 303 pillow six future tensorboard protobuf 292 - ] ++ lib.optionals MPISupport [ mpi ] 293 - ++ lib.optionals rocmSupport [ rocmtoolkit_joined ]; 304 + ] 305 + ++ lib.optionals MPISupport [ mpi ] 306 + ++ lib.optionals rocmSupport [ rocmtoolkit_joined ] 307 + # rocm build requires openai-triton; 308 + # openai-triton currently requires cuda_nvcc, 309 + # so not including it in the cpu-only build; 310 + # torch.compile relies on openai-triton, 311 + # so we include it for the cuda build as well 312 + ++ lib.optionals (rocmSupport || cudaSupport) [ 313 + openai-triton 314 + ]; 294 315 295 316 # Tests take a long time and may be flaky, so just sanity-check imports 296 317 doCheck = false; ··· 316 337 (optionalString (majorMinor version == "1.3" ) "tensorboard") 317 338 ]) 318 339 "runHook postCheck" 340 + ]; 341 + 342 + pythonRemoveDeps = [ 343 + # In our dist-info the name is just "triton" 344 + "pytorch-triton-rocm" 319 345 ]; 320 346 321 347 postInstall = ''
+8 -3
pkgs/development/python-modules/torchinfo/default.nix
··· 9 9 10 10 buildPythonPackage rec { 11 11 pname = "torchinfo"; 12 - version = "1.64"; 12 + version = "1.7.2"; 13 13 format = "setuptools"; 14 14 15 15 disabled = pythonOlder "3.7"; ··· 18 18 owner = "TylerYep"; 19 19 repo = pname; 20 20 rev = "refs/tags/v${version}"; 21 - hash = "sha256-gcl8RxCD017FP4LtB60WVtOh7jg2Otv/vNd9hKneEAU="; 21 + hash = "sha256-O+I7BNQ5moV/ZcbbuP/IFoi0LO0WsGHBbSfgPmFu1Ec="; 22 22 }; 23 23 24 24 propagatedBuildInputs = [ ··· 37 37 "test_google" 38 38 ]; 39 39 40 + disabledTestPaths = [ 41 + # Wants "compressai", which we don't package (2023-03-23) 42 + "tests/torchinfo_xl_test.py" 43 + ]; 44 + 40 45 pythonImportsCheck = [ 41 - "torchvision" 46 + "torchinfo" 42 47 ]; 43 48 44 49 meta = with lib; {
+2
pkgs/tools/audio/tts/default.nix
··· 188 188 "tests/vocoder_tests/test_multiband_melgan_train.py" 189 189 "tests/vocoder_tests/test_melgan_train.py" 190 190 "tests/vocoder_tests/test_wavernn_train.py" 191 + # only a feed forward test, but still takes too long 192 + "tests/tts_tests/test_overflow.py" 191 193 ]; 192 194 193 195 passthru = {
+7
pkgs/tools/misc/partition-manager/default.nix
··· 62 62 63 63 meta = with lib; { 64 64 description = "KDE Partition Manager"; 65 + longDescription = '' 66 + KDE Partition Manager is a utility to help you manage the disks, partitions, and file systems on your computer. 67 + It allows you to easily create, copy, move, delete, back up, restore, and resize them without losing data. 68 + It supports a large number of file systems, including ext2/3/4, btrfs, reiserfs, NTFS, FAT16/32, JFS, XFS and more. 69 + 70 + To install on NixOS, use the option `programs.partition-manager.enable = true`. 71 + ''; 65 72 license = with licenses; [ cc-by-40 cc0 gpl3Plus lgpl3Plus mit ]; 66 73 homepage = "https://www.kde.org/applications/system/kdepartitionmanager/"; 67 74 maintainers = with maintainers; [ peterhoeg oxalica ];
+60 -20
pkgs/tools/networking/ntopng/default.nix
··· 1 - { lib, stdenv, fetchFromGitHub, fetchpatch, pkg-config, autoreconfHook 2 - , zeromq, ndpi, json_c, openssl, libpcap, libcap, curl, libmaxminddb 3 - , rrdtool, sqlite, libmysqlclient, expat, net-snmp 1 + { lib 2 + , stdenv 3 + , autoreconfHook 4 + , curl 5 + , expat 6 + , fetchFromGitHub 7 + , git 8 + , json_c 9 + , libcap 10 + , libmaxminddb 11 + , libmysqlclient 12 + , libpcap 13 + , libsodium 14 + , ndpi 15 + , net-snmp 16 + , openssl 17 + , pkg-config 18 + , rdkafka 19 + , gtest 20 + , rrdtool 21 + , hiredis 22 + , sqlite 23 + , which 24 + , zeromq 4 25 }: 5 26 6 27 stdenv.mkDerivation rec { 7 28 pname = "ntopng"; 8 - version = "5.2.1"; 29 + version = "5.6"; 9 30 10 31 src = fetchFromGitHub { 11 32 owner = "ntop"; 12 33 repo = "ntopng"; 13 - rev = version; 14 - sha256 = "sha256-FeRERSq8F3HEelUCkA6pgNNcP94xrWy6EbJgk+cEdqc="; 34 + rev = "refs/tags/${version}"; 35 + hash = "sha256-iGqrS0AneKYwGMEpbKy9if8bnaEu6aEV+QaH+JrF9xs="; 15 36 }; 16 37 17 - patches = [ 18 - (fetchpatch { 19 - url = "https://github.com/ntop/ntopng/commit/0aa580e1a45f248fffe6d11729ce40571f08e187.patch"; 20 - sha256 = "sha256-xqEVwfGgkNS+akbJnLZsVvEQdp9GxxUen8VkFomtcPI="; 21 - }) 22 - ]; 38 + preConfigure = '' 39 + substituteInPlace Makefile.in \ 40 + --replace "/bin/rm" "rm" 41 + ''; 23 42 24 - nativeBuildInputs = [ autoreconfHook pkg-config ]; 43 + nativeBuildInputs = [ 44 + autoreconfHook 45 + git 46 + pkg-config 47 + which 48 + ]; 25 49 26 50 buildInputs = [ 27 - zeromq ndpi json_c openssl libpcap curl libmaxminddb rrdtool sqlite 28 - libmysqlclient expat net-snmp libcap 51 + curl 52 + expat 53 + json_c 54 + libcap 55 + libmaxminddb 56 + libmysqlclient 57 + libpcap 58 + gtest 59 + hiredis 60 + libsodium 61 + net-snmp 62 + openssl 63 + rdkafka 64 + rrdtool 65 + sqlite 66 + zeromq 29 67 ]; 30 68 31 69 autoreconfPhase = "bash autogen.sh"; 32 70 33 - preConfigure = '' 34 - substituteInPlace Makefile.in --replace "/bin/rm" "rm" 35 - ''; 71 + configureFlags = [ 72 + "--with-ndpi-includes=${ndpi}/include/ndpi" 73 + "--with-ndpi-static-lib=${ndpi}/lib/" 74 + ]; 36 75 37 76 preBuild = '' 38 77 sed -e "s|\(#define CONST_BIN_DIR \).*|\1\"$out/bin\"|g" \ ··· 44 83 45 84 meta = with lib; { 46 85 description = "High-speed web-based traffic analysis and flow collection tool"; 47 - homepage = "http://www.ntop.org/products/ntop/"; 86 + homepage = "https://www.ntop.org/products/traffic-analysis/ntop/"; 87 + changelog = "https://github.com/ntop/ntopng/blob/${version}/CHANGELOG.md"; 48 88 license = licenses.gpl3Plus; 49 89 platforms = platforms.linux ++ platforms.darwin; 50 - maintainers = [ maintainers.bjornfor ]; 90 + maintainers = with maintainers; [ bjornfor ]; 51 91 }; 52 92 }
+2 -2
pkgs/tools/text/goawk/default.nix
··· 2 2 3 3 buildGoModule rec { 4 4 pname = "goawk"; 5 - version = "1.21.0"; 5 + version = "1.22.0"; 6 6 7 7 src = fetchFromGitHub { 8 8 owner = "benhoyt"; 9 9 repo = "goawk"; 10 10 rev = "v${version}"; 11 - sha256 = "sha256-I6KmNPFD8kkYDyek8lR1ZS7biPA/LYGwJqMoA2fG7Wg="; 11 + sha256 = "sha256-8UoYGAmYmC0hcxLfkNac3flKyPBT/xsykuy+TvVosBI="; 12 12 }; 13 13 14 14 vendorSha256 = null;
+4
pkgs/top-level/python-packages.nix
··· 6809 6809 6810 6810 open-meteo = callPackage ../development/python-modules/open-meteo { }; 6811 6811 6812 + openai-triton = callPackage ../development/python-modules/openai-triton { llvmPackages = pkgs.llvmPackages_rocm; }; 6813 + 6812 6814 openai-whisper = callPackage ../development/python-modules/openai-whisper { }; 6813 6815 6814 6816 openant = callPackage ../development/python-modules/openant { }; ··· 11864 11866 torchWithCuda = self.torch.override { 11865 11867 magma = pkgs.magma-cuda; 11866 11868 cudaSupport = true; 11869 + rocmSupport = false; 11867 11870 }; 11868 11871 11869 11872 torchWithoutCuda = self.torch.override { ··· 11873 11876 torchWithRocm = self.torch.override { 11874 11877 magma = pkgs.magma-hip; 11875 11878 rocmSupport = true; 11879 + cudaSupport = false; 11876 11880 }; 11877 11881 11878 11882 torchWithoutRocm = self.torch.override {