From 9435abcaee6e68ba746deef96bde5846e0abe46e Mon Sep 17 00:00:00 2001 From: Mike Purvis Date: Tue, 7 Mar 2023 11:01:24 -0500 Subject: [PATCH 01/32] Expose the buildTensorRTPackage function. --- .../development/libraries/science/math/tensorrt/extension.nix | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/pkgs/development/libraries/science/math/tensorrt/extension.nix b/pkgs/development/libraries/science/math/tensorrt/extension.nix index b4f47a8969c..507dc0cfd50 100644 --- a/pkgs/development/libraries/science/math/tensorrt/extension.nix +++ b/pkgs/development/libraries/science/math/tensorrt/extension.nix @@ -27,7 +27,9 @@ final: prev: let defaultBuild = { "tensorrt" = if allBuilds ? ${computeName tensorRTDefaultVersion} then allBuilds.${computeName tensorRTDefaultVersion} else throw "tensorrt-${tensorRTDefaultVersion} does not support your cuda version ${cudaVersion}"; }; - in allBuilds // defaultBuild; + in { + inherit buildTensorRTPackage; + } // allBuilds // defaultBuild; tensorRTVersions = { "8.4.0" = [ From 0f57b4bd736cbb3237e2b221d4fcbb742d24b527 Mon Sep 17 00:00:00 2001 From: Felix Buehler Date: Sun, 5 Mar 2023 17:42:40 +0100 Subject: [PATCH 02/32] fetchMavenArtifact: deprecate phases & use pname+version --- .../fetchmavenartifact/default.nix | 17 ++++++----------- 1 file changed, 6 insertions(+), 11 deletions(-) diff --git a/pkgs/build-support/fetchmavenartifact/default.nix b/pkgs/build-support/fetchmavenartifact/default.nix index b1d8593e72d..0f3cd4e64dd 100644 --- a/pkgs/build-support/fetchmavenartifact/default.nix +++ b/pkgs/build-support/fetchmavenartifact/default.nix @@ -37,13 +37,8 @@ assert (url == "") || (urls == []); assert (repos != []) || (url != "") || (urls != []); let - name_ = - lib.concatStrings [ - (lib.replaceStrings ["."] ["_"] groupId) "_" - (lib.replaceStrings ["."] ["_"] artifactId) "-" - version - ]; - suffix = if isNull classifier then "" else "-${classifier}"; + pname = (lib.replaceStrings [ "." ] [ "_" ] groupId) + "_" + (lib.replaceStrings [ "." ] [ "_" ] artifactId); + suffix = lib.optionalString (classifier != null) "-${classifier}"; filename = "${artifactId}-${version}${suffix}.jar"; mkJarUrl = repoUrl: lib.concatStringsSep "/" [ @@ -59,13 +54,13 @@ let else map mkJarUrl repos; jar = fetchurl ( - builtins.removeAttrs args ["groupId" "artifactId" "version" "classifier" "repos" "url" ] - // { urls = urls_; name = "${name_}.jar"; } + builtins.removeAttrs args [ "groupId" "artifactId" "version" "classifier" "repos" "url" ] + // { urls = urls_; name = "${pname}-${version}.jar"; } ); in stdenv.mkDerivation { - name = name_; - phases = "installPhase fixupPhase"; + inherit pname version; + dontUnpack = true; # By moving the jar to $out/share/java we make it discoverable by java # packages packages that mention this derivation in their buildInputs. installPhase = '' From 26773687e155a68ecafb8a333a660f7e339347ce Mon Sep 17 00:00:00 2001 From: Fabian Affolter Date: Thu, 30 Mar 2023 16:46:01 +0200 Subject: [PATCH 03/32] ndpi: add changelog to meta --- pkgs/development/libraries/ndpi/default.nix | 27 ++++++++++++++------- 1 file changed, 18 insertions(+), 9 deletions(-) diff --git a/pkgs/development/libraries/ndpi/default.nix b/pkgs/development/libraries/ndpi/default.nix index b798320c691..a3b247d08ca 100644 --- a/pkgs/development/libraries/ndpi/default.nix +++ b/pkgs/development/libraries/ndpi/default.nix @@ -1,13 +1,14 @@ { lib , stdenv -, fetchFromGitHub -, which , autoconf , automake -, libtool -, libpcap +, fetchFromGitHub , json_c -, pkg-config }: +, libpcap +, libtool +, pkg-config +, which +}: stdenv.mkDerivation rec { pname = "ndpi"; @@ -16,16 +17,23 @@ stdenv.mkDerivation rec { src = fetchFromGitHub { owner = "ntop"; repo = "nDPI"; - rev = version; - sha256 = "sha256-ZWWuyPGl+hbrfXdtPvCBqMReuJ4FiGx+qiI7qCz6wtQ="; + rev = "refs/tags/${version}"; + hash = "sha256-ZWWuyPGl+hbrfXdtPvCBqMReuJ4FiGx+qiI7qCz6wtQ="; }; configureScript = "./autogen.sh"; - nativeBuildInputs = [ which autoconf automake libtool pkg-config ]; + nativeBuildInputs = [ + autoconf + automake + libtool + pkg-config + which + ]; + buildInputs = [ - libpcap json_c + libpcap ]; meta = with lib; { @@ -34,6 +42,7 @@ stdenv.mkDerivation rec { nDPI is a library for deep-packet inspection based on OpenDPI. ''; homepage = "https://www.ntop.org/products/deep-packet-inspection/ndpi/"; + changelog = "https://github.com/ntop/nDPI/blob/${version}/CHANGELOG.md"; license = with licenses; [ lgpl3Plus bsd3 ]; maintainers = with maintainers; [ takikawa ]; mainProgram = "ndpiReader"; From a9c29a7af72e8f7ea1f193a52d60431a9a1aed09 Mon Sep 17 00:00:00 2001 From: Fabian Affolter Date: Thu, 30 Mar 2023 16:47:26 +0200 Subject: [PATCH 04/32] ndpi: 4.2 -> 4.6 Diff: https://github.com/ntop/nDPI/compare/refs/tags/4.2...4.6 Changelog: https://github.com/ntop/nDPI/blob/4.6/CHANGELOG.md --- pkgs/development/libraries/ndpi/default.nix | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/pkgs/development/libraries/ndpi/default.nix b/pkgs/development/libraries/ndpi/default.nix index a3b247d08ca..d9d8f5895a2 100644 --- a/pkgs/development/libraries/ndpi/default.nix +++ b/pkgs/development/libraries/ndpi/default.nix @@ -12,13 +12,13 @@ stdenv.mkDerivation rec { pname = "ndpi"; - version = "4.2"; + version = "4.6"; src = fetchFromGitHub { owner = "ntop"; repo = "nDPI"; rev = "refs/tags/${version}"; - hash = "sha256-ZWWuyPGl+hbrfXdtPvCBqMReuJ4FiGx+qiI7qCz6wtQ="; + hash = "sha256-S0lVh5FZewPbYG/1ikI2RroCSC7OI8Xmfeq73hYCHnY="; }; configureScript = "./autogen.sh"; From 0d88b0dac68abc0e9ebc4fb36a6fecbe818d3c5f Mon Sep 17 00:00:00 2001 From: Fabian Affolter Date: Thu, 30 Mar 2023 22:25:36 +0200 Subject: [PATCH 05/32] ntopng: add changelog to meta --- pkgs/tools/networking/ntopng/default.nix | 1 + 1 file changed, 1 insertion(+) diff --git a/pkgs/tools/networking/ntopng/default.nix b/pkgs/tools/networking/ntopng/default.nix index d5700c42534..73944732aa4 100644 --- a/pkgs/tools/networking/ntopng/default.nix +++ b/pkgs/tools/networking/ntopng/default.nix @@ -45,6 +45,7 @@ stdenv.mkDerivation rec { meta = with lib; { description = "High-speed web-based traffic analysis and flow collection tool"; homepage = "http://www.ntop.org/products/ntop/"; + changelog = "https://github.com/ntop/ntopng/blob/${version}/CHANGELOG.md"; license = licenses.gpl3Plus; platforms = platforms.linux ++ platforms.darwin; maintainers = [ maintainers.bjornfor ]; From 31cb9b08409ea3d4c0fc0eaaaa2af600bc8eb9fa Mon Sep 17 00:00:00 2001 From: Fabian Affolter Date: Thu, 30 Mar 2023 22:29:35 +0200 Subject: [PATCH 06/32] ntopng: 5.2.1 -> 5.6 Changelog: https://github.com/ntop/ntopng/blob/5.6/CHANGELOG.md --- pkgs/tools/networking/ntopng/default.nix | 79 ++++++++++++++++++------ 1 file changed, 59 insertions(+), 20 deletions(-) diff --git a/pkgs/tools/networking/ntopng/default.nix b/pkgs/tools/networking/ntopng/default.nix index 73944732aa4..c1feb3126d7 100644 --- a/pkgs/tools/networking/ntopng/default.nix +++ b/pkgs/tools/networking/ntopng/default.nix @@ -1,38 +1,77 @@ -{ lib, stdenv, fetchFromGitHub, fetchpatch, pkg-config, autoreconfHook -, zeromq, ndpi, json_c, openssl, libpcap, libcap, curl, libmaxminddb -, rrdtool, sqlite, libmysqlclient, expat, net-snmp +{ lib +, stdenv +, autoreconfHook +, curl +, expat +, fetchFromGitHub +, git +, json_c +, libcap +, libmaxminddb +, libmysqlclient +, libpcap +, libsodium +, ndpi +, net-snmp +, openssl +, pkg-config +, rdkafka +, gtest +, rrdtool +, hiredis +, sqlite +, which +, zeromq }: stdenv.mkDerivation rec { pname = "ntopng"; - version = "5.2.1"; + version = "5.6"; src = fetchFromGitHub { owner = "ntop"; repo = "ntopng"; - rev = version; - sha256 = "sha256-FeRERSq8F3HEelUCkA6pgNNcP94xrWy6EbJgk+cEdqc="; + rev = "refs/tags/${version}"; + hash = "sha256-iGqrS0AneKYwGMEpbKy9if8bnaEu6aEV+QaH+JrF9xs="; }; - patches = [ - (fetchpatch { - url = "https://github.com/ntop/ntopng/commit/0aa580e1a45f248fffe6d11729ce40571f08e187.patch"; - sha256 = "sha256-xqEVwfGgkNS+akbJnLZsVvEQdp9GxxUen8VkFomtcPI="; - }) + preConfigure = '' + substituteInPlace Makefile.in \ + --replace "/bin/rm" "rm" + ''; + + nativeBuildInputs = [ + autoreconfHook + git + pkg-config + which ]; - nativeBuildInputs = [ autoreconfHook pkg-config ]; - buildInputs = [ - zeromq ndpi json_c openssl libpcap curl libmaxminddb rrdtool sqlite - libmysqlclient expat net-snmp libcap + curl + expat + json_c + libcap + libmaxminddb + libmysqlclient + libpcap + gtest + hiredis + libsodium + net-snmp + openssl + rdkafka + rrdtool + sqlite + zeromq ]; autoreconfPhase = "bash autogen.sh"; - preConfigure = '' - substituteInPlace Makefile.in --replace "/bin/rm" "rm" - ''; + configureFlags = [ + "--with-ndpi-includes=${ndpi}/include/ndpi" + "--with-ndpi-static-lib=${ndpi}/lib/" + ]; preBuild = '' sed -e "s|\(#define CONST_BIN_DIR \).*|\1\"$out/bin\"|g" \ @@ -44,10 +83,10 @@ stdenv.mkDerivation rec { meta = with lib; { description = "High-speed web-based traffic analysis and flow collection tool"; - homepage = "http://www.ntop.org/products/ntop/"; + homepage = "https://www.ntop.org/products/traffic-analysis/ntop/"; changelog = "https://github.com/ntop/ntopng/blob/${version}/CHANGELOG.md"; license = licenses.gpl3Plus; platforms = platforms.linux ++ platforms.darwin; - maintainers = [ maintainers.bjornfor ]; + maintainers = with maintainers; [ bjornfor ]; }; } From bcd2d49d8512a8530def8fe999d6183e0fff3f09 Mon Sep 17 00:00:00 2001 From: Robert Hensing Date: Fri, 7 Apr 2023 01:09:05 +0200 Subject: [PATCH 07/32] nixos: Make services.resolved discoverable via "systemd-resolved" search This query yielded no results on search.nixos.org. I don't think I can make all options magically appear, but you can the other options by reading the text. --- nixos/modules/system/boot/resolved.nix | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/nixos/modules/system/boot/resolved.nix b/nixos/modules/system/boot/resolved.nix index 0ab2a875975..4e7201833db 100644 --- a/nixos/modules/system/boot/resolved.nix +++ b/nixos/modules/system/boot/resolved.nix @@ -16,7 +16,9 @@ in default = false; type = types.bool; description = lib.mdDoc '' - Whether to enable the systemd DNS resolver daemon. + Whether to enable the systemd DNS resolver daemon, `systemd-resolved`. + + Search for `services.resolved` to see all options. ''; }; From 68b4014b6508c4b1555eae7191509fea1218e7e6 Mon Sep 17 00:00:00 2001 From: "R. Ryantm" Date: Thu, 6 Apr 2023 23:40:18 +0000 Subject: [PATCH 08/32] tauon: 7.6.2 -> 7.6.3 --- pkgs/applications/audio/tauon/default.nix | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/pkgs/applications/audio/tauon/default.nix b/pkgs/applications/audio/tauon/default.nix index 6441182e72a..427c1e9e5dc 100644 --- a/pkgs/applications/audio/tauon/default.nix +++ b/pkgs/applications/audio/tauon/default.nix @@ -25,13 +25,13 @@ stdenv.mkDerivation rec { pname = "tauon"; - version = "7.6.2"; + version = "7.6.3"; src = fetchFromGitHub { owner = "Taiko2k"; repo = "TauonMusicBox"; rev = "v${version}"; - hash = "sha256-x/tHCDplC45XEaBaf0aQ0w8AS1SorXtYilJoiOcBDtM="; + hash = "sha256-cNR4Ffn9HvgL5KV4FUSnbzEh6VfoKaIbfpb18/qKEns="; }; postUnpack = '' From 837b784d15c6d415cf84955b2fc4aaba55d56352 Mon Sep 17 00:00:00 2001 From: "R. Ryantm" Date: Fri, 7 Apr 2023 00:55:33 +0000 Subject: [PATCH 09/32] binocle: 0.3.0 -> 0.3.1 --- pkgs/applications/misc/binocle/default.nix | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/pkgs/applications/misc/binocle/default.nix b/pkgs/applications/misc/binocle/default.nix index 6a999790635..8ec1490bdcc 100644 --- a/pkgs/applications/misc/binocle/default.nix +++ b/pkgs/applications/misc/binocle/default.nix @@ -16,16 +16,16 @@ rustPlatform.buildRustPackage rec { pname = "binocle"; - version = "0.3.0"; + version = "0.3.1"; src = fetchFromGitHub { owner = "sharkdp"; repo = pname; rev = "v${version}"; - sha256 = "0b0hf2aq34kxxj0la0yar5sp44k6mqcbyailp6j6q0mksf1l74bc"; + sha256 = "sha256-L4l8Gl7Ok/TWqHjNujPx8Qk3UWebs0SgOQNyBNtpnZo="; }; - cargoSha256 = "sha256-CZWAHWZYaL54Rl6Jrp8B6w6HK+2fIKQle2x4mGHv2/o="; + cargoHash = "sha256-9d0MNQ7jEJKpGbjVtl1XBoOBEVNKDgFouSMrcZ7tXNU="; nativeBuildInputs = [ makeWrapper From 136e3dbdae827ca91aa5559f9edbce5191ded602 Mon Sep 17 00:00:00 2001 From: Aroun Date: Tue, 21 Mar 2023 23:49:27 +0530 Subject: [PATCH 10/32] docker: add pointer to related nix option Co-authored-by: Sandro --- pkgs/applications/virtualization/docker/default.nix | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/pkgs/applications/virtualization/docker/default.nix b/pkgs/applications/virtualization/docker/default.nix index 17b3982a1f7..0d846e9a696 100644 --- a/pkgs/applications/virtualization/docker/default.nix +++ b/pkgs/applications/virtualization/docker/default.nix @@ -254,6 +254,11 @@ rec { meta = with lib; { homepage = "https://www.docker.com/"; description = "An open source project to pack, ship and run any application as a lightweight container"; + longDescription = '' + Docker is a platform designed to help developers build, share, and run modern applications. + + To enable the docker daemon on NixOS, set the `virtualisation.docker.enable` option to `true`. + ''; license = licenses.asl20; maintainers = with maintainers; [ offline tailhook vdemeester periklis mikroskeem maxeaubrey ]; }; From 00afa18aff64de113ade8c70bb4e85def296b387 Mon Sep 17 00:00:00 2001 From: Aroun Date: Tue, 21 Mar 2023 23:07:52 +0530 Subject: [PATCH 11/32] partition-manager: add indication for the nix option Co-authored-by: Sandro --- pkgs/tools/misc/partition-manager/default.nix | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/pkgs/tools/misc/partition-manager/default.nix b/pkgs/tools/misc/partition-manager/default.nix index 0d566eb11b2..4587830e6cd 100644 --- a/pkgs/tools/misc/partition-manager/default.nix +++ b/pkgs/tools/misc/partition-manager/default.nix @@ -62,6 +62,13 @@ in mkDerivation rec { meta = with lib; { description = "KDE Partition Manager"; + longDescription = '' + KDE Partition Manager is a utility to help you manage the disks, partitions, and file systems on your computer. + It allows you to easily create, copy, move, delete, back up, restore, and resize them without losing data. + It supports a large number of file systems, including ext2/3/4, btrfs, reiserfs, NTFS, FAT16/32, JFS, XFS and more. + + To install on NixOS, use the option `programs.partition-manager.enable = true`. + ''; license = with licenses; [ cc-by-40 cc0 gpl3Plus lgpl3Plus mit ]; homepage = "https://www.kde.org/applications/system/kdepartitionmanager/"; maintainers = with maintainers; [ peterhoeg oxalica ]; From 09d5d6b198a6af9bdab0ee5c211a459070f71b88 Mon Sep 17 00:00:00 2001 From: Someone Serge Date: Tue, 21 Mar 2023 03:29:07 +0200 Subject: [PATCH 12/32] python3Packages.torch: 1.13.1 -> 2.0.0 --- .../python-modules/torch/default.nix | 32 +++++++++---------- 1 file changed, 16 insertions(+), 16 deletions(-) diff --git a/pkgs/development/python-modules/torch/default.nix b/pkgs/development/python-modules/torch/default.nix index 062fcea4334..3373c3d3d3b 100644 --- a/pkgs/development/python-modules/torch/default.nix +++ b/pkgs/development/python-modules/torch/default.nix @@ -12,6 +12,7 @@ Accelerate, CoreServices, libobjc, # Propagated build inputs + sympy, numpy, pyyaml, cffi, click, typing-extensions, # Unit tests @@ -49,9 +50,7 @@ let inherit (cudaPackages) cudatoolkit cudaFlags cudnn nccl; in -# assert that everything needed for cuda is present and that the correct cuda versions are used -assert !cudaSupport || (let majorIs = lib.versions.major cudatoolkit.version; - in majorIs == "9" || majorIs == "10" || majorIs == "11"); +assert cudaSupport -> (cudaPackages.cudaMajorVersion == "11"); # confirm that cudatoolkits are sync'd across dependencies assert !(MPISupport && cudaSupport) || mpi.cudatoolkit == cudatoolkit; @@ -129,10 +128,10 @@ let in buildPythonPackage rec { pname = "torch"; # Don't forget to update torch-bin to the same version. - version = "1.13.1"; + version = "2.0.0"; format = "setuptools"; - disabled = pythonOlder "3.7.0"; + disabled = pythonOlder "3.8.0"; outputs = [ "out" # output standard python package @@ -145,7 +144,7 @@ in buildPythonPackage rec { repo = "pytorch"; rev = "refs/tags/v${version}"; fetchSubmodules = true; - hash = "sha256-yQz+xHPw9ODRBkV9hv1th38ZmUr/fXa+K+d+cvmX3Z8="; + hash = "sha256-cSw7+AYBUcZLz3UyK/+JWWjQxKwVBXcFvBq0XAcL3tE="; }; patches = lib.optionals (stdenv.isDarwin && stdenv.isx86_64) [ @@ -155,15 +154,6 @@ in buildPythonPackage rec { # base is 10.12. Until we upgrade, we can fall back on the older # pthread support. ./pthreadpool-disable-gcd.diff - ] ++ [ - # PyTorch fails to build on gcc 12 due to gloo - # https://github.com/pytorch/pytorch/issues/77614 - (fetchpatch { - url = "https://github.com/facebookincubator/gloo/commit/4a5e339b764261d20fc409071dc7a8b8989aa195.patch"; - stripLen = 1; - extraPrefix = "third_party/gloo/"; - hash = "sha256-UxR1r7F6g76BWj3GBIrSy5t+YZDCWy6mMddwx+hon5w="; - }) ]; postPatch = lib.optionalString rocmSupport '' @@ -261,7 +251,16 @@ in buildPythonPackage rec { # Suppress gcc regression: avx512 math function raises uninitialized variable warning # https://gcc.gnu.org/bugzilla/show_bug.cgi?id=105593 # See also: Fails to compile with GCC 12.1.0 https://github.com/pytorch/pytorch/issues/77939 - ++ lib.optionals stdenv.cc.isGNU [ "-Wno-error=maybe-uninitialized" "-Wno-error=uninitialized" ])); + ++ lib.optionals stdenv.cc.isGNU [ + "-Wno-error=maybe-uninitialized" + "-Wno-error=uninitialized" + ] + # Since pytorch 2.0: + # gcc-12.2.0/include/c++/12.2.0/bits/new_allocator.h:158:33: error: ‘void operator delete(void*, std::size_t)’ + # ... called on pointer ‘’ with nonzero offset [1, 9223372036854775800] [-Werror=free-nonheap-object] + ++ lib.optionals (stdenv.cc.isGNU && lib.versions.major stdenv.cc.version == "12" ) [ + "-Wno-error=free-nonheap-object" + ])); nativeBuildInputs = [ cmake @@ -287,6 +286,7 @@ in buildPythonPackage rec { numpy pyyaml typing-extensions + sympy # the following are required for tensorboard support pillow six future tensorboard protobuf ] ++ lib.optionals MPISupport [ mpi ] From a9faf1b9efee58458e0174e696471a7db3ff97ed Mon Sep 17 00:00:00 2001 From: Someone Serge Date: Tue, 21 Mar 2023 03:34:20 +0200 Subject: [PATCH 13/32] python3Packages.torch: add missing install_requires --- pkgs/development/python-modules/torch/default.nix | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/pkgs/development/python-modules/torch/default.nix b/pkgs/development/python-modules/torch/default.nix index 3373c3d3d3b..9f701493f07 100644 --- a/pkgs/development/python-modules/torch/default.nix +++ b/pkgs/development/python-modules/torch/default.nix @@ -12,7 +12,10 @@ Accelerate, CoreServices, libobjc, # Propagated build inputs + filelock, sympy, + networkx, + jinja2, numpy, pyyaml, cffi, click, typing-extensions, # Unit tests @@ -285,8 +288,14 @@ in buildPythonPackage rec { click numpy pyyaml + + # From install_requires: + filelock typing-extensions sympy + networkx + jinja2 + # the following are required for tensorboard support pillow six future tensorboard protobuf ] ++ lib.optionals MPISupport [ mpi ] From 0f76efb48143e4f38f7714588e3cb38d055404b2 Mon Sep 17 00:00:00 2001 From: Someone Serge Date: Tue, 21 Mar 2023 16:30:06 +0200 Subject: [PATCH 14/32] python3Packages.torchWithRocm: ignore config.cudaSupport --- pkgs/top-level/python-packages.nix | 2 ++ 1 file changed, 2 insertions(+) diff --git a/pkgs/top-level/python-packages.nix b/pkgs/top-level/python-packages.nix index 1e9075ace6a..5349c922a4e 100644 --- a/pkgs/top-level/python-packages.nix +++ b/pkgs/top-level/python-packages.nix @@ -11856,6 +11856,7 @@ self: super: with self; { torchWithCuda = self.torch.override { magma = pkgs.magma-cuda; cudaSupport = true; + rocmSupport = false; }; torchWithoutCuda = self.torch.override { @@ -11865,6 +11866,7 @@ self: super: with self; { torchWithRocm = self.torch.override { magma = pkgs.magma-hip; rocmSupport = true; + cudaSupport = false; }; torchWithoutRocm = self.torch.override { From 378c0c69832c350f672a21128b5d2782a29c2d6e Mon Sep 17 00:00:00 2001 From: Someone Serge Date: Wed, 22 Mar 2023 05:34:47 +0200 Subject: [PATCH 15/32] python3Packages.openai-triton: init at 2.0.0 --- pkgs/development/compilers/llvm/rocm/llvm.nix | 6 +- .../python-modules/openai-triton/default.nix | 246 + .../python-modules/openai-triton/llvm15.patch | 4617 +++++++++++++++++ .../python-modules/torch/default.nix | 25 +- pkgs/tools/audio/tts/default.nix | 2 + pkgs/top-level/python-packages.nix | 2 + 6 files changed, 4893 insertions(+), 5 deletions(-) create mode 100644 pkgs/development/python-modules/openai-triton/default.nix create mode 100644 pkgs/development/python-modules/openai-triton/llvm15.patch diff --git a/pkgs/development/compilers/llvm/rocm/llvm.nix b/pkgs/development/compilers/llvm/rocm/llvm.nix index 1f1add5cf67..6092bc1a9fc 100644 --- a/pkgs/development/compilers/llvm/rocm/llvm.nix +++ b/pkgs/development/compilers/llvm/rocm/llvm.nix @@ -24,6 +24,8 @@ , targetDir ? "llvm" , targetProjects ? [ ] , targetRuntimes ? [ ] +# "NATIVE" resolves into x86 or aarch64 depending on stdenv +, llvmTargetsToBuild ? [ "NATIVE" ] , extraPatches ? [ ] , extraNativeBuildInputs ? [ ] , extraBuildInputs ? [ ] @@ -46,6 +48,8 @@ let if stdenv.isx86_64 then "X86" else if stdenv.isAarch64 then "AArch64" else throw "Unsupported ROCm LLVM platform"; + inferNativeTarget = t: if t == "NATIVE" then llvmNativeTarget else t; + llvmTargetsToBuild' = [ "AMDGPU" ] ++ builtins.map inferNativeTarget llvmTargetsToBuild; in stdenv.mkDerivation (finalAttrs: { pname = "rocm-llvm-${targetName}"; version = "5.4.4"; @@ -98,7 +102,7 @@ in stdenv.mkDerivation (finalAttrs: { sourceRoot = "${finalAttrs.src.name}/${targetDir}"; cmakeFlags = [ - "-DLLVM_TARGETS_TO_BUILD=AMDGPU;${llvmNativeTarget}" + "-DLLVM_TARGETS_TO_BUILD=${builtins.concatStringsSep ";" llvmTargetsToBuild'}" ] ++ lib.optionals (finalAttrs.passthru.isLLVM && targetProjects != [ ]) [ "-DLLVM_ENABLE_PROJECTS=${lib.concatStringsSep ";" targetProjects}" ] ++ lib.optionals ((finalAttrs.passthru.isLLVM || targetDir == "runtimes") && targetRuntimes != [ ]) [ diff --git a/pkgs/development/python-modules/openai-triton/default.nix b/pkgs/development/python-modules/openai-triton/default.nix new file mode 100644 index 00000000000..9340aad3a95 --- /dev/null +++ b/pkgs/development/python-modules/openai-triton/default.nix @@ -0,0 +1,246 @@ +{ lib +, buildPythonPackage +, python +, fetchpatch +, fetchFromGitHub +, addOpenGLRunpath +, cmake +, cudaPackages +, llvmPackages +, pybind11 +, gtest +, zlib +, ncurses +, libxml2 +, lit +, filelock +, torchWithRocm +, pytest +, pytestCheckHook +, pythonRelaxDepsHook +, pkgsTargetTarget +}: + +let + pname = "triton"; + version = "2.0.0"; + + inherit (cudaPackages) cuda_cudart backendStdenv; + ptxas = "${pkgsTargetTarget.cudaPackages.cuda_nvcc}/bin/ptxas"; + + llvm = (llvmPackages.llvm.override { + llvmTargetsToBuild = [ "NATIVE" "NVPTX" ]; + # Upstream CI sets these too: + # targetProjects = [ "mlir" ]; + extraCMakeFlags = [ + "-DLLVM_INSTALL_UTILS=ON" + ]; + }); +in +buildPythonPackage { + inherit pname version; + + format = "setuptools"; + + src = fetchFromGitHub { + owner = "openai"; + repo = pname; + rev = "v${version}"; + hash = "sha256-9GZzugab+Pdt74Dj6zjlEzjj4BcJ69rzMJmqcVMxsKU="; + }; + + patches = [ + # Prerequisite for llvm15 patch + (fetchpatch { + url = "https://github.com/openai/triton/commit/2aba985daaa70234823ea8f1161da938477d3e02.patch"; + hash = "sha256-LGv0+Ut2WYPC4Ksi4803Hwmhi3FyQOF9zElJc/JCobk="; + }) + (fetchpatch { + url = "https://github.com/openai/triton/commit/e3941f9d09cdd31529ba4a41018cfc0096aafea6.patch"; + hash = "sha256-A+Gor6qzFlGQhVVhiaaYOzqqx8yO2MdssnQS6TIfUWg="; + }) + + # Source: https://github.com/openai/triton/commit/fc7a8e35819bda632bdcf1cf75fd9abe4d4e077a.patch + # The original patch adds ptxas binary, so we include our own clean copy + # Drop with the next update + ./llvm15.patch + + # TODO: there have been commits upstream aimed at removing the "torch" + # circular dependency, but the patches fail to apply on the release + # revision. Keeping the link for future reference + # Also cf. https://github.com/openai/triton/issues/1374 + + # (fetchpatch { + # url = "https://github.com/openai/triton/commit/fc7c0b0e437a191e421faa61494b2ff4870850f1.patch"; + # hash = "sha256-f0shIqHJkVvuil2Yku7vuqWFn7VCRKFSFjYRlwx25ig="; + # }) + ]; + + postPatch = '' + substituteInPlace python/setup.py \ + --replace \ + '= get_thirdparty_packages(triton_cache_path)' \ + '= os.environ["cmakeFlags"].split()' + '' + # Wiring triton=2.0.0 with llcmPackages_rocm.llvm=5.4.3 + # Revisit when updating either triton or llvm + + '' + substituteInPlace CMakeLists.txt \ + --replace "nvptx" "NVPTX" \ + --replace "LLVM 11" "LLVM" + sed -i '/AddMLIR/a set(MLIR_TABLEGEN_EXE "${llvmPackages.mlir}/bin/mlir-tblgen")' CMakeLists.txt + sed -i '/AddMLIR/a set(MLIR_INCLUDE_DIR ''${MLIR_INCLUDE_DIRS})' CMakeLists.txt + find -iname '*.td' -exec \ + sed -i \ + -e '\|include "mlir/IR/OpBase.td"|a include "mlir/IR/AttrTypeBase.td"' \ + -e 's|include "mlir/Dialect/StandardOps/IR/Ops.td"|include "mlir/Dialect/Func/IR/FuncOps.td"|' \ + '{}' ';' + substituteInPlace unittest/CMakeLists.txt --replace "include(GoogleTest)" "find_package(GTest REQUIRED)" + sed -i 's/^include.*$//' unittest/CMakeLists.txt + sed -i '/LINK_LIBS/i NVPTXInfo' lib/Target/PTX/CMakeLists.txt + sed -i '/LINK_LIBS/i NVPTXCodeGen' lib/Target/PTX/CMakeLists.txt + '' + # TritonMLIRIR already links MLIRIR. Not transitive? + # + '' + # echo "target_link_libraries(TritonPTX PUBLIC MLIRIR)" >> lib/Target/PTX/CMakeLists.txt + # '' + # Already defined in llvm, when built with -DLLVM_INSTALL_UTILS + + '' + substituteInPlace bin/CMakeLists.txt \ + --replace "add_subdirectory(FileCheck)" "" + + rm cmake/FindLLVM.cmake + '' + + + ( + let + # Bash was getting weird without linting, + # but basically upstream contains [cc, ..., "-lcuda", ...] + # and we replace it with [..., "-lcuda", "-L/run/opengl-driver/lib", "-L$stubs", ...] + old = [ "-lcuda" ]; + new = [ "-lcuda" "-L${addOpenGLRunpath.driverLink}" "-L${cuda_cudart}/lib/stubs/" ]; + + quote = x: ''"${x}"''; + oldStr = lib.concatMapStringsSep ", " quote old; + newStr = lib.concatMapStringsSep ", " quote new; + in + '' + substituteInPlace python/triton/compiler.py \ + --replace '${oldStr}' '${newStr}' + '' + ) + # Triton seems to be looking up cuda.h + + '' + sed -i 's|cu_include_dir = os.path.join.*$|cu_include_dir = "${cuda_cudart}/include"|' python/triton/compiler.py + ''; + + nativeBuildInputs = [ + cmake + pythonRelaxDepsHook + + # Requires torch (circular dependency) and probably needs GPUs: + # pytestCheckHook + + # Note for future: + # These *probably* should go in depsTargetTarget + # ...but we cannot test cross right now anyway + # because we only support cudaPackages on x86_64-linux atm + lit + llvm + llvmPackages.mlir + ]; + + buildInputs = [ + gtest + libxml2.dev + ncurses + pybind11 + zlib + ]; + + propagatedBuildInputs = [ + filelock + ]; + + # Avoid GLIBCXX mismatch with other cuda-enabled python packages + preConfigure = + '' + export CC="${backendStdenv.cc}/bin/cc"; + export CXX="${backendStdenv.cc}/bin/c++"; + '' + # Upstream's setup.py tries to write cache somewhere in ~/ + + '' + export HOME=$TMPDIR + '' + # Upstream's github actions patch setup.cfg to write base-dir. May be redundant + + '' + echo "" >> python/setup.cfg + echo "[build_ext]" >> python/setup.cfg + echo "base-dir=$PWD" >> python/setup.cfg + '' + # The rest (including buildPhase) is relative to ./python/ + + '' + cd python/ + '' + # Work around download_and_copy_ptxas() + + '' + dst_cuda="$PWD/triton/third_party/cuda/bin" + mkdir -p "$dst_cuda" + ln -s "${ptxas}" "$dst_cuda/" + ''; + + # CMake is run by setup.py instead + dontUseCmakeConfigure = true; + cmakeFlags = [ + "-DMLIR_DIR=${llvmPackages.mlir}/lib/cmake/mlir" + ]; + + postFixup = + let + ptxasDestination = "$out/${python.sitePackages}/triton/third_party/cuda/bin/ptxas"; + in + # Setuptools (?) strips runpath and +x flags. Let's just restore the symlink + '' + rm -f ${ptxasDestination} + ln -s ${ptxas} ${ptxasDestination} + ''; + + checkInputs = [ + cmake # ctest + ]; + dontUseSetuptoolsCheck = true; + preCheck = + # build/temp* refers to build_ext.build_temp (looked up in the build logs) + '' + (cd /build/source/python/build/temp* ; ctest) + '' # For pytestCheckHook + + '' + cd test/unit + ''; + pythonImportsCheck = [ + # Circular dependency on torch + # "triton" + # "triton.language" + ]; + + # Ultimately, torch is our test suite: + passthru.tests = { + inherit torchWithRocm; + }; + + pythonRemoveDeps = [ + # Circular dependency, cf. https://github.com/openai/triton/issues/1374 + "torch" + + # CLI tools without dist-info + "cmake" + "lit" + ]; + meta = with lib; { + description = "Development repository for the Triton language and compiler"; + homepage = "https://github.com/openai/triton/"; + platforms = lib.platforms.unix; + license = licenses.mit; + maintainers = with maintainers; [ SomeoneSerge ]; + }; +} diff --git a/pkgs/development/python-modules/openai-triton/llvm15.patch b/pkgs/development/python-modules/openai-triton/llvm15.patch new file mode 100644 index 00000000000..3e20cce2380 --- /dev/null +++ b/pkgs/development/python-modules/openai-triton/llvm15.patch @@ -0,0 +1,4617 @@ +From fc7a8e35819bda632bdcf1cf75fd9abe4d4e077a Mon Sep 17 00:00:00 2001 +From: Christian Sigg +Date: Thu, 16 Feb 2023 15:40:53 +0100 +Subject: [PATCH] Rebase Triton to LLVM-15. (#1070) + +This PR rebases Triton from LLVM-14 to LLVM-15. Most changes are +mechanical, except for the analysis framework changes. +--- + CMakeLists.txt | 6 +- + bin/CMakeLists.txt | 2 +- + bin/FileCheck/FileCheck.cpp | 3 + + bin/triton-opt.cpp | 6 +- + bin/triton-translate.cpp | 7 +- + include/triton/Analysis/Alias.h | 21 +- + include/triton/Analysis/Allocation.h | 2 + + include/triton/Analysis/AxisInfo.h | 56 ++- + include/triton/Analysis/Utility.h | 6 +- + include/triton/Conversion/Passes.td | 4 +- + include/triton/Dialect/Triton/IR/Dialect.h | 7 +- + .../triton/Dialect/Triton/IR/TritonDialect.td | 8 +- + include/triton/Dialect/Triton/IR/TritonOps.td | 12 +- + .../triton/Dialect/Triton/IR/TritonTypes.td | 2 + + .../Dialect/Triton/Transforms/Passes.td | 3 +- + include/triton/Dialect/TritonGPU/IR/Dialect.h | 4 +- + .../Dialect/TritonGPU/IR/TritonGPUAttrDefs.td | 7 + + .../Dialect/TritonGPU/IR/TritonGPUDialect.td | 2 +- + .../Dialect/TritonGPU/IR/TritonGPUOps.td | 13 +- + lib/Analysis/Alias.cpp | 14 +- + lib/Analysis/Allocation.cpp | 30 +- + lib/Analysis/AxisInfo.cpp | 79 ++-- + lib/Analysis/CMakeLists.txt | 2 +- + lib/Analysis/Membar.cpp | 2 +- + lib/Analysis/Utility.cpp | 54 +++ + .../TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp | 3 - + lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h | 10 +- + .../TritonGPUToLLVM/DotOpToLLVM.cpp | 5 - + .../TritonGPUToLLVM/ElementwiseOpToLLVM.cpp | 2 - + .../TritonGPUToLLVM/LoadStoreOpToLLVM.cpp | 5 +- + .../TritonGPUToLLVM/ReduceOpToLLVM.cpp | 2 - + .../TritonGPUToLLVM/TritonGPUToLLVM.cpp | 7 +- + .../TritonGPUToLLVM/TritonGPUToLLVMBase.h | 26 +- + .../TritonGPUToLLVM/TritonGPUToLLVMPass.cpp | 52 +-- + lib/Conversion/TritonGPUToLLVM/Utility.h | 5 +- + .../TritonToTritonGPUPass.cpp | 69 ++-- + lib/Dialect/Triton/IR/CMakeLists.txt | 10 +- + lib/Dialect/Triton/IR/Ops.cpp | 34 +- + lib/Dialect/Triton/Transforms/Combine.cpp | 6 +- + lib/Dialect/Triton/Transforms/Combine.td | 2 +- + lib/Dialect/TritonGPU/IR/Dialect.cpp | 27 +- + lib/Dialect/TritonGPU/Transforms/Coalesce.cpp | 20 +- + lib/Dialect/TritonGPU/Transforms/Combine.cpp | 2 +- + lib/Dialect/TritonGPU/Transforms/Combine.td | 1 + + .../Transforms/DecomposeConversions.cpp | 2 +- + lib/Dialect/TritonGPU/Transforms/Pipeline.cpp | 10 +- + .../Transforms/ReorderInstructions.cpp | 2 +- + .../Transforms/TritonGPUConversion.cpp | 12 +- + .../Transforms/UpdateMmaForVolta.cpp | 6 +- + lib/Dialect/TritonGPU/Transforms/Utility.cpp | 2 +- + lib/Target/LLVMIR/CMakeLists.txt | 3 +- + lib/Target/PTX/PTXTranslation.cpp | 3 + + python/setup.py | 15 +- + python/src/triton.cc | 85 +++-- + python/test/unit/language/test_core.py | 2 +- + python/triton/compiler.py | 4 +- + test/Analysis/test-alias.mlir | 24 +- + test/Analysis/test-alignment.mlir | 344 +++++++++--------- + test/Analysis/test-allocation.mlir | 32 +- + test/Analysis/test-membar.mlir | 38 +- + test/Conversion/triton_ops.mlir | 10 +- + test/Conversion/triton_to_tritongpu.mlir | 6 +- + test/Conversion/tritongpu_to_llvm.mlir | 94 ++--- + test/Target/tritongpu_to_llvmir.mlir | 4 +- + test/Target/tritongpu_to_ptx.mlir | 2 +- + test/Triton/combine.mlir | 40 +- + test/Triton/vecadd.mlir | 4 +- + test/TritonGPU/coalesce.mlir | 2 +- + test/TritonGPU/combine.mlir | 38 +- + test/TritonGPU/loop-pipeline.mlir | 22 +- + test/TritonGPU/matmul.mlir | 4 +- + test/TritonGPU/prefetch.mlir | 4 +- + test/TritonGPU/update-mma-for-volta.mlir | 4 +- + test/lib/Analysis/TestAlias.cpp | 29 +- + test/lib/Analysis/TestAllocation.cpp | 5 +- + test/lib/Analysis/TestAxisInfo.cpp | 51 +-- + test/lib/Analysis/TestMembar.cpp | 7 +- + 78 files changed, 808 insertions(+), 742 deletions(-) + +diff --git a/CMakeLists.txt b/CMakeLists.txt +index d0d361fc7c..b281a28400 100644 +--- a/CMakeLists.txt ++++ b/CMakeLists.txt +@@ -1,4 +1,7 @@ + cmake_minimum_required(VERSION 3.6) ++ ++cmake_policy(SET CMP0116 OLD) ++ + include(ExternalProject) + + set(CMAKE_CXX_STANDARD 17) +@@ -155,7 +158,6 @@ if(TRITON_BUILD_PYTHON_MODULE) + endif() + endif() + +- + # # Triton + # file(GLOB_RECURSE LIBTRITON_SRC lib/*.cc) + # if (WIN32 AND TRITON_BUILD_PYTHON_MODULE) +@@ -212,7 +214,7 @@ if(TRITON_BUILD_PYTHON_MODULE) + # optimizations + MLIRPass + MLIRTransforms +- MLIRLLVMIR ++ MLIRLLVMDialect + MLIRSupport + MLIRTargetLLVMIRExport + MLIRExecutionEngine +diff --git a/bin/CMakeLists.txt b/bin/CMakeLists.txt +index 906f635f8b..695b3479fd 100644 +--- a/bin/CMakeLists.txt ++++ b/bin/CMakeLists.txt +@@ -48,7 +48,7 @@ llvm_update_compile_flags(triton-translate) + # MLIR core + MLIROptLib + MLIRIR +- MLIRLLVMIR ++ MLIRLLVMDialect + MLIRPass + MLIRSupport + MLIRTransforms +diff --git a/bin/FileCheck/FileCheck.cpp b/bin/FileCheck/FileCheck.cpp +index 819efc3541..9ac6f1b277 100644 +--- a/bin/FileCheck/FileCheck.cpp ++++ b/bin/FileCheck/FileCheck.cpp +@@ -19,6 +19,7 @@ + #include "llvm/Support/CommandLine.h" + #include "llvm/Support/InitLLVM.h" + #include "llvm/Support/Process.h" ++#include "llvm/Support/SourceMgr.h" + #include "llvm/Support/WithColor.h" + #include "llvm/Support/raw_ostream.h" + #include +@@ -360,6 +361,8 @@ static std::string GetCheckTypeAbbreviation(Check::FileCheckType Ty) { + return "bad-not"; + case Check::CheckBadCount: + return "bad-count"; ++ case Check::CheckMisspelled: ++ return "misspelled"; + case Check::CheckNone: + llvm_unreachable("invalid FileCheckType"); + } +diff --git a/bin/triton-opt.cpp b/bin/triton-opt.cpp +index 9f3b53b7ae..f96232e1b0 100644 +--- a/bin/triton-opt.cpp ++++ b/bin/triton-opt.cpp +@@ -8,7 +8,7 @@ + + #include "mlir/IR/Dialect.h" + #include "mlir/InitAllPasses.h" +-#include "mlir/Support/MlirOptMain.h" ++#include "mlir/Tools/mlir-opt/MlirOptMain.h" + + namespace mlir { + namespace test { +@@ -33,8 +33,8 @@ int main(int argc, char **argv) { + // TODO: register Triton & TritonGPU passes + mlir::DialectRegistry registry; + registry.insert(); + + return mlir::asMainReturnCode(mlir::MlirOptMain( +diff --git a/bin/triton-translate.cpp b/bin/triton-translate.cpp +index 05ba15e453..56b5d65857 100644 +--- a/bin/triton-translate.cpp ++++ b/bin/triton-translate.cpp +@@ -3,7 +3,7 @@ + #include "mlir/IR/AsmState.h" + #include "mlir/IR/BuiltinOps.h" + #include "mlir/IR/Dialect.h" +-#include "mlir/Parser.h" ++#include "mlir/Parser/Parser.h" + #include "mlir/Pass/Pass.h" + #include "mlir/Pass/PassManager.h" + #include "mlir/Support/FileUtilities.h" +@@ -38,7 +38,7 @@ OwningOpRef loadMLIRModule(llvm::StringRef inputFilename, + mlir::DialectRegistry registry; + registry.insert(); ++ scf::SCFDialect>(); + + context.appendDialectRegistry(registry); + +@@ -50,7 +50,8 @@ OwningOpRef loadMLIRModule(llvm::StringRef inputFilename, + context.loadAllAvailableDialects(); + context.allowUnregisteredDialects(); + +- OwningOpRef module(parseSourceFile(sourceMgr, &context)); ++ OwningOpRef module = ++ parseSourceFile(sourceMgr, &context); + if (!module) { + llvm::errs() << "Parse MLIR file failed."; + return nullptr; +diff --git a/include/triton/Analysis/Alias.h b/include/triton/Analysis/Alias.h +index fa6b906fc9..631df518bc 100644 +--- a/include/triton/Analysis/Alias.h ++++ b/include/triton/Analysis/Alias.h +@@ -2,7 +2,7 @@ + #define TRITON_ANALYSIS_ALIAS_H + + #include "mlir/Analysis/AliasAnalysis.h" +-#include "mlir/Analysis/DataFlowAnalysis.h" ++#include "mlir/Analysis/DataFlow/SparseAnalysis.h" + #include "llvm/ADT/DenseSet.h" + + namespace mlir { +@@ -21,7 +21,7 @@ class AliasInfo { + } + + /// The pessimistic value state of a value without alias +- static AliasInfo getPessimisticValueState(MLIRContext *context) { ++ static AliasInfo getPessimisticValueState(MLIRContext *context = nullptr) { + return AliasInfo(); + } + static AliasInfo getPessimisticValueState(Value value) { return AliasInfo(); } +@@ -29,6 +29,10 @@ class AliasInfo { + /// The union of both arguments + static AliasInfo join(const AliasInfo &lhs, const AliasInfo &rhs); + ++ void print(raw_ostream &os) const { ++ llvm::interleaveComma(allocs, os, [&](Value alloc) { alloc.print(os); }); ++ } ++ + private: + /// The set of allocated values that are aliased by this lattice. + /// For now, we only consider aliased value produced by the following +@@ -58,9 +62,13 @@ class AliasInfo { + //===----------------------------------------------------------------------===// + // Shared Memory Alias Analysis + //===----------------------------------------------------------------------===// +-class SharedMemoryAliasAnalysis : public ForwardDataFlowAnalysis { ++class SharedMemoryAliasAnalysis ++ : public dataflow::SparseDataFlowAnalysis> { + public: +- using ForwardDataFlowAnalysis::ForwardDataFlowAnalysis; ++ using dataflow::SparseDataFlowAnalysis< ++ dataflow::Lattice>::SparseDataFlowAnalysis; ++ using dataflow::SparseDataFlowAnalysis< ++ dataflow::Lattice>::getLatticeElement; + + /// XXX(Keren): Compatible interface with MLIR AliasAnalysis for future use. + /// Given two values, returns their aliasing behavior. +@@ -70,9 +78,10 @@ class SharedMemoryAliasAnalysis : public ForwardDataFlowAnalysis { + ModRefResult getModRef(Operation *op, Value location); + + /// Computes if the alloc set of the results are changed. +- ChangeResult ++ void + visitOperation(Operation *op, +- ArrayRef *> operands) override; ++ ArrayRef *> operands, ++ ArrayRef *> results) override; + }; + + } // namespace mlir +diff --git a/include/triton/Analysis/Allocation.h b/include/triton/Analysis/Allocation.h +index b7c136d602..89b77034cc 100644 +--- a/include/triton/Analysis/Allocation.h ++++ b/include/triton/Analysis/Allocation.h +@@ -188,6 +188,8 @@ class Allocation { + friend class triton::AllocationAnalysis; + }; + ++template Interval(T, T) -> Interval; ++ + } // namespace mlir + + #endif // TRITON_ANALYSIS_ALLOCATION_H +diff --git a/include/triton/Analysis/AxisInfo.h b/include/triton/Analysis/AxisInfo.h +index fdfbd8fbb3..7083b9c43b 100644 +--- a/include/triton/Analysis/AxisInfo.h ++++ b/include/triton/Analysis/AxisInfo.h +@@ -1,9 +1,10 @@ + #ifndef TRITON_ANALYSIS_AXISINFO_H + #define TRITON_ANALYSIS_AXISINFO_H + +-#include "mlir/Analysis/DataFlowAnalysis.h" ++#include "mlir/Analysis/DataFlow/SparseAnalysis.h" + #include "llvm/Support/raw_ostream.h" + ++#include "mlir/Support/LLVM.h" + #include "triton/Analysis/Utility.h" + #include "triton/Dialect/Triton/IR/Dialect.h" + #include "triton/Dialect/TritonGPU/IR/Dialect.h" +@@ -62,7 +63,7 @@ class AxisInfo { + } + + /// The pessimistic value state of the contiguity is unknown. +- static AxisInfo getPessimisticValueState(MLIRContext *context) { ++ static AxisInfo getPessimisticValueState(MLIRContext *context = nullptr) { + return AxisInfo(); + } + static AxisInfo getPessimisticValueState(Value value); +@@ -70,6 +71,22 @@ class AxisInfo { + /// The gcd of both arguments for each dimension + static AxisInfo join(const AxisInfo &lhs, const AxisInfo &rhs); + ++ void print(raw_ostream &os) const { ++ auto print = [&](StringRef name, DimVectorT vec) { ++ os << name << " = ["; ++ llvm::interleaveComma(vec, os); ++ os << "]"; ++ }; ++ print("contiguity", contiguity); ++ print(", divisibility", divisibility); ++ print(", constancy", constancy); ++ os << ", constant_value = "; ++ if (constantValue) ++ os << *constantValue; ++ else ++ os << ""; ++ } ++ + private: + /// The _contiguity_ information maps the `d`-th + /// dimension to the length of the shortest +@@ -147,7 +164,8 @@ class AxisInfoVisitor { + } + + virtual AxisInfo +- getAxisInfo(Operation *op, ArrayRef *> operands) = 0; ++ getAxisInfo(Operation *op, ++ ArrayRef *> operands) = 0; + + virtual bool match(Operation *op) = 0; + }; +@@ -157,15 +175,16 @@ template class AxisInfoVisitorImpl : public AxisInfoVisitor { + public: + using AxisInfoVisitor::AxisInfoVisitor; + +- AxisInfo getAxisInfo(Operation *op, +- ArrayRef *> operands) final { ++ AxisInfo ++ getAxisInfo(Operation *op, ++ ArrayRef *> operands) final { + return getAxisInfo(cast(op), operands); + } + + bool match(Operation *op) final { return isa(op); } + +- virtual AxisInfo getAxisInfo(OpTy op, +- ArrayRef *> operands) { ++ virtual AxisInfo ++ getAxisInfo(OpTy op, ArrayRef *> operands) { + llvm_unreachable("Unimplemented getAxisInfo"); + } + }; +@@ -176,8 +195,9 @@ class BinaryOpVisitorImpl : public AxisInfoVisitorImpl { + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(OpTy op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(OpTy op, ++ ArrayRef *> operands) override { + auto lhsInfo = operands[0]->getValue(); + auto rhsInfo = operands[1]->getValue(); + auto rank = lhsInfo.getRank(); +@@ -230,7 +250,8 @@ class AxisInfoVisitorList { + (visitors.emplace_back(std::make_unique()), ...); + } + +- AxisInfo apply(Operation *op, ArrayRef *> operands) { ++ AxisInfo apply(Operation *op, ++ ArrayRef *> operands) { + for (auto &visitor : visitors) + if (visitor->match(op)) + return visitor->getAxisInfo(op, operands); +@@ -241,16 +262,19 @@ class AxisInfoVisitorList { + std::vector> visitors; + }; + +-class AxisInfoAnalysis : public ForwardDataFlowAnalysis { ++class AxisInfoAnalysis ++ : public dataflow::SparseDataFlowAnalysis> { + private: + AxisInfoVisitorList visitors; + + public: +- AxisInfoAnalysis(MLIRContext *context); ++ AxisInfoAnalysis(DataFlowSolver &solver); ++ using dataflow::SparseDataFlowAnalysis< ++ dataflow::Lattice>::getLatticeElement; + +- ChangeResult +- visitOperation(Operation *op, +- ArrayRef *> operands) override; ++ void visitOperation(Operation *op, ++ ArrayRef *> operands, ++ ArrayRef *> results) override; + + unsigned getPtrContiguity(Value ptr); + +@@ -261,4 +285,4 @@ class AxisInfoAnalysis : public ForwardDataFlowAnalysis { + + } // namespace mlir + +-#endif +\ No newline at end of file ++#endif +diff --git a/include/triton/Analysis/Utility.h b/include/triton/Analysis/Utility.h +index c5ac137dc1..ee7fadb59d 100644 +--- a/include/triton/Analysis/Utility.h ++++ b/include/triton/Analysis/Utility.h +@@ -1,6 +1,7 @@ + #ifndef TRITON_ANALYSIS_UTILITY_H + #define TRITON_ANALYSIS_UTILITY_H + ++#include "mlir/Analysis/DataFlowFramework.h" + #include "mlir/Analysis/SliceAnalysis.h" + #include "triton/Dialect/TritonGPU/IR/Dialect.h" + #include +@@ -12,7 +13,7 @@ namespace mlir { + class ReduceOpHelper { + public: + explicit ReduceOpHelper(triton::ReduceOp op) : op(op) { +- srcTy = op.operand().getType().cast(); ++ srcTy = op.getOperand().getType().cast(); + } + + ArrayRef getSrcShape() { return srcTy.getShape(); } +@@ -103,6 +104,9 @@ SetVector + multiRootGetSlice(Operation *op, TransitiveFilter backwardFilter = nullptr, + TransitiveFilter forwardFilter = nullptr); + ++// Create a basic DataFlowSolver with constant and dead code analysis included. ++std::unique_ptr createDataFlowSolver(); ++ + } // namespace mlir + + #endif // TRITON_ANALYSIS_UTILITY_H +diff --git a/include/triton/Conversion/Passes.td b/include/triton/Conversion/Passes.td +index 70bb20b78e..be00eb2dac 100644 +--- a/include/triton/Conversion/Passes.td ++++ b/include/triton/Conversion/Passes.td +@@ -12,7 +12,6 @@ def ConvertTritonToTritonGPU: Pass<"convert-triton-to-tritongpu", "mlir::ModuleO + + let dependentDialects = ["mlir::arith::ArithmeticDialect", + "mlir::math::MathDialect", +- "mlir::StandardOpsDialect", + // TODO: Does this pass depend on SCF? + "mlir::scf::SCFDialect", + "mlir::triton::TritonDialect", +@@ -41,8 +40,7 @@ def ConvertTritonGPUToLLVM : Pass<"convert-triton-gpu-to-llvm", "mlir::ModuleOp" + "mlir::tensor::TensorDialect", + "mlir::triton::TritonDialect", + "mlir::triton::gpu::TritonGPUDialect", +- "mlir::NVVM::NVVMDialect", +- "mlir::StandardOpsDialect"]; ++ "mlir::NVVM::NVVMDialect"]; + + let options = [ + Option<"computeCapability", "compute-capability", +diff --git a/include/triton/Dialect/Triton/IR/Dialect.h b/include/triton/Dialect/Triton/IR/Dialect.h +index e8012a51df..15869e262e 100644 +--- a/include/triton/Dialect/Triton/IR/Dialect.h ++++ b/include/triton/Dialect/Triton/IR/Dialect.h +@@ -1,14 +1,15 @@ + #ifndef TRITON_DIALECT_TRITON_IR_DIALECT_H_ + #define TRITON_DIALECT_TRITON_IR_DIALECT_H_ + ++#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" ++#include "mlir/Dialect/ControlFlow/IR/ControlFlow.h" ++#include "mlir/Dialect/Func/IR/FuncOps.h" + #include "mlir/Dialect/Math/IR/Math.h" +-#include "mlir/Dialect/SCF/SCF.h" +-#include "mlir/Dialect/StandardOps/IR/Ops.h" ++#include "mlir/Dialect/SCF/IR/SCF.h" + #include "mlir/Dialect/Tensor/IR/Tensor.h" + #include "mlir/IR/BuiltinOps.h" + #include "mlir/IR/Dialect.h" + #include "mlir/Interfaces/ControlFlowInterfaces.h" +- + #include "triton/Dialect/Triton/IR/Dialect.h.inc" + #include "triton/Dialect/Triton/IR/OpsEnums.h.inc" + #include "triton/Dialect/Triton/IR/Traits.h" +diff --git a/include/triton/Dialect/Triton/IR/TritonDialect.td b/include/triton/Dialect/Triton/IR/TritonDialect.td +index 07b069e14f..d98ce73884 100644 +--- a/include/triton/Dialect/Triton/IR/TritonDialect.td ++++ b/include/triton/Dialect/Triton/IR/TritonDialect.td +@@ -25,12 +25,9 @@ def Triton_Dialect : Dialect { + let dependentDialects = [ + "arith::ArithmeticDialect", + "math::MathDialect", +- "StandardOpsDialect", + "scf::SCFDialect", +- +- // Since LLVM 15 +- // "cf::ControlFlowDialect", +- // "func::FuncDialect" ++ "cf::ControlFlowDialect", ++ "func::FuncDialect" + ]; + + let extraClassDeclaration = [{ +@@ -38,6 +35,7 @@ def Triton_Dialect : Dialect { + }]; + + let hasConstantMaterializer = 1; ++ let useDefaultTypePrinterParser = 1; + } + + include "triton/Dialect/Triton/IR/TritonTypes.td" +diff --git a/include/triton/Dialect/Triton/IR/TritonOps.td b/include/triton/Dialect/Triton/IR/TritonOps.td +index 779e0b648c..0a69211179 100644 +--- a/include/triton/Dialect/Triton/IR/TritonOps.td ++++ b/include/triton/Dialect/Triton/IR/TritonOps.td +@@ -141,11 +141,7 @@ def TT_LoadOp : TT_Op<"load", + "triton::EvictionPolicy":$evict, "bool":$isVolatile)>, + ]; + +- // let assemblyFormat = "operands attr-dict `:` type($result)"; +- let parser = [{ return mlir::triton::parseLoadOp(parser, result); }]; +- +- let printer = [{ return mlir::triton::printLoadOp(p, *this); }]; +- ++ let hasCustomAssemblyFormat = 1; + let hasCanonicalizer = 1; + } + +@@ -170,11 +166,7 @@ def TT_StoreOp : TT_Op<"store", + "triton::EvictionPolicy":$evict)>, + ]; + +- // let assemblyFormat = "operands attr-dict `:` type($value)"; +- let parser = [{ return mlir::triton::parseStoreOp(parser, result); }]; +- +- let printer = [{ return mlir::triton::printStoreOp(p, *this); }]; +- ++ let hasCustomAssemblyFormat = 1; + let hasCanonicalizer = 1; + } + +diff --git a/include/triton/Dialect/Triton/IR/TritonTypes.td b/include/triton/Dialect/Triton/IR/TritonTypes.td +index 66d2a7b9a9..2fe2fd077d 100644 +--- a/include/triton/Dialect/Triton/IR/TritonTypes.td ++++ b/include/triton/Dialect/Triton/IR/TritonTypes.td +@@ -1,6 +1,7 @@ + #ifndef TRITON_TYPES + #define TRITON_TYPES + ++include "mlir/IR/AttrTypeBase.td" + include "triton/Dialect/Triton/IR/TritonDialect.td" + + // +@@ -58,6 +59,7 @@ def TT_Ptr : TritonTypeDef<"Pointer", "ptr"> { + }]> + ]; + ++ let hasCustomAssemblyFormat = 1; + let skipDefaultBuilders = 1; + } + def TT_PtrTensor : TensorOf<[TT_Ptr]>; +diff --git a/include/triton/Dialect/Triton/Transforms/Passes.td b/include/triton/Dialect/Triton/Transforms/Passes.td +index 8f77aed774..a25cdc5680 100644 +--- a/include/triton/Dialect/Triton/Transforms/Passes.td ++++ b/include/triton/Dialect/Triton/Transforms/Passes.td +@@ -16,8 +16,7 @@ def TritonCombineOps : Pass + + let constructor = "mlir::triton::createCombineOpsPass()"; + +- let dependentDialects = ["mlir::arith::ArithmeticDialect", +- /*SelectOp*/"mlir::StandardOpsDialect"]; ++ let dependentDialects = ["mlir::arith::ArithmeticDialect"]; + } + + #endif +diff --git a/include/triton/Dialect/TritonGPU/IR/Dialect.h b/include/triton/Dialect/TritonGPU/IR/Dialect.h +index b4c8daec7b..dfc5f53ab1 100644 +--- a/include/triton/Dialect/TritonGPU/IR/Dialect.h ++++ b/include/triton/Dialect/TritonGPU/IR/Dialect.h +@@ -1,19 +1,17 @@ + #ifndef TRITON_DIALECT_TRITONGPU_IR_DIALECT_H_ + #define TRITON_DIALECT_TRITONGPU_IR_DIALECT_H_ + +-#include "mlir/Dialect/GPU/GPUDialect.h" ++#include "mlir/Dialect/GPU/IR/GPUDialect.h" + #include "mlir/Dialect/Tensor/IR/Tensor.h" + #include "mlir/IR/BuiltinOps.h" + #include "mlir/IR/Dialect.h" + + // TritonGPU depends on Triton + #include "triton/Dialect/Triton/IR/Dialect.h" +- + #include "triton/Dialect/TritonGPU/IR/Dialect.h.inc" + #include "triton/Dialect/TritonGPU/IR/Traits.h" + + #define GET_ATTRDEF_CLASSES +-#include "triton/Dialect/Triton/IR/AttrInterfaces.h.inc" + #include "triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.h.inc" + + #define GET_OP_CLASSES +diff --git a/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td b/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td +index 0242c3cc17..af2aeb03a8 100644 +--- a/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td ++++ b/include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td +@@ -1,6 +1,7 @@ + #ifndef TRITONGPU_ATTRDEFS + #define TRITONGPU_ATTRDEFS + ++include "mlir/IR/AttrTypeBase.td" + include "triton/Dialect/TritonGPU/IR/TritonGPUDialect.td" + include "triton/Dialect/Triton/IR/TritonInterfaces.td" + +@@ -136,6 +137,7 @@ A_{3, 2} A_{3, 3} A_{3, 0} A_{3, 1} ... [phase 1] / + ]; + + let extraClassDeclaration = extraBaseClassDeclaration; ++ let hasCustomAssemblyFormat = 1; + } + + //===----------------------------------------------------------------------===// +@@ -273,6 +275,7 @@ for + // ArrayRefParameter<"unsigned">:$sizePerCTA + ); + ++ let hasCustomAssemblyFormat = 1; + } + + //===----------------------------------------------------------------------===// +@@ -422,6 +425,7 @@ For example, the matrix L corresponding to blockTileSize=[32,16] is: + static constexpr int numBitsToHoldMmaV1ID{5}; + }]; + ++ let hasCustomAssemblyFormat = 1; + } + + def SliceEncodingAttr : DistributedEncoding<"SliceEncoding"> { +@@ -456,6 +460,8 @@ def SliceEncodingAttr : DistributedEncoding<"SliceEncoding"> { + template + SmallVector paddedShape(ArrayRef shape) const; + }]; ++ ++ let hasCustomAssemblyFormat = 1; + } + + def DotOperandEncodingAttr : DistributedEncoding<"DotOperandEncoding"> { +@@ -492,6 +498,7 @@ section 9.7.13.4.1 for more details. + + ]; + ++ let hasCustomAssemblyFormat = 1; + let extraClassDeclaration = extraBaseClassDeclaration; + } + +diff --git a/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td b/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td +index 87ec1d36c6..6489a721b4 100644 +--- a/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td ++++ b/include/triton/Dialect/TritonGPU/IR/TritonGPUDialect.td +@@ -30,7 +30,7 @@ def TritonGPU_Dialect : Dialect { + } + }]; + +- ++ let useDefaultAttributePrinterParser = 1; + } + + #endif +diff --git a/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td b/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td +index 510f8d0183..7aba11dc75 100644 +--- a/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td ++++ b/include/triton/Dialect/TritonGPU/IR/TritonGPUOps.td +@@ -59,7 +59,7 @@ def TTG_AsyncCommitGroupOp : TTG_Op<"async_commit_group"> { + // This is needed because these ops don't + // handle encodings + // e.g., https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/Dialect/Arith/IR/ArithOps.td#L111 +-def TTG_CmpIOp : TTG_Op<"cmpi", [NoSideEffect, Elementwise, ++def TTG_CmpIOp : TTG_Op<"cmpi", [NoSideEffect, Elementwise, + SameOperandsAndResultShape, + SameOperandsAndResultEncoding]> { + let summary = "integer comparison operation"; +@@ -73,7 +73,7 @@ def TTG_CmpIOp : TTG_Op<"cmpi", [NoSideEffect, Elementwise, + let results = (outs TT_BoolLike:$result); + } + +-def TTG_CmpFOp : TTG_Op<"cmpf", [NoSideEffect, Elementwise, ++def TTG_CmpFOp : TTG_Op<"cmpf", [NoSideEffect, Elementwise, + SameOperandsAndResultShape, + SameOperandsAndResultEncoding]> { + let summary = "floating-point comparison operation"; +@@ -88,8 +88,8 @@ def TTG_CmpFOp : TTG_Op<"cmpf", [NoSideEffect, Elementwise, + } + + // TODO: migrate to arith::SelectOp on LLVM16 +-def TTG_SelectOp : TTG_Op<"select", [NoSideEffect, Elementwise, +- SameOperandsAndResultShape, ++def TTG_SelectOp : TTG_Op<"select", [NoSideEffect, Elementwise, ++ SameOperandsAndResultShape, + SameOperandsAndResultEncoding]> { + let summary = "select operation"; + +@@ -188,10 +188,7 @@ def TTG_InsertSliceAsyncOp : TTG_Op<"insert_slice_async", + } + }]; + +- // The custom parser could be replaced with oilist in LLVM-16 +- let parser = [{ return parseInsertSliceAsyncOp(parser, result); }]; +- +- let printer = [{ return printInsertSliceAsyncOp(p, *this); }]; ++ let hasCustomAssemblyFormat = 1; + } + + def TTG_AllocTensorOp : TTG_Op<"alloc_tensor", [MemoryEffects<[MemAlloc]>, // Allocate shared memory +diff --git a/lib/Analysis/Alias.cpp b/lib/Analysis/Alias.cpp +index a39e4de9aa..208fdd4afc 100644 +--- a/lib/Analysis/Alias.cpp ++++ b/lib/Analysis/Alias.cpp +@@ -18,8 +18,9 @@ AliasInfo AliasInfo::join(const AliasInfo &lhs, const AliasInfo &rhs) { + return ret; + } + +-ChangeResult SharedMemoryAliasAnalysis::visitOperation( +- Operation *op, ArrayRef *> operands) { ++void SharedMemoryAliasAnalysis::visitOperation( ++ Operation *op, ArrayRef *> operands, ++ ArrayRef *> results) { + AliasInfo aliasInfo; + bool pessimistic = true; + if (maybeSharedAllocationOp(op)) { +@@ -44,14 +45,11 @@ ChangeResult SharedMemoryAliasAnalysis::visitOperation( + } + + if (pessimistic) { +- return markAllPessimisticFixpoint(op->getResults()); ++ return markAllPessimisticFixpoint(results); + } + // Join all lattice elements +- ChangeResult result = ChangeResult::NoChange; +- for (Value value : op->getResults()) { +- result |= getLatticeElement(value).join(aliasInfo); +- } +- return result; ++ for (auto *result : results) ++ propagateIfChanged(result, result->join(aliasInfo)); + } + + AliasResult SharedMemoryAliasAnalysis::alias(Value lhs, Value rhs) { +diff --git a/lib/Analysis/Allocation.cpp b/lib/Analysis/Allocation.cpp +index 712c08c475..b4de8dcd9d 100644 +--- a/lib/Analysis/Allocation.cpp ++++ b/lib/Analysis/Allocation.cpp +@@ -1,4 +1,5 @@ + #include "triton/Analysis/Allocation.h" ++#include "mlir/Analysis/DataFlowFramework.h" + #include "mlir/Analysis/Liveness.h" + #include "mlir/Analysis/SliceAnalysis.h" + #include "mlir/Dialect/Tensor/IR/Tensor.h" +@@ -33,10 +34,8 @@ constexpr int kPtrBitWidth = 64; + + static std::pair, SmallVector> + getCvtOrder(const Attribute &srcLayout, const Attribute &dstLayout) { +- auto srcBlockedLayout = srcLayout.dyn_cast(); + auto srcMmaLayout = srcLayout.dyn_cast(); + auto srcDotLayout = srcLayout.dyn_cast(); +- auto dstBlockedLayout = dstLayout.dyn_cast(); + auto dstMmaLayout = dstLayout.dyn_cast(); + auto dstDotLayout = dstLayout.dyn_cast(); + assert(!(srcMmaLayout && dstMmaLayout) && +@@ -224,14 +223,12 @@ class AllocationAnalysis { + } + + void getValueAlias(Value value, SharedMemoryAliasAnalysis &analysis) { +- LatticeElement *latticeElement = +- analysis.lookupLatticeElement(value); +- if (latticeElement) { +- auto &info = latticeElement->getValue(); +- if (!info.getAllocs().empty()) { +- for (auto alloc : info.getAllocs()) { +- allocation->addAlias(value, alloc); +- } ++ dataflow::Lattice *latticeElement = ++ analysis.getLatticeElement(value); ++ if (latticeElement && !latticeElement->isUninitialized()) { ++ AliasInfo &info = latticeElement->getValue(); ++ for (auto alloc : info.getAllocs()) { ++ allocation->addAlias(value, alloc); + } + } + } +@@ -244,14 +241,19 @@ class AllocationAnalysis { + getScratchValueSize(op); + }); + // Get the alias values +- SharedMemoryAliasAnalysis aliasAnalysis(operation->getContext()); +- aliasAnalysis.run(operation); ++ std::unique_ptr solver = createDataFlowSolver(); ++ SharedMemoryAliasAnalysis *aliasAnalysis = ++ solver->load(); ++ if (failed(solver->initializeAndRun(operation))) { ++ // TODO: return error instead of bailing out.. ++ llvm_unreachable("failed to run SharedMemoryAliasAnalysis"); ++ } + operation->walk([&](Operation *op) { + for (auto operand : op->getOperands()) { +- getValueAlias(operand, aliasAnalysis); ++ getValueAlias(operand, *aliasAnalysis); + } + for (auto value : op->getResults()) { +- getValueAlias(value, aliasAnalysis); ++ getValueAlias(value, *aliasAnalysis); + } + }); + } +diff --git a/lib/Analysis/AxisInfo.cpp b/lib/Analysis/AxisInfo.cpp +index 0b7142b04d..4af46c3fbb 100644 +--- a/lib/Analysis/AxisInfo.cpp ++++ b/lib/Analysis/AxisInfo.cpp +@@ -1,4 +1,4 @@ +-#include "mlir/Analysis/DataFlowAnalysis.h" ++#include "mlir/Analysis/DataFlowFramework.h" + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" + #include "llvm/Support/raw_ostream.h" + +@@ -52,7 +52,7 @@ AxisInfo AxisInfo::getPessimisticValueState(Value value) { + BlockArgument blockArg = value.dyn_cast(); + if (blockArg && blockArg.getOwner()->isEntryBlock()) { + Operation *op = blockArg.getOwner()->getParentOp(); +- if (FuncOp fun = dyn_cast(op)) { ++ if (func::FuncOp fun = dyn_cast(op)) { + Attribute attr = + fun.getArgAttr(blockArg.getArgNumber(), "tt.divisibility"); + if (attr) +@@ -136,8 +136,9 @@ class CastOpAxisInfoVisitor final : public AxisInfoVisitorImpl { + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(OpTy op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(OpTy op, ++ ArrayRef *> operands) override { + return operands[0]->getValue(); + } + }; +@@ -147,8 +148,9 @@ class MakeRangeOpAxisInfoVisitor final + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(triton::MakeRangeOp op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(triton::MakeRangeOp op, ++ ArrayRef *> operands) override { + auto start = op.start(); + auto end = op.end(); + return AxisInfo(/*contiguity=*/{end - start}, +@@ -162,8 +164,9 @@ class ConstantOpAxisInfoVisitor final + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(arith::ConstantOp op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(arith::ConstantOp op, ++ ArrayRef *> operands) override { + auto intAttr = op.getValue().dyn_cast(); + auto boolAttr = op.getValue().dyn_cast(); + if (intAttr || boolAttr) { +@@ -416,8 +419,9 @@ class SplatOpAxisInfoVisitor final + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(triton::SplatOp op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(triton::SplatOp op, ++ ArrayRef *> operands) override { + Type _retTy = *op->result_type_begin(); + TensorType retTy = _retTy.cast(); + AxisInfo opInfo = operands[0]->getValue(); +@@ -439,8 +443,9 @@ class ExpandDimsOpAxisInfoVisitor final + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(triton::ExpandDimsOp op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(triton::ExpandDimsOp op, ++ ArrayRef *> operands) override { + AxisInfo opInfo = operands[0]->getValue(); + AxisInfo::DimVectorT contiguity = opInfo.getContiguity(); + AxisInfo::DimVectorT divisibility = opInfo.getDivisibility(); +@@ -458,8 +463,9 @@ class BroadcastOpAxisInfoVisitor final + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(triton::BroadcastOp op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(triton::BroadcastOp op, ++ ArrayRef *> operands) override { + Type _retTy = *op->result_type_begin(); + Type _opTy = *op->operand_type_begin(); + TensorType retTy = _retTy.cast(); +@@ -486,8 +492,9 @@ class CmpOpAxisInfoVisitor final : public AxisInfoVisitorImpl { + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(OpTy op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(OpTy op, ++ ArrayRef *> operands) override { + auto resTy = op.getResult().getType().template dyn_cast(); + if (!resTy) + return AxisInfo(); +@@ -596,8 +603,9 @@ class SelectOpAxisInfoVisitor final : public AxisInfoVisitorImpl { + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(OpTy op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(OpTy op, ++ ArrayRef *> operands) override { + auto resTy = op.getResult().getType().template dyn_cast(); + if (!resTy) + return AxisInfo(); +@@ -757,8 +765,9 @@ class MaxMinOpAxisInfoVisitor final : public AxisInfoVisitorImpl { + public: + using AxisInfoVisitorImpl::AxisInfoVisitorImpl; + +- AxisInfo getAxisInfo(OpTy op, +- ArrayRef *> operands) override { ++ AxisInfo ++ getAxisInfo(OpTy op, ++ ArrayRef *> operands) override { + auto lhsInfo = operands[0]->getValue(); + auto rhsInfo = operands[1]->getValue(); + std::optional constantValue; +@@ -786,8 +795,8 @@ class MaxMinOpAxisInfoVisitor final : public AxisInfoVisitorImpl { + // AxisInfoAnalysis + //===----------------------------------------------------------------------===// + +-AxisInfoAnalysis::AxisInfoAnalysis(MLIRContext *context) +- : ForwardDataFlowAnalysis(context) { ++AxisInfoAnalysis::AxisInfoAnalysis(DataFlowSolver &solver) ++ : dataflow::SparseDataFlowAnalysis>(solver) { + // UnrealizedConversionCast: + // This is needed by TritonGPUToLLVM, to get AxisInfo when the graph is + // in the process of a PartialConversion, where UnrealizedConversionCast +@@ -819,7 +828,7 @@ AxisInfoAnalysis::AxisInfoAnalysis(MLIRContext *context) + visitors.append, + LogicalOpAxisInfoVisitor, + LogicalOpAxisInfoVisitor>(); +- visitors.append, ++ visitors.append, + SelectOpAxisInfoVisitor>(); + visitors.append, + ShROpAxisInfoVisitor>(); +@@ -829,11 +838,12 @@ AxisInfoAnalysis::AxisInfoAnalysis(MLIRContext *context) + MaxMinOpAxisInfoVisitor>(); + } + +-ChangeResult AxisInfoAnalysis::visitOperation( +- Operation *op, ArrayRef *> operands) { ++void AxisInfoAnalysis::visitOperation( ++ Operation *op, ArrayRef *> operands, ++ ArrayRef *> results) { + AxisInfo curr = visitors.apply(op, operands); + if (curr.getRank() == 0) { +- return markAllPessimisticFixpoint(op->getResults()); ++ return markAllPessimisticFixpoint(results); + } + // override with hint + auto newContiguity = curr.getContiguity(); +@@ -854,11 +864,8 @@ ChangeResult AxisInfoAnalysis::visitOperation( + curr = mlir::AxisInfo(newContiguity, newDivisibility, newConstancy, + curr.getConstantValue()); + // join all lattice elements +- ChangeResult result = ChangeResult::NoChange; +- for (Value value : op->getResults()) { +- result |= getLatticeElement(value).join(curr); +- } +- return result; ++ for (auto *result : results) ++ propagateIfChanged(result, result->join(curr)); + } + + unsigned AxisInfoAnalysis::getPtrContiguity(Value ptr) { +@@ -884,7 +891,10 @@ unsigned AxisInfoAnalysis::getPtrAlignment(Value ptr) { + auto tensorTy = ptr.getType().dyn_cast(); + if (!tensorTy) + return 1; +- auto axisInfo = lookupLatticeElement(ptr)->getValue(); ++ dataflow::Lattice *latticeElement = getLatticeElement(ptr); ++ if (!latticeElement || latticeElement->isUninitialized()) ++ return 1; ++ auto axisInfo = latticeElement->getValue(); + auto layout = tensorTy.getEncoding(); + auto order = triton::gpu::getOrder(layout); + auto maxMultipleBytes = axisInfo.getDivisibility(order[0]); +@@ -900,8 +910,11 @@ unsigned AxisInfoAnalysis::getMaskAlignment(Value mask) { + auto tensorTy = mask.getType().dyn_cast(); + if (!tensorTy) + return 1; ++ dataflow::Lattice *latticeElement = getLatticeElement(mask); ++ if (!latticeElement || latticeElement->isUninitialized()) ++ return 1; ++ auto maskAxis = latticeElement->getValue(); + auto maskOrder = triton::gpu::getOrder(tensorTy.getEncoding()); +- auto maskAxis = lookupLatticeElement(mask)->getValue(); + auto alignment = std::max(maskAxis.getConstancy(maskOrder[0]), 1); + return alignment; + } +diff --git a/lib/Analysis/CMakeLists.txt b/lib/Analysis/CMakeLists.txt +index afbc692510..1f761f845c 100644 +--- a/lib/Analysis/CMakeLists.txt ++++ b/lib/Analysis/CMakeLists.txt +@@ -8,7 +8,7 @@ add_mlir_library(TritonAnalysis + DEPENDS + TritonTableGen + TritonGPUAttrDefsIncGen +- ++ + LINK_LIBS PUBLIC + MLIRAnalysis + ) +diff --git a/lib/Analysis/Membar.cpp b/lib/Analysis/Membar.cpp +index acc885e827..910274b2ac 100644 +--- a/lib/Analysis/Membar.cpp ++++ b/lib/Analysis/Membar.cpp +@@ -2,7 +2,7 @@ + #include "triton/Analysis/Alias.h" + #include "triton/Dialect/TritonGPU/IR/Dialect.h" + +-#include "mlir/Dialect/GPU/GPUDialect.h" ++#include "mlir/Dialect/GPU/IR/GPUDialect.h" + #include "mlir/Dialect/Tensor/IR/Tensor.h" + + namespace mlir { +diff --git a/lib/Analysis/Utility.cpp b/lib/Analysis/Utility.cpp +index d9e917e731..6ea52df272 100644 +--- a/lib/Analysis/Utility.cpp ++++ b/lib/Analysis/Utility.cpp +@@ -1,5 +1,8 @@ + #include "triton/Analysis/Utility.h" ++#include "mlir/Analysis/DataFlow/ConstantPropagationAnalysis.h" ++#include "mlir/Analysis/DataFlow/DeadCodeAnalysis.h" + #include "mlir/IR/Dialect.h" ++#include "mlir/IR/Matchers.h" + #include "triton/Dialect/Triton/IR/Dialect.h" + #include "triton/Dialect/TritonGPU/IR/Dialect.h" + #include +@@ -325,4 +328,55 @@ SetVector multiRootGetSlice(Operation *op, + return multiRootTopologicalSort(slice); + } + ++namespace { ++// Copied from TestDeadCodeAnalysis.cpp, because some dead code analysis ++// interacts with constant propagation, but SparseConstantPropagation ++// doesn't seem to be sufficient. ++struct ConstantAnalysis : public DataFlowAnalysis { ++ using DataFlowAnalysis::DataFlowAnalysis; ++ ++ LogicalResult initialize(Operation *top) override { ++ WalkResult result = top->walk([&](Operation *op) { ++ if (failed(visit(op))) ++ return WalkResult::interrupt(); ++ return WalkResult::advance(); ++ }); ++ return success(!result.wasInterrupted()); ++ } ++ ++ LogicalResult visit(ProgramPoint point) override { ++ Operation *op = point.get(); ++ Attribute value; ++ if (matchPattern(op, m_Constant(&value))) { ++ auto *constant = getOrCreate>( ++ op->getResult(0)); ++ propagateIfChanged(constant, constant->join(dataflow::ConstantValue( ++ value, op->getDialect()))); ++ return success(); ++ } ++ setAllToUnknownConstants(op->getResults()); ++ for (Region ®ion : op->getRegions()) ++ setAllToUnknownConstants(region.getArguments()); ++ return success(); ++ } ++ ++ /// Set all given values as not constants. ++ void setAllToUnknownConstants(ValueRange values) { ++ dataflow::ConstantValue unknownConstant(nullptr, nullptr); ++ for (Value value : values) { ++ auto *constant = ++ getOrCreate>(value); ++ propagateIfChanged(constant, constant->join(unknownConstant)); ++ } ++ } ++}; ++} // namespace ++ ++std::unique_ptr createDataFlowSolver() { ++ auto solver = std::make_unique(); ++ solver->load(); ++ solver->load(); ++ return solver; ++} ++ + } // namespace mlir +diff --git a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp +index 6a46265bd7..e352eb3698 100644 +--- a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp ++++ b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp +@@ -159,9 +159,6 @@ struct ConvertLayoutOpConversion + Value smemBase) const { + auto accumNumCTAsEachRep = product(numCTAsEachRep); + auto layout = type.getEncoding(); +- auto blockedLayout = layout.dyn_cast(); +- auto sliceLayout = layout.dyn_cast(); +- auto mmaLayout = layout.dyn_cast(); + auto rank = type.getRank(); + auto sizePerThread = getSizePerThread(layout); + auto accumSizePerThread = product(sizePerThread); +diff --git a/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h b/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h +index 4b89965aa9..1d9e00519b 100644 +--- a/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h ++++ b/lib/Conversion/TritonGPUToLLVM/DotOpHelpers.h +@@ -7,10 +7,8 @@ + #include "mlir/Conversion/LLVMCommon/LoweringOptions.h" + #include "mlir/Conversion/LLVMCommon/Pattern.h" + #include "mlir/Conversion/MathToLLVM/MathToLLVM.h" +-#include "mlir/Conversion/SCFToStandard/SCFToStandard.h" +-#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +-#include "mlir/Dialect/GPU/GPUDialect.h" ++#include "mlir/Dialect/GPU/IR/GPUDialect.h" + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" + #include "mlir/Dialect/Tensor/IR/Tensor.h" + #include "mlir/IR/Matchers.h" +@@ -422,9 +420,9 @@ struct MMA16816ConversionHelper { + MMA16816ConversionHelper(Type dotOperand, MmaEncodingAttr mmaLayout, + Value thread, ConversionPatternRewriter &rewriter, + TypeConverter *typeConverter, Location loc) +- : mmaLayout(mmaLayout), thread(thread), helper(mmaLayout), +- rewriter(rewriter), typeConverter(typeConverter), loc(loc), +- ctx(mmaLayout.getContext()), wpt(mmaLayout.getWarpsPerCTA()) { ++ : mmaLayout(mmaLayout), wpt(mmaLayout.getWarpsPerCTA()), thread(thread), ++ helper(mmaLayout), rewriter(rewriter), typeConverter(typeConverter), ++ loc(loc), ctx(mmaLayout.getContext()) { + helper.deduceMmaType(dotOperand); + + Value _32 = i32_val(32); +diff --git a/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp +index 0f8070ca9f..e4bd47c411 100644 +--- a/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp ++++ b/lib/Conversion/TritonGPUToLLVM/DotOpToLLVM.cpp +@@ -115,8 +115,6 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern { + auto DTensorTy = D.getType().cast(); + auto AShape = ATensorTy.getShape(); + auto BShape = BTensorTy.getShape(); +- auto DShape = DTensorTy.getShape(); +- auto wpt = mmaLayout.getWarpsPerCTA(); + + bool isARow = ALayout.getIsMMAv1Row().cast().getValue(); + bool isBRow = BLayout.getIsMMAv1Row().cast().getValue(); +@@ -221,7 +219,6 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern { + ConversionPatternRewriter &rewriter) const { + auto *ctx = rewriter.getContext(); + auto loc = op.getLoc(); +- auto threadId = getThreadId(rewriter, loc); + + auto A = op.a(); + auto B = op.b(); +@@ -230,12 +227,10 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern { + + auto aTensorTy = A.getType().cast(); + auto bTensorTy = B.getType().cast(); +- auto cTensorTy = C.getType().cast(); + auto dTensorTy = D.getType().cast(); + + auto aShape = aTensorTy.getShape(); + auto bShape = bTensorTy.getShape(); +- auto cShape = cTensorTy.getShape(); + + BlockedEncodingAttr dLayout = + dTensorTy.getEncoding().cast(); +diff --git a/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp +index deb71b9597..0b9e67674b 100644 +--- a/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp ++++ b/lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp +@@ -61,7 +61,6 @@ struct FpToFpOpConversion + convertFp16x4ToFp8x4(Location loc, ConversionPatternRewriter &rewriter, + const Value &v0, const Value &v1, const Value &v2, + const Value &v3) { +- auto ctx = rewriter.getContext(); + auto fp16x2VecTy = vec_ty(f16_ty, 2); + Value fp16x2Vec0 = undef(fp16x2VecTy); + Value fp16x2Vec1 = undef(fp16x2VecTy); +@@ -153,7 +152,6 @@ struct FpToFpOpConversion + convertBf16x4ToFp8x4(Location loc, ConversionPatternRewriter &rewriter, + const Value &v0, const Value &v1, const Value &v2, + const Value &v3) { +- auto ctx = rewriter.getContext(); + auto bf16x2VecTy = vec_ty(i16_ty, 2); + Value bf16x2Vec0 = undef(bf16x2VecTy); + Value bf16x2Vec1 = undef(bf16x2VecTy); +diff --git a/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp +index 9a8b4702bc..bae675f0cb 100644 +--- a/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp ++++ b/lib/Conversion/TritonGPUToLLVM/LoadStoreOpToLLVM.cpp +@@ -109,7 +109,8 @@ struct LoadOpConversion + DenseElementsAttr constAttr; + int64_t splatVal = 0; + if (other && valueElemTy.isa() && +- matchPattern(other, m_Constant(&constAttr)) && constAttr.isSplat()) { ++ matchPattern(other, m_Constant(&constAttr)) && constAttr.isSplat() && ++ constAttr.getElementType().isa()) { + otherIsSplatConstInt = true; + splatVal = constAttr.getSplatValue().getSExtValue(); + } +@@ -333,7 +334,6 @@ struct StoreOpConversion + elem = rewriter.create(loc, type::i8Ty(ctx), elem); + elem = bitcast(elem, valueElemTy); + +- Type u32Ty = typeConverter->convertType(type::u32Ty(ctx)); + llWord = insert_element(wordTy, llWord, elem, i32_val(elemIdx)); + } + llWord = bitcast(llWord, valArgTy); +@@ -387,7 +387,6 @@ struct AtomicCASOpConversion + ConversionPatternRewriter &rewriter) const override { + auto loc = op.getLoc(); + MLIRContext *ctx = rewriter.getContext(); +- Value ptr = op.ptr(); + + Value llPtr = adaptor.ptr(); + Value llCmp = adaptor.cmp(); +diff --git a/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp +index 69abd889be..1c973dc196 100644 +--- a/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp ++++ b/lib/Conversion/TritonGPUToLLVM/ReduceOpToLLVM.cpp +@@ -286,7 +286,6 @@ struct ReduceOpConversion + auto srcTy = op.operand().getType().cast(); + auto srcLayout = srcTy.getEncoding(); + auto srcShape = srcTy.getShape(); +- auto srcRank = srcTy.getRank(); + auto order = getOrder(srcLayout); + + auto threadsPerWarp = triton::gpu::getThreadsPerWarp(srcLayout); +@@ -351,7 +350,6 @@ struct ReduceOpConversion + + Value zero = i32_val(0); + Value laneZero = icmp_eq(laneIdAxis, zero); +- Value warpZero = icmp_eq(warpIdAxis, zero); + + for (auto it : accs) { + const SmallVector &key = it.first; +diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp +index 5b77150b1a..78cfa076bd 100644 +--- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp ++++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp +@@ -11,11 +11,11 @@ using ::mlir::LLVM::getStructFromElements; + using ::mlir::triton::gpu::getElemsPerThread; + using ::mlir::triton::gpu::SharedEncodingAttr; + +-struct ReturnOpConversion : public ConvertOpToLLVMPattern<::mlir::ReturnOp> { +- using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern; ++struct ReturnOpConversion : public ConvertOpToLLVMPattern { ++ using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern; + + LogicalResult +- matchAndRewrite(ReturnOp op, OpAdaptor adaptor, ++ matchAndRewrite(func::ReturnOp op, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + unsigned numArguments = op.getNumOperands(); + +@@ -476,7 +476,6 @@ struct ExtractSliceOpConversion + + auto llvmElemTy = getTypeConverter()->convertType(srcTy.getElementType()); + auto elemPtrTy = ptr_ty(llvmElemTy, 3); +- auto resTy = op.getType().dyn_cast(); + smemObj = SharedMemoryObject(gep(elemPtrTy, smemObj.base, offset), + strideVals, offsetVals); + auto retVal = getStructFromSharedMemoryObject(loc, smemObj, rewriter); +diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h +index bb10d5b24a..00e399f848 100644 +--- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h ++++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h +@@ -4,6 +4,7 @@ + // TODO: refactor so that it doesn't fail if Allocation.h + // is included after utility.h (due to conflict in `store` macro + // and ++#include "mlir/Dialect/Func/IR/FuncOps.h" + #include "triton/Analysis/Allocation.h" + + // +@@ -39,15 +40,15 @@ void vprintf_array(Value thread, ArrayRef arr, std::string info, + // TODO(Superjomn): remove the code when MLIR v15.0 is included. + // All the rights are reserved by the LLVM community. + +-struct FuncOpConversionBase : public ConvertOpToLLVMPattern { ++struct FuncOpConversionBase : public ConvertOpToLLVMPattern { + private: + /// Only retain those attributes that are not constructed by + /// `LLVMFuncOp::build`. If `filterArgAttrs` is set, also filter out argument + /// attributes. +- static void filterFuncAttributes(ArrayRef attrs, +- bool filterArgAttrs, ++ static void filterFuncAttributes(func::FuncOp op, bool filterArgAttrs, + SmallVectorImpl &result) { +- for (const auto &attr : attrs) { ++ ++ for (const auto &attr : op->getAttrs()) { + if (attr.getName() == SymbolTable::getSymbolAttrName() || + attr.getName() == FunctionOpInterface::getTypeAttrName() || + attr.getName() == "std.varargs" || +@@ -65,27 +66,27 @@ struct FuncOpConversionBase : public ConvertOpToLLVMPattern { + } + + protected: +- using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern; ++ using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern; + + // Convert input FuncOp to LLVMFuncOp by using the LLVMTypeConverter provided + // to this legalization pattern. + LLVM::LLVMFuncOp +- convertFuncOpToLLVMFuncOp(FuncOp funcOp, ++ convertFuncOpToLLVMFuncOp(func::FuncOp funcOp, + ConversionPatternRewriter &rewriter) const { + // Convert the original function arguments. They are converted using the + // LLVMTypeConverter provided to this legalization pattern. + auto varargsAttr = funcOp->getAttrOfType("func.varargs"); + TypeConverter::SignatureConversion result(funcOp.getNumArguments()); + auto llvmType = getTypeConverter()->convertFunctionSignature( +- funcOp.getType(), varargsAttr && varargsAttr.getValue(), result); ++ funcOp.getFunctionType(), varargsAttr && varargsAttr.getValue(), ++ result); + if (!llvmType) + return nullptr; + + // Propagate argument/result attributes to all converted arguments/result + // obtained after converting a given original argument/result. + SmallVector attributes; +- filterFuncAttributes(funcOp->getAttrs(), /*filterArgAttrs=*/true, +- attributes); ++ filterFuncAttributes(funcOp, /*filterArgAttrs=*/true, attributes); + if (ArrayAttr resAttrDicts = funcOp.getAllResultAttrs()) { + assert(!resAttrDicts.empty() && "expected array to be non-empty"); + auto newResAttrDicts = +@@ -131,7 +132,7 @@ struct FuncOpConversionBase : public ConvertOpToLLVMPattern { + } + auto newFuncOp = rewriter.create( + funcOp.getLoc(), funcOp.getName(), llvmType, linkage, +- /*dsoLocal*/ false, attributes); ++ /*dsoLocal*/ false, LLVM::CConv::C, attributes); + rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(), + newFuncOp.end()); + if (failed(rewriter.convertRegionTypes(&newFuncOp.getBody(), *typeConverter, +@@ -191,8 +192,8 @@ class ConvertTritonGPUOpToLLVMPatternBase { + const Allocation *allocation, + Value smem, + IndexCacheInfo indexCacheInfo) +- : converter(&typeConverter), indexCacheInfo(indexCacheInfo), +- allocation(allocation), smem(smem) {} ++ : converter(&typeConverter), allocation(allocation), smem(smem), ++ indexCacheInfo(indexCacheInfo) {} + + LLVMTypeConverter *getTypeConverter() const { return converter; } + +@@ -861,7 +862,6 @@ class ConvertTritonGPUOpToLLVMPatternBase { + ArrayRef shape) const { + auto parent = sliceLayout.getParent(); + unsigned dim = sliceLayout.getDim(); +- size_t rank = shape.size(); + auto parentIndices = + emitIndices(loc, rewriter, parent, sliceLayout.paddedShape(shape)); + unsigned numIndices = parentIndices.size(); +diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp +index ff1af09835..6f66af4e34 100644 +--- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp ++++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.cpp +@@ -1,10 +1,11 @@ + #include "triton/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.h" + ++#include "mlir/Analysis/DataFlowFramework.h" + #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" ++#include "mlir/Conversion/ControlFlowToLLVM//ControlFlowToLLVM.h" + #include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h" + #include "mlir/Conversion/MathToLLVM/MathToLLVM.h" +-#include "mlir/Conversion/SCFToStandard/SCFToStandard.h" +-#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h" ++#include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" + #include "mlir/Dialect/LLVMIR/NVVMDialect.h" + #include "mlir/Pass/Pass.h" +@@ -40,7 +41,6 @@ class TritonLLVMConversionTarget : public ConversionTarget { + addIllegalDialect(); + addIllegalDialect(); + addIllegalDialect(); +- addIllegalDialect(); + addLegalOp(); + } + }; +@@ -51,7 +51,7 @@ class TritonLLVMFunctionConversionTarget : public ConversionTarget { + : ConversionTarget(ctx) { + addLegalDialect(); + addLegalDialect(); +- addIllegalOp(); ++ addIllegalOp(); + addLegalOp(); + } + }; +@@ -69,7 +69,7 @@ struct FuncOpConversion : public FuncOpConversionBase { + : FuncOpConversionBase(converter, benefit), numWarps(numWarps) {} + + LogicalResult +- matchAndRewrite(FuncOp funcOp, OpAdaptor adaptor, ++ matchAndRewrite(func::FuncOp funcOp, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + auto newFuncOp = convertFuncOpToLLVMFuncOp(funcOp, rewriter); + if (!newFuncOp) +@@ -133,7 +133,8 @@ class ConvertTritonGPUToLLVM + decomposeBlockedToDotOperand(mod); + + // Step 2 +- decomposeInsertSliceAsyncOp(mod); ++ if (failed(decomposeInsertSliceAsyncOp(mod))) ++ return signalPassFailure(); + + // Step 3 + Allocation allocation(mod); +@@ -142,7 +143,7 @@ class ConvertTritonGPUToLLVM + + // Step 4 + RewritePatternSet scf_patterns(context); +- mlir::populateLoopToStdConversionPatterns(scf_patterns); ++ mlir::populateSCFToControlFlowConversionPatterns(scf_patterns); + mlir::ConversionTarget scf_target(*context); + scf_target.addIllegalOp(); +@@ -159,8 +160,10 @@ class ConvertTritonGPUToLLVM + return signalPassFailure(); + + // Step 6 - get axis and shared memory info +- AxisInfoAnalysis axisInfoAnalysis(mod.getContext()); +- axisInfoAnalysis.run(mod); ++ std::unique_ptr solver = createDataFlowSolver(); ++ AxisInfoAnalysis *axisInfoAnalysis = solver->load(); ++ if (failed(solver->initializeAndRun(mod))) ++ return signalPassFailure(); + initSharedMemory(allocation.getSharedMemorySize(), typeConverter); + mod->setAttr("triton_gpu.shared", + mlir::IntegerAttr::get(mlir::IntegerType::get(context, 32), +@@ -178,38 +181,39 @@ class ConvertTritonGPUToLLVM + + // Normal conversions + populateTritonGPUToLLVMPatterns(typeConverter, patterns, numWarps, +- axisInfoAnalysis, &allocation, smem, ++ *axisInfoAnalysis, &allocation, smem, + indexCacheInfo, /*benefit=*/10); + // ConvertLayoutOp + populateConvertLayoutOpToLLVMPatterns(typeConverter, patterns, numWarps, +- axisInfoAnalysis, &allocation, smem, ++ *axisInfoAnalysis, &allocation, smem, + indexCacheInfo, /*benefit=*/10); + // DotOp + populateDotOpToLLVMPatterns(typeConverter, patterns, numWarps, +- axisInfoAnalysis, &allocation, smem, ++ *axisInfoAnalysis, &allocation, smem, + /*benefit=*/10); + // ElementwiseOp + populateElementwiseOpToLLVMPatterns(typeConverter, patterns, numWarps, +- axisInfoAnalysis, &allocation, smem, ++ *axisInfoAnalysis, &allocation, smem, + /*benefit=*/10); + // LoadStoreOp + populateLoadStoreOpToLLVMPatterns(typeConverter, patterns, numWarps, +- axisInfoAnalysis, &allocation, smem, ++ *axisInfoAnalysis, &allocation, smem, + indexCacheInfo, /*benefit=*/10); + // ReduceOp + populateReduceOpToLLVMPatterns(typeConverter, patterns, numWarps, +- axisInfoAnalysis, &allocation, smem, ++ *axisInfoAnalysis, &allocation, smem, + indexCacheInfo, /*benefit=*/10); + // ViewOp + populateViewOpToLLVMPatterns(typeConverter, patterns, numWarps, +- axisInfoAnalysis, &allocation, smem, ++ *axisInfoAnalysis, &allocation, smem, + /*benefit=*/10); + + // Add arith/math's patterns to help convert scalar expression to LLVM. + mlir::arith::populateArithmeticToLLVMConversionPatterns(typeConverter, + patterns); + mlir::populateMathToLLVMConversionPatterns(typeConverter, patterns); +- mlir::populateStdToLLVMConversionPatterns(typeConverter, patterns); ++ mlir::cf::populateControlFlowToLLVMConversionPatterns(typeConverter, ++ patterns); + mlir::populateGpuToNVVMConversionPatterns(typeConverter, patterns); + + if (failed(applyPartialConversion(mod, target, std::move(patterns)))) +@@ -306,9 +310,11 @@ class ConvertTritonGPUToLLVM + }); + } + +- void decomposeInsertSliceAsyncOp(ModuleOp mod) const { +- AxisInfoAnalysis axisInfoAnalysis(mod.getContext()); +- axisInfoAnalysis.run(mod); ++ LogicalResult decomposeInsertSliceAsyncOp(ModuleOp mod) const { ++ std::unique_ptr solver = createDataFlowSolver(); ++ AxisInfoAnalysis *axisInfoAnalysis = solver->load(); ++ if (failed(solver->initializeAndRun(mod))) ++ return failure(); + // TODO(Keren): This is a hacky knob that may cause performance regression + // when decomposition has been performed. We should remove this knob once we + // have thorough analysis on async wait. Currently, we decompose +@@ -342,7 +348,7 @@ class ConvertTritonGPUToLLVM + auto resSharedLayout = + dstTy.getEncoding().dyn_cast(); + auto resElemTy = dstTy.getElementType(); +- unsigned inVec = axisInfoAnalysis.getPtrContiguity(src); ++ unsigned inVec = axisInfoAnalysis->getPtrContiguity(src); + unsigned outVec = resSharedLayout.getVec(); + unsigned minVec = std::min(outVec, inVec); + auto maxBitWidth = +@@ -400,11 +406,11 @@ class ConvertTritonGPUToLLVM + } else if (decomposed) { + // Wait for all previous async ops + OpBuilder builder(asyncWaitOp); +- auto newAsyncWaitOp = +- builder.create(asyncWaitOp.getLoc(), 0); ++ builder.create(asyncWaitOp.getLoc(), 0); + asyncWaitOp.erase(); + } + }); ++ return success(); + } + }; + +diff --git a/lib/Conversion/TritonGPUToLLVM/Utility.h b/lib/Conversion/TritonGPUToLLVM/Utility.h +index d35dac28c5..11976908cf 100644 +--- a/lib/Conversion/TritonGPUToLLVM/Utility.h ++++ b/lib/Conversion/TritonGPUToLLVM/Utility.h +@@ -220,10 +220,7 @@ struct SharedMemoryObject { + ConversionPatternRewriter &rewriter) + : base(base) { + strides = getStridesFromShapeAndOrder(shape, order, loc, rewriter); +- +- for (auto idx : order) { +- offsets.emplace_back(i32_val(0)); +- } ++ offsets.append(order.size(), i32_val(0)); + } + + SmallVector getElems() const { +diff --git a/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp b/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp +index fe42202c34..5f230f787f 100644 +--- a/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp ++++ b/lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp +@@ -1,10 +1,10 @@ + #include "triton/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.h" + + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +-#include "mlir/Dialect/GPU/GPUDialect.h" ++#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" ++#include "mlir/Dialect/GPU/IR/GPUDialect.h" + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" + #include "mlir/Dialect/LLVMIR/NVVMDialect.h" +-#include "mlir/Dialect/StandardOps/IR/Ops.h" + #include "mlir/Pass/Pass.h" + #include "mlir/Transforms/DialectConversion.h" + #include "triton/Dialect/Triton/IR/Dialect.h" +@@ -59,10 +59,13 @@ class ArithConstantPattern : public OpConversionPattern { + Type retType = getTypeConverter()->convertType(op.getType()); + auto value = adaptor.getValue().dyn_cast(); + assert(value); +- rewriter.replaceOpWithNewOp( +- op, retType, +- value.reshape(retType) // This is a hack. We just want to add encoding +- ); ++ if (value.getElementType().isInteger(1) && value.isSplat()) ++ // Workaround until https://reviews.llvm.org/D133743 is included. ++ value = DenseElementsAttr::get(retType, value.getSplatValue()); ++ else ++ // This is a hack. We just want to add encoding ++ value = value.reshape(retType); ++ rewriter.replaceOpWithNewOp(op, retType, value); + return success(); + } + }; +@@ -127,12 +130,12 @@ void populateArithmeticPatternsAndLegality( + } + + // this shouldn't exist if mlir's SelectOp checked encodings properly +-class StdSelectPattern : public OpConversionPattern { ++class StdSelectPattern : public OpConversionPattern { + public: +- using OpConversionPattern::OpConversionPattern; ++ using OpConversionPattern::OpConversionPattern; + + LogicalResult +- matchAndRewrite(SelectOp op, typename SelectOp::Adaptor adaptor, ++ matchAndRewrite(arith::SelectOp op, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + Type retType = this->getTypeConverter()->convertType(op.getType()); + rewriter.replaceOpWithNewOp( +@@ -148,8 +151,8 @@ void populateStdPatternsAndLegality(TritonGPUTypeConverter &typeConverter, + MLIRContext *context = patterns.getContext(); + // Rewrite rule + patterns.add(typeConverter, context); +- target.addLegalOp(); // this is ok because all functions are inlined +- // by the frontend ++ target.addLegalOp(); // this is ok because all functions are ++ // inlined by the frontend + } + + void populateMathPatternsAndLegality(TritonGPUTypeConverter &typeConverter, +@@ -455,18 +458,19 @@ struct TritonPrintfPattern : public OpConversionPattern { + void populateTritonPatterns(TritonGPUTypeConverter &typeConverter, + RewritePatternSet &patterns) { + MLIRContext *context = patterns.getContext(); +- patterns.add< // TODO: view should have custom pattern that views the layout +- TritonGenericPattern, +- TritonGenericPattern, +- TritonGenericPattern, +- TritonGenericPattern, +- TritonGenericPattern, +- TritonGenericPattern, TritonBroadcastPattern, +- TritonGenericPattern, TritonCatPattern, +- TritonReducePattern, TritonTransPattern, TritonExpandDimsPattern, +- TritonMakeRangePattern, TritonDotPattern, TritonLoadPattern, +- TritonStorePattern, TritonExtElemwisePattern, TritonPrintfPattern, +- TritonAtomicRMWPattern>(typeConverter, context); ++ patterns ++ .insert< // TODO: view should have custom pattern that views the layout ++ TritonGenericPattern, ++ TritonGenericPattern, ++ TritonGenericPattern, ++ TritonGenericPattern, ++ TritonGenericPattern, ++ TritonGenericPattern, TritonBroadcastPattern, ++ TritonGenericPattern, TritonCatPattern, ++ TritonReducePattern, TritonTransPattern, TritonExpandDimsPattern, ++ TritonMakeRangePattern, TritonDotPattern, TritonLoadPattern, ++ TritonStorePattern, TritonExtElemwisePattern, TritonPrintfPattern, ++ TritonAtomicRMWPattern>(typeConverter, context); + } + + // +@@ -623,29 +627,28 @@ void populateSCFPatterns(TritonGPUTypeConverter &typeConverter, + + // CF + +-class CFBranchPattern : public OpConversionPattern { ++class CFBranchPattern : public OpConversionPattern { + public: +- using OpConversionPattern::OpConversionPattern; ++ using OpConversionPattern::OpConversionPattern; + + LogicalResult +- matchAndRewrite(BranchOp op, BranchOp::Adaptor adaptor, ++ matchAndRewrite(cf::BranchOp op, cf::BranchOp::Adaptor adaptor, + ConversionPatternRewriter &rewriter) const override { +- auto converter = getTypeConverter(); +- auto newOp = rewriter.replaceOpWithNewOp(op, op.getSuccessor(), +- adaptor.getOperands()); ++ auto newOp = rewriter.replaceOpWithNewOp( ++ op, op.getSuccessor(), adaptor.getOperands()); + return success(); + } + }; + +-class CFCondBranchPattern : public OpConversionPattern { ++class CFCondBranchPattern : public OpConversionPattern { + public: +- using OpConversionPattern::OpConversionPattern; ++ using OpConversionPattern::OpConversionPattern; + + LogicalResult +- matchAndRewrite(CondBranchOp op, CondBranchOp::Adaptor adaptor, ++ matchAndRewrite(cf::CondBranchOp op, cf::CondBranchOp::Adaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + auto converter = getTypeConverter(); +- auto newOp = rewriter.replaceOpWithNewOp( ++ auto newOp = rewriter.replaceOpWithNewOp( + op, adaptor.getCondition(), op.getTrueDest(), + adaptor.getTrueDestOperands(), op.getFalseDest(), + adaptor.getFalseDestOperands()); +diff --git a/lib/Dialect/Triton/IR/CMakeLists.txt b/lib/Dialect/Triton/IR/CMakeLists.txt +index 2d679b21fd..705554ba6b 100644 +--- a/lib/Dialect/Triton/IR/CMakeLists.txt ++++ b/lib/Dialect/Triton/IR/CMakeLists.txt +@@ -10,11 +10,7 @@ add_mlir_dialect_library(TritonIR + + LINK_LIBS PUBLIC + MLIRIR +- MLIRArithmetic +- MLIRSCF +- +- # Since LLVM 15 +- # MLIRFunc +- # else +- MLIRStandard ++ MLIRArithmeticDialect ++ MLIRSCFDialect ++ MLIRFuncDialect + ) +diff --git a/lib/Dialect/Triton/IR/Ops.cpp b/lib/Dialect/Triton/IR/Ops.cpp +index 3aadbfa0c0..86570359c5 100644 +--- a/lib/Dialect/Triton/IR/Ops.cpp ++++ b/lib/Dialect/Triton/IR/Ops.cpp +@@ -1,10 +1,9 @@ +-#include "triton/Dialect/Triton/IR/Dialect.h" +-#include "triton/Dialect/Triton/IR/Types.h" +- + #include "mlir/IR/Builders.h" + #include "mlir/IR/BuiltinAttributes.h" + #include "mlir/IR/BuiltinTypes.h" + #include "mlir/IR/OperationSupport.h" ++#include "triton/Dialect/Triton/IR/Dialect.h" ++#include "triton/Dialect/Triton/IR/Types.h" + + namespace mlir { + namespace triton { +@@ -38,8 +37,8 @@ static Type getPointerTypeSameShape(Type type) { + } + + // Parser & printer for assembly forms +-ParseResult parseLoadOp(OpAsmParser &parser, OperationState &result) { +- SmallVector allOperands; ++ParseResult LoadOp::parse(OpAsmParser &parser, OperationState &result) { ++ SmallVector allOperands; + Type resultTypes[1]; + SMLoc allOperandLoc = parser.getCurrentLocation(); + if (parser.parseOperandList(allOperands) || +@@ -73,18 +72,18 @@ ParseResult parseLoadOp(OpAsmParser &parser, OperationState &result) { + return success(); + } + +-void printLoadOp(OpAsmPrinter &printer, LoadOp loadOp) { ++void LoadOp::print(OpAsmPrinter &printer) { + printer << " "; +- printer << loadOp.getOperation()->getOperands(); ++ printer << getOperation()->getOperands(); + // "operand_segment_sizes" can be deduced, so we don't print it. +- printer.printOptionalAttrDict(loadOp->getAttrs(), +- {loadOp.operand_segment_sizesAttrName()}); ++ printer.printOptionalAttrDict(getOperation()->getAttrs(), ++ {operand_segment_sizesAttrName()}); + printer << " : "; +- printer.printStrippedAttrOrType(loadOp.result().getType()); ++ printer.printStrippedAttrOrType(getResult().getType()); + } + +-ParseResult parseStoreOp(OpAsmParser &parser, OperationState &result) { +- SmallVector allOperands; ++ParseResult StoreOp::parse(OpAsmParser &parser, OperationState &result) { ++ SmallVector allOperands; + Type valueType; + SMLoc allOperandLoc = parser.getCurrentLocation(); + if (parser.parseOperandList(allOperands) || +@@ -104,12 +103,12 @@ ParseResult parseStoreOp(OpAsmParser &parser, OperationState &result) { + return success(); + } + +-void printStoreOp(OpAsmPrinter &printer, StoreOp storeOp) { ++void StoreOp::print(OpAsmPrinter &printer) { + printer << " "; +- printer << storeOp.getOperation()->getOperands(); +- printer.printOptionalAttrDict(storeOp->getAttrs(), /*elidedAttrs=*/{}); ++ printer << getOperation()->getOperands(); ++ printer.printOptionalAttrDict(getOperation()->getAttrs(), /*elidedAttrs=*/{}); + printer << " : "; +- printer.printStrippedAttrOrType(storeOp.value().getType()); ++ printer.printStrippedAttrOrType(value().getType()); + } + + } // namespace triton +@@ -319,7 +318,8 @@ OpFoldResult SplatOp::fold(ArrayRef operands) { + if (!constOperand) + return {}; + auto shapedType = getType().cast(); +- auto ret = SplatElementsAttr::get(shapedType, {constOperand.getValue()}); ++ auto ret = SplatElementsAttr::get( ++ shapedType, ArrayRef(constOperand.getValue())); + return ret; + } + +diff --git a/lib/Dialect/Triton/Transforms/Combine.cpp b/lib/Dialect/Triton/Transforms/Combine.cpp +index 2261472170..11570283d6 100644 +--- a/lib/Dialect/Triton/Transforms/Combine.cpp ++++ b/lib/Dialect/Triton/Transforms/Combine.cpp +@@ -57,13 +57,13 @@ DenseElementsAttr getConstantValue(Builder &builder, Attribute value, + class CombineSelectMaskedLoadPattern : public mlir::RewritePattern { + public: + CombineSelectMaskedLoadPattern(mlir::MLIRContext *context) +- : mlir::RewritePattern(mlir::SelectOp::getOperationName(), 3, context, +- {triton::LoadOp::getOperationName()}) {} ++ : mlir::RewritePattern(mlir::arith::SelectOp::getOperationName(), 3, ++ context, {triton::LoadOp::getOperationName()}) {} + + mlir::LogicalResult + matchAndRewrite(mlir::Operation *op, + mlir::PatternRewriter &rewriter) const override { +- auto selectOp = llvm::dyn_cast(op); ++ auto selectOp = llvm::dyn_cast(op); + if (!selectOp) + return mlir::failure(); + +diff --git a/lib/Dialect/Triton/Transforms/Combine.td b/lib/Dialect/Triton/Transforms/Combine.td +index 14f286b26e..ded0e346e6 100644 +--- a/lib/Dialect/Triton/Transforms/Combine.td ++++ b/lib/Dialect/Triton/Transforms/Combine.td +@@ -1,9 +1,9 @@ + #ifndef TRITON_PATTERNS + #define TRITON_PATTERNS + +-include "mlir/Dialect/StandardOps/IR/Ops.td" + include "mlir/Dialect/Arithmetic/IR/ArithmeticOps.td" + include "triton/Dialect/Triton/IR/TritonOps.td" ++include "mlir/IR/PatternBase.td" + + + // AddIOp(DotOp(a, b, c), d) and c==0 => DotOp(a, b, d) +diff --git a/lib/Dialect/TritonGPU/IR/Dialect.cpp b/lib/Dialect/TritonGPU/IR/Dialect.cpp +index 1fbc609e88..bfc3f3d3da 100644 +--- a/lib/Dialect/TritonGPU/IR/Dialect.cpp ++++ b/lib/Dialect/TritonGPU/IR/Dialect.cpp +@@ -1,14 +1,14 @@ ++#include "triton/Dialect/Triton/IR/Dialect.h" ++ + #include + + #include "mlir/IR/DialectImplementation.h" + #include "mlir/IR/OpImplementation.h" + #include "triton/Analysis/Utility.h" +-#include "triton/Dialect/Triton/IR/Dialect.h" ++#include "triton/Dialect/TritonGPU/IR/Dialect.cpp.inc" + #include "triton/Dialect/TritonGPU/IR/Dialect.h" + #include "llvm/ADT/TypeSwitch.h" + +-#include "triton/Dialect/TritonGPU/IR/Dialect.cpp.inc" +- + using namespace mlir; + using namespace mlir::triton::gpu; + +@@ -366,7 +366,6 @@ template SmallVector + SliceEncodingAttr::paddedShape(ArrayRef shape) const; + + unsigned SliceEncodingAttr::getElemsPerThread(ArrayRef shape) const { +- size_t rank = shape.size(); + auto parent = getParent(); + return ::getElemsPerThread(parent, paddedShape(shape)); + } +@@ -655,9 +654,9 @@ void DotOperandEncodingAttr::print(mlir::AsmPrinter &printer) const { + // InsertSliceAsyncOp + //===----------------------------------------------------------------------===// + +-ParseResult parseInsertSliceAsyncOp(OpAsmParser &parser, +- OperationState &result) { +- SmallVector allOperands; ++ParseResult InsertSliceAsyncOp::parse(OpAsmParser &parser, ++ OperationState &result) { ++ SmallVector allOperands; + Type srcType, dstType; + SMLoc allOperandLoc = parser.getCurrentLocation(); + if (parser.parseOperandList(allOperands) || +@@ -696,18 +695,16 @@ ParseResult parseInsertSliceAsyncOp(OpAsmParser &parser, + return success(); + } + +-void printInsertSliceAsyncOp(OpAsmPrinter &printer, +- InsertSliceAsyncOp insertSliceAsyncOp) { ++void InsertSliceAsyncOp::print(OpAsmPrinter &printer) { + printer << " "; +- printer << insertSliceAsyncOp.getOperation()->getOperands(); ++ printer << getOperation()->getOperands(); + // "operand_segment_sizes" can be deduced, so we don't print it. +- printer.printOptionalAttrDict( +- insertSliceAsyncOp->getAttrs(), +- {insertSliceAsyncOp.operand_segment_sizesAttrName()}); ++ printer.printOptionalAttrDict(getOperation()->getAttrs(), ++ {operand_segment_sizesAttrName()}); + printer << " : "; +- printer.printStrippedAttrOrType(insertSliceAsyncOp.src().getType()); ++ printer.printStrippedAttrOrType(src().getType()); + printer << " -> "; +- printer.printStrippedAttrOrType(insertSliceAsyncOp.result().getType()); ++ printer.printStrippedAttrOrType(result().getType()); + } + + //===----------------------------------------------------------------------===// +diff --git a/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp b/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp +index 82407980d3..ee6009f44a 100644 +--- a/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp ++++ b/lib/Dialect/TritonGPU/Transforms/Coalesce.cpp +@@ -27,7 +27,11 @@ struct CoalescePass : public TritonGPUCoalesceBase { + auto origType = ptr.getType().cast(); + // Get the shape of the tensor. + size_t rank = origType.getRank(); +- AxisInfo info = axisInfo.lookupLatticeElement(ptr)->getValue(); ++ dataflow::Lattice *latticeElement = ++ axisInfo.getLatticeElement(ptr); ++ AxisInfo info = latticeElement && !latticeElement->isUninitialized() ++ ? latticeElement->getValue() ++ : AxisInfo(); + // Get the contiguity order of `ptr` + auto order = argSort(info.getContiguity()); + // The desired divisibility is the maximum divisibility +@@ -40,7 +44,7 @@ struct CoalescePass : public TritonGPUCoalesceBase { + for (Value val : op->getResults()) { + if (val.getType() != origType) + continue; +- auto valInfo = axisInfo.lookupLatticeElement(val); ++ auto valInfo = axisInfo.getLatticeElement(val); + auto currOrder = argSort(valInfo->getValue().getContiguity()); + if (order == currOrder) + withSameOrder.insert(val); +@@ -55,7 +59,7 @@ struct CoalescePass : public TritonGPUCoalesceBase { + unsigned elemNumBytes = std::max(elemNumBits / 8, 1u); + unsigned perThread = 1; + for (Value val : withSameOrder) { +- AxisInfo info = axisInfo.lookupLatticeElement(val)->getValue(); ++ AxisInfo info = axisInfo.getLatticeElement(val)->getValue(); + unsigned maxMultipleBytes = info.getDivisibility(order[0]); + unsigned maxMultiple = std::max(maxMultipleBytes / elemNumBytes, 1u); + unsigned maxContig = info.getContiguity(order[0]); +@@ -123,8 +127,10 @@ struct CoalescePass : public TritonGPUCoalesceBase { + void runOnOperation() override { + Operation *op = getOperation(); + // Run axis info analysis +- AxisInfoAnalysis axisInfo(&getContext()); +- axisInfo.run(op); ++ std::unique_ptr solver = createDataFlowSolver(); ++ AxisInfoAnalysis *axisInfo = solver->load(); ++ if (failed(solver->initializeAndRun(op))) ++ return signalPassFailure(); + + // For each i/o operation, we determine what layout + // the pointers should have for best memory coalescing +@@ -146,10 +152,10 @@ struct CoalescePass : public TritonGPUCoalesceBase { + RankedTensorType ty = ptr.getType().template dyn_cast(); + if (!ty || !ty.getElementType().isa()) + return; +- AxisInfo info = axisInfo.lookupLatticeElement(ptr)->getValue(); ++ AxisInfo info = axisInfo->getLatticeElement(ptr)->getValue(); + auto mod = curr->getParentOfType(); + int numWarps = triton::gpu::TritonGPUDialect::getNumWarps(mod); +- auto convertType = getTypeConverter(axisInfo, ptr, numWarps); ++ auto convertType = getTypeConverter(*axisInfo, ptr, numWarps); + layoutMap[ptr] = convertType; + }); + +diff --git a/lib/Dialect/TritonGPU/Transforms/Combine.cpp b/lib/Dialect/TritonGPU/Transforms/Combine.cpp +index efa37ff2dc..089ce3996c 100644 +--- a/lib/Dialect/TritonGPU/Transforms/Combine.cpp ++++ b/lib/Dialect/TritonGPU/Transforms/Combine.cpp +@@ -1,6 +1,6 @@ + #include "Utility.h" + #include "mlir/Analysis/SliceAnalysis.h" +-#include "mlir/Dialect/SCF/SCF.h" ++#include "mlir/Dialect/SCF/IR/SCF.h" + #include "mlir/IR/BlockAndValueMapping.h" + #include "mlir/IR/BuiltinAttributes.h" + #include "mlir/IR/Matchers.h" +diff --git a/lib/Dialect/TritonGPU/Transforms/Combine.td b/lib/Dialect/TritonGPU/Transforms/Combine.td +index 6bf1b14866..6a7b10dbcb 100644 +--- a/lib/Dialect/TritonGPU/Transforms/Combine.td ++++ b/lib/Dialect/TritonGPU/Transforms/Combine.td +@@ -3,5 +3,6 @@ + + include "triton/Dialect/TritonGPU/IR/TritonGPUOps.td" + include "triton/Dialect/Triton/IR/TritonOps.td" ++include "mlir/IR/PatternBase.td" + + #endif +diff --git a/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp b/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp +index 4bd3bc76bf..b2f8defd81 100644 +--- a/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp ++++ b/lib/Dialect/TritonGPU/Transforms/DecomposeConversions.cpp +@@ -1,5 +1,5 @@ + #include "mlir/Analysis/SliceAnalysis.h" +-#include "mlir/Dialect/SCF/SCF.h" ++#include "mlir/Dialect/SCF/IR/SCF.h" + #include "mlir/IR/BlockAndValueMapping.h" + #include "mlir/IR/BuiltinAttributes.h" + #include "mlir/IR/Matchers.h" +diff --git a/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp b/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp +index 9b2f42231e..85f746c1dc 100644 +--- a/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp ++++ b/lib/Dialect/TritonGPU/Transforms/Pipeline.cpp +@@ -2,6 +2,7 @@ + #include "mlir/IR/BlockAndValueMapping.h" + #include "mlir/IR/TypeUtilities.h" + #include "triton/Analysis/AxisInfo.h" ++#include "triton/Analysis/Utility.h" + #include "triton/Dialect/TritonGPU/IR/Dialect.h" + #include "triton/Dialect/TritonGPU/Transforms/Passes.h" + +@@ -160,15 +161,18 @@ ttg::AllocTensorOp LoopPipeliner::allocateEmptyBuffer(Operation *op, + LogicalResult LoopPipeliner::initialize() { + Block *loop = forOp.getBody(); + +- AxisInfoAnalysis axisInfoAnalysis(forOp.getContext()); +- axisInfoAnalysis.run(forOp->getParentOfType()); ++ std::unique_ptr solver = createDataFlowSolver(); ++ AxisInfoAnalysis *axisInfoAnalysis = solver->load(); ++ if (failed(solver->initializeAndRun(forOp->getParentOfType()))) { ++ return failure(); ++ } + + // can we use forOp.walk(...) here? + SmallVector allLoads; + for (Operation &op : *loop) + if (auto loadOp = dyn_cast(&op)) { + auto ptr = loadOp.ptr(); +- unsigned vec = axisInfoAnalysis.getPtrContiguity(ptr); ++ unsigned vec = axisInfoAnalysis->getPtrContiguity(ptr); + auto tensorTy = ptr.getType().dyn_cast(); + if (!tensorTy) + continue; +diff --git a/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp b/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp +index 0e7dbe5264..b95a4f50a6 100644 +--- a/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp ++++ b/lib/Dialect/TritonGPU/Transforms/ReorderInstructions.cpp +@@ -1,5 +1,5 @@ + #include "mlir/Analysis/SliceAnalysis.h" +-#include "mlir/Dialect/SCF/SCF.h" ++#include "mlir/Dialect/SCF/IR/SCF.h" + #include "mlir/IR/BlockAndValueMapping.h" + #include "mlir/IR/BuiltinAttributes.h" + #include "mlir/IR/Matchers.h" +diff --git a/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp b/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp +index 37ac710995..762e887f36 100644 +--- a/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp ++++ b/lib/Dialect/TritonGPU/Transforms/TritonGPUConversion.cpp +@@ -82,12 +82,12 @@ TritonGPUConversionTarget::TritonGPUConversionTarget( + scf::ReduceReturnOp>(); + + addDynamicallyLegalDialect([&](Operation *op) { +- if (typeConverter.isLegal(op)) +- return true; +- return false; +- }); ++ triton::TritonDialect, scf::SCFDialect>( ++ [&](Operation *op) { ++ if (typeConverter.isLegal(op)) ++ return true; ++ return false; ++ }); + + // We have requirements for the data layouts + addDynamicallyLegalOp([](triton::DotOp dotOp) -> bool { +diff --git a/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp b/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp +index c229104286..c911fd4a5c 100644 +--- a/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp ++++ b/lib/Dialect/TritonGPU/Transforms/UpdateMmaForVolta.cpp +@@ -1,5 +1,5 @@ + #include "Utility.h" +-#include "mlir/Dialect/SCF/SCF.h" ++#include "mlir/Dialect/SCF/IR/SCF.h" + #include "mlir/IR/Matchers.h" + #include "mlir/IR/PatternMatch.h" + #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +@@ -118,8 +118,8 @@ void setOpResultType(Operation *op, ArrayRef newTypes) { + .get("value") + .dyn_cast(); + if (attr) { +- auto newAttr = mlir::DenseElementsAttr::getFromRawBuffer( +- newType, attr.getRawData(), true); ++ auto newAttr = ++ mlir::DenseElementsAttr::getFromRawBuffer(newType, attr.getRawData()); + op->setAttr("value", newAttr); + } + } +diff --git a/lib/Dialect/TritonGPU/Transforms/Utility.cpp b/lib/Dialect/TritonGPU/Transforms/Utility.cpp +index ed15f02f67..6400f1633a 100644 +--- a/lib/Dialect/TritonGPU/Transforms/Utility.cpp ++++ b/lib/Dialect/TritonGPU/Transforms/Utility.cpp +@@ -1,5 +1,5 @@ + #include "Utility.h" +-#include "mlir/Dialect/SCF/SCF.h" ++#include "mlir/Dialect/SCF/IR/SCF.h" + #include "mlir/IR/BlockAndValueMapping.h" + #include "mlir/Transforms/GreedyPatternRewriteDriver.h" + +diff --git a/lib/Target/LLVMIR/CMakeLists.txt b/lib/Target/LLVMIR/CMakeLists.txt +index f1bbd0bf4e..ac8973ad19 100644 +--- a/lib/Target/LLVMIR/CMakeLists.txt ++++ b/lib/Target/LLVMIR/CMakeLists.txt +@@ -6,8 +6,7 @@ add_mlir_translation_library(TritonLLVMIR + + LINK_LIBS PUBLIC + MLIRIR +- MLIRLLVMIR +- MLIRSCFToStandard ++ MLIRLLVMDialect + MLIRSupport + MLIRTargetLLVMIRExport + ) +diff --git a/lib/Target/PTX/PTXTranslation.cpp b/lib/Target/PTX/PTXTranslation.cpp +index 4cb0d8193c..6a5453a6e7 100644 +--- a/lib/Target/PTX/PTXTranslation.cpp ++++ b/lib/Target/PTX/PTXTranslation.cpp +@@ -1,11 +1,14 @@ + #include "triton/Target/PTX/PTXTranslation.h" + #include "triton/Target/LLVMIR/LLVMIRTranslation.h" ++#include + + #include "llvm/IR/IRBuilder.h" + #include "llvm/IR/LegacyPassManager.h" + #include "llvm/IR/Module.h" + #include "llvm/IR/Verifier.h" + #include "llvm/MC/TargetRegistry.h" ++#include "llvm/Pass.h" ++#include "llvm/Support/CommandLine.h" + #include "llvm/Support/TargetSelect.h" + #include "llvm/Target/TargetMachine.h" + +diff --git a/python/setup.py b/python/setup.py +index 2ac3accd25..4530b36714 100644 +--- a/python/setup.py ++++ b/python/setup.py +@@ -57,19 +57,10 @@ def get_pybind11_package_info(): + def get_llvm_package_info(): + # download if nothing is installed + system = platform.system() +- if system == "Darwin": +- system_suffix = "apple-darwin" +- elif system == "Linux": +- vglibc = tuple(map(int, platform.libc_ver()[1].split('.'))) +- vglibc = vglibc[0] * 100 + vglibc[1] +- linux_suffix = 'ubuntu-18.04' if vglibc > 217 else 'centos-7' +- system_suffix = f"linux-gnu-{linux_suffix}" +- else: +- raise RuntimeError(f"unsupported system: {system}") ++ system_suffix = {"Linux": "linux-gnu-ubuntu-18.04", "Darwin": "apple-darwin"}[system] + use_assert_enabled_llvm = check_env_flag("TRITON_USE_ASSERT_ENABLED_LLVM", "False") +- release_suffix = "assert" if use_assert_enabled_llvm else "release" +- name = f'llvm+mlir-14.0.6-x86_64-{system_suffix}-{release_suffix}' +- url = f"https://github.com/ptillet/triton-llvm-releases/releases/download/llvm-14.0.6-f28c006a5895/{name}.tar.xz" ++ name = 'llvm+mlir-15.0.7-x86_64-{}-{}'.format(system_suffix, "assert" if use_assert_enabled_llvm else "release") ++ url = "https://github.com/ptillet/triton-llvm-releases/releases/download/llvm-15.0.7-8dfdcc7b7bf6/{}.tar.xz".format(name) + return Package("llvm", name, url, "lib", "LLVM_INCLUDE_DIRS", "LLVM_LIBRARY_DIR", "LLVM_SYSPATH") + + +diff --git a/python/src/triton.cc b/python/src/triton.cc +index c40b117a55..f190eacc34 100644 +--- a/python/src/triton.cc ++++ b/python/src/triton.cc +@@ -8,9 +8,10 @@ + #include "mlir/Pass/PassManager.h" + #include "mlir/Transforms/Passes.h" + +-#include "mlir/Parser.h" ++#include "mlir/Parser/Parser.h" + #include "mlir/Support/FileUtilities.h" + ++#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" + #include "triton/Analysis/Allocation.h" + #include "triton/Conversion/TritonGPUToLLVM/TritonGPUToLLVMPass.h" +@@ -195,7 +196,7 @@ void init_triton_ir(py::module &&m) { + std::string attrName = name + "_arg" + std::to_string(id); + mlir::Block *owner = arg.getOwner(); + if (owner->isEntryBlock() && +- !mlir::isa(owner->getParentOp())) { ++ !mlir::isa(owner->getParentOp())) { + owner->getParentOp()->setAttr(attrName, attr); + } + } +@@ -348,7 +349,7 @@ void init_triton_ir(py::module &&m) { + return str; + }) + .def("push_back", +- [](mlir::ModuleOp &self, mlir::FuncOp &funcOp) -> void { ++ [](mlir::ModuleOp &self, mlir::func::FuncOp &funcOp) -> void { + self.push_back(funcOp); + }) + .def("has_function", +@@ -358,16 +359,18 @@ void init_triton_ir(py::module &&m) { + return false; + }) + .def("get_function", +- [](mlir::ModuleOp &self, std::string &funcName) -> mlir::FuncOp { +- return self.lookupSymbol(funcName); +- }) +- .def("get_single_function", [](mlir::ModuleOp &self) -> mlir::FuncOp { +- llvm::SmallVector funcs; +- self.walk([&](mlir::FuncOp func) { funcs.push_back(func); }); +- if (funcs.size() != 1) +- throw std::runtime_error("Expected a single function"); +- return funcs[0]; +- }); ++ [](mlir::ModuleOp &self, ++ std::string &funcName) -> mlir::func::FuncOp { ++ return self.lookupSymbol(funcName); ++ }) ++ .def("get_single_function", ++ [](mlir::ModuleOp &self) -> mlir::func::FuncOp { ++ llvm::SmallVector funcs; ++ self.walk([&](mlir::func::FuncOp func) { funcs.push_back(func); }); ++ if (funcs.size() != 1) ++ throw std::runtime_error("Expected a single function"); ++ return funcs[0]; ++ }); + + m.def("make_attr", + [](const std::vector &values, mlir::MLIRContext &context) { +@@ -388,47 +391,48 @@ void init_triton_ir(py::module &&m) { + registry.insert(); ++ mlir::func::FuncDialect, mlir::scf::SCFDialect>(); + context.appendDialectRegistry(registry); + context.loadAllAvailableDialects(); + + // parse module +- mlir::OwningOpRef module( +- mlir::parseSourceFile(inputFilename, &context)); ++ mlir::OwningOpRef module = ++ mlir::parseSourceFile(inputFilename, &context); ++ if (!module) ++ throw std::runtime_error("Parse MLIR file failed."); + // locations are incompatible with ptx < 7.5 ! + module->walk([](mlir::Operation *op) { + op->setLoc(mlir::UnknownLoc::get(op->getContext())); + }); +- if (!module) +- throw std::runtime_error("Parse MLIR file failed."); + + return module->clone(); + }, + ret::take_ownership); + +- py::class_(m, "function") ++ py::class_(m, "function") + // .def_property_readonly("attrs", &ir::function::attrs) + // .def("add_attr", &ir::function::add_attr); + .def("args", +- [](mlir::FuncOp &self, unsigned idx) -> mlir::BlockArgument { ++ [](mlir::func::FuncOp &self, unsigned idx) -> mlir::BlockArgument { + return self.getArgument(idx); + }) + .def( + "add_entry_block", +- [](mlir::FuncOp &self) -> mlir::Block * { ++ [](mlir::func::FuncOp &self) -> mlir::Block * { + return self.addEntryBlock(); + }, + ret::reference) + .def( + "set_arg_attr", +- [](mlir::FuncOp &self, int arg_no, const std::string &name, int val) { ++ [](mlir::func::FuncOp &self, int arg_no, const std::string &name, ++ int val) { + // set arg attributes "name" to value "val" + auto attrTy = mlir::IntegerType::get(self.getContext(), 32); + self.setArgAttr(arg_no, name, mlir::IntegerAttr::get(attrTy, val)); + }, + ret::reference) +- .def_property_readonly("type", &mlir::FuncOp::getType) +- .def("reset_type", &mlir::FuncOp::setType); ++ .def_property_readonly("type", &mlir::func::FuncOp::getFunctionType) ++ .def("reset_type", &mlir::func::FuncOp::setType); + + py::class_(m, "InsertPoint"); + +@@ -445,13 +449,13 @@ void init_triton_ir(py::module &&m) { + .def("ret", + [](mlir::OpBuilder &self, std::vector &vals) -> void { + auto loc = self.getUnknownLoc(); +- self.create(loc, vals); ++ self.create(loc, vals); + }) + .def("call", +- [](mlir::OpBuilder &self, mlir::FuncOp &func, ++ [](mlir::OpBuilder &self, mlir::func::FuncOp &func, + std::vector &args) -> mlir::OpState { + auto loc = self.getUnknownLoc(); +- return self.create(loc, func, args); ++ return self.create(loc, func, args); + }) + // insertion block/point + .def("set_insertion_point_to_start", +@@ -618,15 +622,16 @@ void init_triton_ir(py::module &&m) { + .def("get_or_insert_function", + [](mlir::OpBuilder &self, mlir::ModuleOp &module, + std::string &funcName, mlir::Type &funcType, +- std::string &visibility) -> mlir::FuncOp { ++ std::string &visibility) -> mlir::func::FuncOp { + if (mlir::Operation *funcOperation = module.lookupSymbol(funcName)) +- return llvm::dyn_cast(funcOperation); ++ return llvm::dyn_cast(funcOperation); + auto loc = self.getUnknownLoc(); + if (auto funcTy = funcType.dyn_cast()) { + llvm::SmallVector attrs = { + mlir::NamedAttribute(self.getStringAttr("sym_visibility"), + self.getStringAttr(visibility))}; +- return self.create(loc, funcName, funcTy, attrs); ++ return self.create(loc, funcName, funcTy, ++ attrs); + } + throw std::runtime_error("invalid function type"); + }) +@@ -658,15 +663,15 @@ void init_triton_ir(py::module &&m) { + [](mlir::OpBuilder &self, mlir::Value condition, + mlir::Block *trueDest, mlir::Block *falseDest) { + auto loc = self.getUnknownLoc(); +- self.create(loc, condition, trueDest, +- falseDest); ++ self.create(loc, condition, trueDest, ++ falseDest); + return; + }) + .def("create_branch", + [](mlir::OpBuilder &self, mlir::Block *dest, + std::vector &args) { + auto loc = self.getUnknownLoc(); +- self.create(loc, dest, args); ++ self.create(loc, dest, args); + return; + }) + // Structured control flow +@@ -792,14 +797,14 @@ void init_triton_ir(py::module &&m) { + .def("create_to_index", + [](mlir::OpBuilder &self, mlir::Value &input) -> mlir::Value { + auto loc = self.getUnknownLoc(); +- return self.create(loc, input, +- self.getIndexType()); ++ return self.create( ++ loc, self.getIndexType(), input); + }) + .def("create_index_to_si", + [](mlir::OpBuilder &self, mlir::Value &input) -> mlir::Value { + auto loc = self.getUnknownLoc(); +- return self.create(loc, input, +- self.getI32Type()); ++ return self.create( ++ loc, self.getI32Type(), input); + }) + .def("create_fmul", + [](mlir::OpBuilder &self, mlir::Value &lhs, +@@ -1316,8 +1321,8 @@ void init_triton_ir(py::module &&m) { + [](mlir::OpBuilder &self, mlir::Value &condition, + mlir::Value &trueValue, mlir::Value &falseValue) -> mlir::Value { + auto loc = self.getUnknownLoc(); +- return self.create(loc, condition, trueValue, +- falseValue); ++ return self.create(loc, condition, ++ trueValue, falseValue); + }) + .def("create_printf", + [](mlir::OpBuilder &self, const std::string &prefix, +@@ -1429,7 +1434,7 @@ void init_triton_ir(py::module &&m) { + self.addPass(mlir::triton::createConvertTritonGPUToLLVMPass()); + }) + .def("add_scf_to_cfg", [](mlir::PassManager &self) { +- self.addPass(mlir::createLowerToCFGPass()); ++ self.addPass(mlir::createConvertSCFToCFPass()); + }); + } + +diff --git a/python/test/unit/language/test_core.py b/python/test/unit/language/test_core.py +index 432544a8a4..018f544714 100644 +--- a/python/test/unit/language/test_core.py ++++ b/python/test/unit/language/test_core.py +@@ -1918,7 +1918,7 @@ def test_convert2d(dtype, shape, src_layout, dst_layout, device='cuda'): + #dst = {dst_layout} + """ + """ + module attributes {"triton_gpu.num-warps" = 4 : i32} { +- func public @kernel_0d1d(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { ++ func.func public @kernel_0d1d(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { + %cst = arith.constant dense<128> : tensor<128x1xi32, #src> + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #triton_gpu.slice<{dim = 1, parent = #src}>> + %1 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #triton_gpu.slice<{dim = 0, parent = #src}>> +diff --git a/python/triton/compiler.py b/python/triton/compiler.py +index 5d167634df..c36589037c 100644 +--- a/python/triton/compiler.py ++++ b/python/triton/compiler.py +@@ -1514,14 +1514,14 @@ def make_hash(fn, **kwargs): + return hashlib.md5((Path(fn).read_text() + triton.runtime.jit.version_key()).encode("utf-8")).hexdigest() + + +-# - ^\s*func\s+ : match the start of the string, any leading whitespace, the keyword func, ++# - ^\s*func\.func\s+ : match the start of the string, any leading whitespace, the keyword func, + # and any following whitespace + # - (public\s+)? : optionally match the keyword public and any following whitespace + # - (@\w+) : match an @ symbol followed by one or more word characters + # (letters, digits, or underscores), and capture it as group 1 (the function name) + # - (\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\)) : match a pair of parentheses enclosing + # zero or more arguments separated by commas, and capture it as group 2 (the argument list) +-mlir_prototype_pattern = r'^\s*func\s+(?:public\s+)?(@\w+)(\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\))\s*\{\s*$' ++mlir_prototype_pattern = r'^\s*func\.func\s+(?:public\s+)?(@\w+)(\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\))\s*\{\s*$' + ptx_prototype_pattern = r"\.(?:visible|extern)\s+\.(?:entry|func)\s+(\w+)\s*\(([^)]*)\)" + prototype_pattern = { + "ttir": mlir_prototype_pattern, +diff --git a/test/Analysis/test-alias.mlir b/test/Analysis/test-alias.mlir +index b3d5673f85..bb21615e68 100644 +--- a/test/Analysis/test-alias.mlir ++++ b/test/Analysis/test-alias.mlir +@@ -11,7 +11,7 @@ + + // CHECK-LABEL: matmul_loop + // There shouldn't be any aliasing with the dot op encoding. +-func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + %a_ptr_init = tt.broadcast %A : (!tt.ptr) -> tensor<128x32x!tt.ptr, #AL> + %b_ptr_init = tt.broadcast %B : (!tt.ptr) -> tensor<32x128x!tt.ptr, #BL> + %a_mask = arith.constant dense : tensor<128x32xi1, #AL> +@@ -36,7 +36,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B + } + + // CHECK-LABEL: alloc +-func @alloc(%A : !tt.ptr) { ++func.func @alloc(%A : !tt.ptr) { + // CHECK: %cst -> %cst + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #AL> +@@ -46,7 +46,7 @@ func @alloc(%A : !tt.ptr) { + } + + // CHECK-LABEL: convert +-func @convert(%A : !tt.ptr) { ++func.func @convert(%A : !tt.ptr) { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> + // CHECK: %0 -> %0 + %cst1 = triton_gpu.convert_layout %cst0 : (tensor<16x16xf16, #AL>) -> tensor<16x16xf16, #A_SHARED> +@@ -54,7 +54,7 @@ func @convert(%A : !tt.ptr) { + } + + // CHECK-LABEL: trans +-func @trans(%A : !tt.ptr) { ++func.func @trans(%A : !tt.ptr) { + // CHECK: %cst -> %cst + %tensor = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #A_SHARED> + // CHECK: %0 -> %cst +@@ -63,7 +63,7 @@ func @trans(%A : !tt.ptr) { + } + + // CHECK-LABEL: insert_slice_async +-func @insert_slice_async(%A : !tt.ptr, %i1 : i1) { ++func.func @insert_slice_async(%A : !tt.ptr, %i1 : i1) { + %a_ptr = tt.broadcast %A : (!tt.ptr) -> tensor<16x16x!tt.ptr, #AL> + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> +@@ -76,7 +76,7 @@ func @insert_slice_async(%A : !tt.ptr, %i1 : i1) { + } + + // CHECK-LABEL: insert_slice +-func @insert_slice(%A : !tt.ptr, %i1 : i1) { ++func.func @insert_slice(%A : !tt.ptr, %i1 : i1) { + %a_ptr = tt.broadcast %A : (!tt.ptr) -> tensor<16x16x!tt.ptr, #AL> + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> +@@ -90,7 +90,7 @@ func @insert_slice(%A : !tt.ptr, %i1 : i1) { + } + + // CHECK-LABEL: extract_slice +-func @extract_slice(%A : !tt.ptr) { ++func.func @extract_slice(%A : !tt.ptr) { + // CHECK: %cst -> %cst + %cst0 = arith.constant dense<0.000000e+00> : tensor<1x16x16xf16, #A_SHARED> + %index = arith.constant 0 : index +@@ -100,7 +100,7 @@ func @extract_slice(%A : !tt.ptr) { + } + + // CHECK-LABEL: if_cat +-func @if_cat(%i1 : i1) { ++func.func @if_cat(%i1 : i1) { + // CHECK: %cst -> %cst + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + // CHECK: %cst_0 -> %cst_0 +@@ -119,7 +119,7 @@ func @if_cat(%i1 : i1) { + } + + // CHECK-LABEL: if_alias +-func @if_alias(%i1 : i1) { ++func.func @if_alias(%i1 : i1) { + // CHECK: %cst -> %cst + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + // CHECK-NEXT: %cst_0 -> %cst_0 +@@ -134,7 +134,7 @@ func @if_alias(%i1 : i1) { + } + + // CHECK-LABEL: for +-func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + // CHECK: %cst -> %cst + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: %cst_0 -> %cst_0 +@@ -154,7 +154,7 @@ func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.p + } + + // CHECK-LABEL: for_if +-func @for_if(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { ++func.func @for_if(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { + // CHECK: %cst -> %cst + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: %cst_0 -> %cst_0 +@@ -180,7 +180,7 @@ func @for_if(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !t + } + + // CHECK-LABEL: for_if_for +-func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { ++func.func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { + // CHECK: %cst -> %cst + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: %cst_0 -> %cst_0 +diff --git a/test/Analysis/test-alignment.mlir b/test/Analysis/test-alignment.mlir +index 0ab34c7a78..af8ea6f856 100644 +--- a/test/Analysis/test-alignment.mlir ++++ b/test/Analysis/test-alignment.mlir +@@ -1,288 +1,288 @@ +-// RUN: triton-opt %s -test-print-alignment -split-input-file 2>&1 | FileCheck %s ++// RUN: triton-opt %s -test-print-alignment -split-input-file -o %t 2>&1 | FileCheck %s + +-// CHECK-LABEL: cast +-func @cast() { +- // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [1] ++// CHECK-LABEL: @cast ++func.func @cast() { ++ // CHECK: contiguity = [1], divisibility = [1], constancy = [1], constant_value = 1 + %cst = arith.constant 1 : i32 +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [1] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = 1 + %0 = arith.extsi %cst : i32 to i64 +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 + %cst_tensor = arith.constant dense<1> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 + %1 = tt.bitcast %cst_tensor : tensor<128xi32> -> tensor<128xi64> + return + } + + // ----- + +-// CHECK-LABEL: add +-func @add() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++// CHECK-LABEL: @add ++func.func @add() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 + %1 = arith.constant dense<1> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [1], constancy = [1], constant_value = + %2 = arith.addi %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [127] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 127 + %3 = arith.constant dense<127> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 + %4 = arith.addi %1, %3 : tensor<128xi32> + return + } + + // ----- + +-// CHECK-LABEL: sub +-func @sub() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++// CHECK-LABEL: @sub ++func.func @sub() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 + %1 = arith.constant dense<1> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [1], constancy = [1], constant_value = + %2 = arith.subi %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [129] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 129 + %3 = arith.constant dense<129> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 + %4 = arith.subi %3, %1 : tensor<128xi32> + return + } + + // ----- + +-// CHECK-LABEL: mul +-func @mul() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++// CHECK-LABEL: @mul ++func.func @mul() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 + %1 = arith.constant dense<1> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %2 = arith.muli %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 + %3 = arith.constant dense<128> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 + %4 = arith.muli %3, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [2] ; Constancy: [128] ; ConstantValue: [2] ++ // CHECK-NEXT: contiguity = [1], divisibility = [2], constancy = [128], constant_value = 2 + %5 = arith.constant dense<2> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [256] ; Constancy: [128] ; ConstantValue: [256] ++ // CHECK-NEXT: contiguity = [1], divisibility = [256], constancy = [128], constant_value = 256 + %6 = arith.muli %4, %5 : tensor<128xi32> + return + } + + // ----- + +-// CHECK-LABEL: div +-func @div() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++// CHECK-LABEL: @div ++func.func @div() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 + %1 = arith.constant dense<1> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %2 = arith.divsi %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %3 = arith.divui %1, %0 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] ++ // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 + %4 = arith.constant dense<64> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16777216] ; Constancy: [64] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [16777216], constancy = [64], constant_value = + %5 = arith.divsi %0, %4 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %6 = arith.divsi %4, %0 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] ++ // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 + %7 = arith.divsi %4, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [2] ; Constancy: [128] ; ConstantValue: [66] ++ // CHECK-NEXT: contiguity = [1], divisibility = [2], constancy = [128], constant_value = 66 + %8 = arith.constant dense<66> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [2] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [2], constant_value = + %9 = arith.divui %0, %8 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [8192] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [8192], constancy = [1], constant_value = + %10 = tt.make_range {end = 8320 : i32, start = 8192 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [64] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [64], constant_value = + %11 = arith.divsi %10, %4 : tensor<128xi32> +- return ++ return + } + + // ----- + +-// CHECK-LABEL: rem +-func @rem() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++// CHECK-LABEL: @rem ++func.func @rem() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [1] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 1 + %1 = arith.constant dense<1> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] ++ // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 + %2 = arith.remsi %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %3 = arith.remui %1, %0 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] ++ // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 + %4 = arith.constant dense<64> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [64] ; Divisibility: [64] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [64], divisibility = [64], constancy = [1], constant_value = + %5 = arith.remsi %0, %4 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [1], constant_value = + %6 = arith.remsi %4, %0 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [2] ; Constancy: [128] ; ConstantValue: [66] ++ // CHECK-NEXT: contiguity = [1], divisibility = [2], constancy = [128], constant_value = 66 + %7 = arith.constant dense<66> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [2] ; Divisibility: [2] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [2], divisibility = [2], constancy = [1], constant_value = + %8 = arith.remui %0, %7 : tensor<128xi32> +- return ++ return + } + + // ----- + +-// CHECK-LABEL: broadcast +-func @broadcast() { +- // CHECK: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] ++// CHECK-LABEL: @broadcast ++func.func @broadcast() { ++ // CHECK: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 + %0 = arith.constant dense<64> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [64, 1] ; Constancy: [128, 1] ; ConstantValue: [64] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [64, 1], constancy = [128, 1], constant_value = 64 + %1 = tt.expand_dims %0 {axis = 1 : i32} : (tensor<128xi32>) -> tensor<128x1xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [64, 1] ; Constancy: [128, 128] ; ConstantValue: [64] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [64, 1], constancy = [128, 128], constant_value = 64 + %2 = tt.broadcast %1 : (tensor<128x1xi32>) -> tensor<128x128xi32> + return + } + + // ----- + +-// CHECK-LABEL: splat +-func @splat(%arg0: !tt.ptr {tt.divisibility = 16 : i32}) { +- // CHECK: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 128] ; ConstantValue: [None] ++// CHECK-LABEL: @splat ++func.func @splat(%arg0: !tt.ptr {tt.divisibility = 16 : i32}) { ++ // CHECK: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 128], constant_value = + %0 = tt.splat %arg0 : (!tt.ptr) -> tensor<128x128x!tt.ptr> + return + } + + // ----- + +-// CHECK-LABEL: cmp +-func @cmp() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++// CHECK-LABEL: @cmp ++func.func @cmp() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] ++ // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 + %1 = arith.constant dense<0> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = + %2 = arith.cmpi eq, %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = + %3 = arith.cmpi slt, %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %4 = arith.cmpi sle, %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = + %5 = arith.cmpi sge, %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8] ++ // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8 + %6 = arith.constant dense<8> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = + %7 = arith.cmpi sgt, %0, %6 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [0] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = 0 + %8 = arith.cmpi sgt, %1, %6 : tensor<128xi32> + return + } + + // ----- + +-// CHECK-LABEL: logic +-func @logic() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++// CHECK-LABEL: @logic ++func.func @logic() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [64] ; Constancy: [128] ; ConstantValue: [64] ++ // CHECK-NEXT: contiguity = [1], divisibility = [64], constancy = [128], constant_value = 64 + %1 = arith.constant dense<64> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16777216] ; Constancy: [64] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [16777216], constancy = [64], constant_value = + %2 = arith.divsi %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8] ++ // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8 + %3 = arith.constant dense<8> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [134217728] ; Constancy: [8] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [134217728], constancy = [8], constant_value = + %4 = arith.divsi %0, %3 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %5 = arith.andi %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %6 = arith.ori %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %7 = arith.xori %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = + %8 = arith.andi %2, %4 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = + %9 = arith.ori %2, %4 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [8] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [8], constant_value = + %10 = arith.xori %2, %4 : tensor<128xi32> + return + } + + // ----- + +-// CHECK-LABEL: select +-func @select() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++// CHECK-LABEL: @select ++func.func @select() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] ++ // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 + %1 = arith.constant dense<0> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = + %2 = arith.cmpi eq, %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = + %3 = arith.cmpi slt, %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [1] ; ConstantValue: [0] ++ // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [1], constant_value = 0 + %4 = arith.constant 0 : i1 +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] ++ // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 + %7 = tt.splat %4 : (i1) -> tensor<128xi1> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [128] ; ConstantValue: [0] +- %5 = select %4, %3, %7 : tensor<128xi1> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [128], constant_value = 0 ++ %5 = arith.select %4, %3, %7 : tensor<128xi1> ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [128], constant_value = + %8 = "triton_gpu.select"(%7, %3, %2) : (tensor<128xi1>, tensor<128xi1>, tensor<128xi1>) -> tensor<128xi1> + return + } + + // ----- + +-func @shift() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++func.func @shift() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8] ++ // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8 + %1 = arith.constant dense<8> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4] ; Constancy: [128] ; ConstantValue: [4] ++ // CHECK-NEXT: contiguity = [1], divisibility = [4], constancy = [128], constant_value = 4 + %2 = arith.constant dense<4> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [274877906944] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [274877906944], constancy = [1], constant_value = + %3 = arith.shli %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [67108864] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [67108864], constancy = [1], constant_value = + %4 = arith.shrsi %0, %2 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [128] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = 128 + %5 = arith.shli %1, %2 : tensor<128xi32> + return + } + + // ----- + +-func @max_min() { +- // CHECK: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++func.func @max_min() { ++ // CHECK: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [64] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [64], constancy = [1], constant_value = + %1 = tt.make_range {end = 192 : i32, start = 64 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %2 = arith.maxsi %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %3 = arith.minsi %0, %1 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [8] ; Constancy: [128] ; ConstantValue: [8] ++ // CHECK-NEXT: contiguity = [1], divisibility = [8], constancy = [128], constant_value = 8 + %4 = arith.constant dense<8> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4] ; Constancy: [128] ; ConstantValue: [4] ++ // CHECK-NEXT: contiguity = [1], divisibility = [4], constancy = [128], constant_value = 4 + %5 = arith.constant dense<4> : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [8] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = 8 + %6 = arith.maxsi %4, %5 : tensor<128xi32> + return + } + + // ----- + +-// CHECK-LABEL: for +-func @for() { +- // CHECK: Contiguity: [1, 1] ; Divisibility: [4611686018427387904, 4611686018427387904] ; Constancy: [128, 32] ; ConstantValue: [0] ++// CHECK-LABEL: @for ++func.func @for() { ++ // CHECK: contiguity = [1, 1], divisibility = [4611686018427387904, 4611686018427387904], constancy = [128, 32], constant_value = 0 + %a_init = arith.constant dense<0> : tensor<128x32xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 32] ; ConstantValue: [1] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 32], constant_value = 1 + %b_init = arith.constant dense<1> : tensor<128x32xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [4, 4] ; Constancy: [128, 32] ; ConstantValue: [4] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [4, 4], constancy = [128, 32], constant_value = 4 + %c_init = arith.constant dense<4> : tensor<128x32xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [128] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [1], constant_value = 128 + %ub = arith.constant 128 : index +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [4611686018427387904] ; Constancy: [1] ; ConstantValue: [0] ++ // CHECK-NEXT: contiguity = [1], divisibility = [4611686018427387904], constancy = [1], constant_value = 0 + %lb = arith.constant 0 : index +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [16] ++ // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [1], constant_value = 16 + %step = arith.constant 16 : index + %a, %b, %c = scf.for %iv = %lb to %ub step %step iter_args(%a = %a_init, %b = %b_init, %c = %c_init) -> (tensor<128x32xi32>, tensor<128x32xi32>, tensor<128x32xi32>) { +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [1], constant_value = + %t = arith.index_cast %iv : index to i32 +- // CHECK: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 32] ; ConstantValue: [None] +- // CHECK: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 32] ; ConstantValue: [None] +- // CHECK: Contiguity: [1, 1] ; Divisibility: [4, 4] ; Constancy: [128, 32] ; ConstantValue: [4] ++ // CHECK: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 32], constant_value = ++ // CHECK: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 32], constant_value = ++ // CHECK: contiguity = [1, 1], divisibility = [4, 4], constancy = [128, 32], constant_value = 4 + scf.yield %b, %a, %c : tensor<128x32xi32>, tensor<128x32xi32>, tensor<128x32xi32> + } + return +@@ -290,53 +290,53 @@ func @for() { + + // ----- + +-// CHECK-LABEL: permute_2d +-func @permute_2d(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32}) { +- // CHECK: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [128, 128] ; ConstantValue: [1] ++// CHECK-LABEL: @permute_2d ++func.func @permute_2d(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32}) { ++ // CHECK: contiguity = [1, 1], divisibility = [1, 1], constancy = [128, 128], constant_value = 1 + %cst = arith.constant dense : tensor<128x128xi1> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [1, 1], constancy = [1, 1], constant_value = + %cst_0 = arith.constant dense<0.000000e+00> : tensor<128x128xf32> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %0 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %1 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [1073741824, 1] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128, 1], divisibility = [1073741824, 1], constancy = [1, 1], constant_value = + %2 = tt.expand_dims %0 {axis = 1 : i32} : (tensor<128xi32>) -> tensor<128x1xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 1], constant_value = + %3 = tt.splat %arg1 : (i32) -> tensor<128x1xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [17179869184, 16] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [17179869184, 16], constancy = [1, 1], constant_value = + %4 = arith.muli %2, %3 : tensor<128x1xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 1], constant_value = + %5 = tt.splat %arg0 : (!tt.ptr) -> tensor<128x1x!tt.ptr> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [1, 1], constant_value = + %6 = tt.addptr %5, %4 : tensor<128x1x!tt.ptr>, tensor<128x1xi32> +- // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 1073741824] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 1073741824], constancy = [1, 1], constant_value = + %7 = tt.expand_dims %1 {axis = 0 : i32}: (tensor<128xi32>) -> tensor<1x128xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [1, 128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [1, 128], constant_value = + %8 = tt.broadcast %6 : (tensor<128x1x!tt.ptr>) -> tensor<128x128x!tt.ptr> +- // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 1073741824] ; Constancy: [128, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 1073741824], constancy = [128, 1], constant_value = + %9 = tt.broadcast %7 : (tensor<1x128xi32>) -> tensor<128x128xi32> +- // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 16] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 16], constancy = [1, 1], constant_value = + %10 = tt.addptr %8, %9 : tensor<128x128x!tt.ptr>, tensor<128x128xi32> +- // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [1073741824, 1] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128, 1], divisibility = [1073741824, 1], constancy = [1, 1], constant_value = + %11 = tt.expand_dims %0 {axis = 1 : i32}: (tensor<128xi32>) -> tensor<128x1xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [128, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [128, 1], constant_value = + %12 = tt.splat %arg2 : (!tt.ptr) -> tensor<128x1x!tt.ptr> +- // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [16, 1] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128, 1], divisibility = [16, 1], constancy = [1, 1], constant_value = + %13 = tt.addptr %12, %11 : tensor<128x1x!tt.ptr>, tensor<128x1xi32> +- // CHECK-NEXT: Contiguity: [1, 128] ; Divisibility: [1, 1073741824] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 128], divisibility = [1, 1073741824], constancy = [1, 1], constant_value = + %14 = tt.expand_dims %1 {axis = 0 : i32} : (tensor<128xi32>) -> tensor<1x128xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 16] ; Constancy: [1, 128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 16], constancy = [1, 128], constant_value = + %15 = tt.splat %arg3 : (i32) -> tensor<1x128xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 17179869184] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 17179869184], constancy = [1, 1], constant_value = + %16 = arith.muli %14, %15 : tensor<1x128xi32> +- // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [16, 1] ; Constancy: [1, 128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128, 1], divisibility = [16, 1], constancy = [1, 128], constant_value = + %17 = tt.broadcast %13 : (tensor<128x1x!tt.ptr>) -> tensor<128x128x!tt.ptr> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [16, 17179869184] ; Constancy: [128, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [16, 17179869184], constancy = [128, 1], constant_value = + %18 = tt.broadcast %16 : (tensor<1x128xi32>) -> tensor<128x128xi32> +- // CHECK-NEXT: Contiguity: [128, 1] ; Divisibility: [16, 1] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128, 1], divisibility = [16, 1], constancy = [1, 1], constant_value = + %19 = tt.addptr %17, %18 : tensor<128x128x!tt.ptr>, tensor<128x128xi32> +- // CHECK-NEXT: Contiguity: [1, 1] ; Divisibility: [1, 1] ; Constancy: [1, 1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1, 1], divisibility = [1, 1], constancy = [1, 1], constant_value = + %20 = tt.load %10, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<128x128xf32> + tt.store %19, %20, %cst : tensor<128x128xf32> + return +@@ -347,29 +347,29 @@ func @permute_2d(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: i32 {t + module { + + // This is a tiny test for verifying StoreOp-related alignment, It simply store a constant to a buffer. +-// CHECK-LABEL: store_constant_align +-func @store_constant_align(%addr: !tt.ptr {tt.divisibility = 16 : i32}, %n: i32 {tt.divisibility = 16 : i32}) { +- // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++// CHECK-LABEL: @store_constant_align ++func.func @store_constant_align(%addr: !tt.ptr {tt.divisibility = 16 : i32}, %n: i32 {tt.divisibility = 16 : i32}) { ++ // CHECK: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %pid = tt.get_program_id {axis = 0 : i32} : i32 +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [128] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [1], constant_value = 128 + %c128_i32 = arith.constant 128 : i32 +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [1], constant_value = + %1 = arith.muli %pid, %c128_i32 : i32 +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [1073741824] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [1073741824], constancy = [1], constant_value = + %2 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [128] ; Constancy: [128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [128], constancy = [128], constant_value = + %3 = tt.splat %1 : (i32) -> tensor<128xi32> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [128] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [128], constancy = [1], constant_value = + %4 = arith.addi %3, %2 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [128], constant_value = + %5 = tt.splat %addr : (!tt.ptr) -> tensor<128x!tt.ptr> +- // CHECK-NEXT: Contiguity: [128] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [128], divisibility = [16], constancy = [1], constant_value = + %6 = tt.addptr %5, %4 : tensor<128x!tt.ptr>, tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [16] ; Constancy: [128] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [16], constancy = [128], constant_value = + %9 = tt.splat %n : (i32) -> tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [16] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [16], constant_value = + %mask = arith.cmpi slt, %4, %9 : tensor<128xi32> +- // CHECK-NEXT: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ++ // CHECK-NEXT: contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %cst = arith.constant dense<0.0> : tensor<128xf32> + tt.store %5, %cst, %mask : tensor<128xf32> + return +@@ -381,8 +381,8 @@ func @store_constant_align(%addr: !tt.ptr {tt.divisibility = 16 : i32}, %n: + + // This IR is dumped from vecadd test. + // Note, the hint {tt.divisibility = 16 : i32} for %n_elements affects the alignment of mask. +-// CHECK-LABEL: vecadd_mask_align_16 +-func @vecadd_mask_align_16(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %n_elements: i32 {tt.divisibility = 16 : i32}) { ++// CHECK-LABEL: @vecadd_mask_align_16 ++func.func @vecadd_mask_align_16(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %n_elements: i32 {tt.divisibility = 16 : i32}) { + %c64_i32 = arith.constant 64 : i32 + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = arith.muli %0, %c64_i32 : i32 +@@ -394,13 +394,13 @@ func @vecadd_mask_align_16(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %ar + %7 = tt.splat %arg1 : (!tt.ptr) -> tensor<64x!tt.ptr> + %8 = tt.addptr %7, %4 : tensor<64x!tt.ptr>, tensor<64xi32> + %9 = tt.splat %n_elements : (i32) -> tensor<64xi32> +- // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [16] ; ConstantValue: [None] ( %{{.*}} = arith.cmpi slt, %{{.*}}, %{{.*}} : tensor<64xi32> ) ++ // CHECK: arith.cmpi slt, %{{.*}} => contiguity = [1], divisibility = [1], constancy = [16], constant_value = + %mask = arith.cmpi slt, %4, %9 : tensor<64xi32> + %11 = tt.load %6, %mask {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32> + %12 = tt.load %8, %mask {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32> + %13 = arith.addf %11, %12 : tensor<64xf32> + %14 = tt.splat %arg2 : (!tt.ptr) -> tensor<64x!tt.ptr> +- // CHECK: Contiguity: [64] ; Divisibility: [16] ; Constancy: [1] ; ConstantValue: [None] ( %{{.*}} = tt.addptr %{{.*}}, %{{.*}} : tensor<64x!tt.ptr>, tensor<64xi32> ) ++ // CHECK: tt.addptr %{{.*}} => contiguity = [64], divisibility = [16], constancy = [1], constant_value = + %15 = tt.addptr %14, %4 : tensor<64x!tt.ptr>, tensor<64xi32> + tt.store %15, %13, %mask : tensor<64xf32> + return +@@ -410,8 +410,8 @@ func @vecadd_mask_align_16(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %ar + + // This IR is dumped from vecadd test. + // Note, there is no divisibility hint for %n_elements, Triton should assume its divisibility to be 1 by default. +-// CHECK-LABEL: vecadd_mask_align_1 +-func @vecadd_mask_align_1(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %n_elements: i32) { ++// CHECK-LABEL: @vecadd_mask_align_1 ++func.func @vecadd_mask_align_1(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %n_elements: i32) { + %c64_i32 = arith.constant 64 : i32 + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = arith.muli %0, %c64_i32 : i32 +@@ -423,7 +423,7 @@ func @vecadd_mask_align_1(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg + %7 = tt.splat %arg1 : (!tt.ptr) -> tensor<64x!tt.ptr> + %8 = tt.addptr %7, %4 : tensor<64x!tt.ptr>, tensor<64xi32> + %9 = tt.splat %n_elements : (i32) -> tensor<64xi32> +- // CHECK: Contiguity: [1] ; Divisibility: [1] ; Constancy: [1] ; ConstantValue: [None] ( %{{.*}} = arith.cmpi slt, %{{.*}}, %{{.*}} : tensor<64xi32> ) ++ // CHECK: arith.cmpi slt, %{{.*}} => contiguity = [1], divisibility = [1], constancy = [1], constant_value = + %10 = arith.cmpi slt, %4, %9 : tensor<64xi32> + %11 = tt.load %6, %10 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32> + %12 = tt.load %8, %10 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64xf32> +diff --git a/test/Analysis/test-allocation.mlir b/test/Analysis/test-allocation.mlir +index efb00c404d..f79222aa7b 100644 +--- a/test/Analysis/test-allocation.mlir ++++ b/test/Analysis/test-allocation.mlir +@@ -13,7 +13,7 @@ + module attributes {"triton_gpu.num-warps" = 4 : i32} { + + // CHECK-LABEL: matmul_loop +-func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + %a_ptr_init = tt.broadcast %A : (!tt.ptr) -> tensor<128x32x!tt.ptr, #AL> + %b_ptr_init = tt.broadcast %B : (!tt.ptr) -> tensor<32x128x!tt.ptr, #BL> + +@@ -46,7 +46,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B + + // Shared memory is available after a tensor's liveness range ends + // CHECK-LABEL: reusable +-func @reusable(%A : !tt.ptr) { ++func.func @reusable(%A : !tt.ptr) { + %cst1 = arith.constant dense : tensor<128x32xi1, #AL> + %cst2 = arith.constant dense<0.000000e+00> : tensor<128x32xf16, #AL> + %cst3 = arith.constant dense : tensor<32x128xi1, #AL> +@@ -78,7 +78,7 @@ func @reusable(%A : !tt.ptr) { + // %cst1->%cst4 + // %cst3->%g->%h->%i + // CHECK-LABEL: preallocate +-func @preallocate(%A : !tt.ptr) { ++func.func @preallocate(%A : !tt.ptr) { + // CHECK: offset = 0, size = 512 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + // CHECK-NEXT: offset = 1024, size = 512 +@@ -113,7 +113,7 @@ func @preallocate(%A : !tt.ptr) { + + // Unused tensors are immediately released + // CHECK-LABEL: unused +-func @unused(%A : !tt.ptr) { ++func.func @unused(%A : !tt.ptr) { + // CHECK: offset = 0, size = 1024 + %cst0 = arith.constant dense<0.000000e+00> : tensor<32x16xf16, #A_SHARED> + // CHECK-NEXT: offset = 0, size = 512 +@@ -128,7 +128,7 @@ func @unused(%A : !tt.ptr) { + + // cst0 is alive through the entire function, it cannot be released before the end of the function + // CHECK-LABEL: longlive +-func @longlive(%A : !tt.ptr) { ++func.func @longlive(%A : !tt.ptr) { + // CHECK: offset = 0, size = 512 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + // CHECK-NEXT: offset = 512, size = 512 +@@ -156,7 +156,7 @@ func @longlive(%A : !tt.ptr) { + } + + // CHECK-LABEL: alloc +-func @alloc(%A : !tt.ptr) { ++func.func @alloc(%A : !tt.ptr) { + // CHECK: offset = 0, size = 512 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #AL> +@@ -167,7 +167,7 @@ func @alloc(%A : !tt.ptr) { + } + + // CHECK-LABEL: scratch +-func @scratch() { ++func.func @scratch() { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> + // CHECK: scratch offset = 0, size = 512 + %b = tt.reduce %cst0 {redOp = 1 : i32, axis = 0 : i32} : tensor<16x16xf16, #AL> -> tensor<16xf16, #sliceAd0> +@@ -176,7 +176,7 @@ func @scratch() { + } + + // CHECK-LABEL: trans +-func @trans(%A : !tt.ptr) { ++func.func @trans(%A : !tt.ptr) { + // CHECK: offset = 0, size = 1024 + %tensor = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #A_SHARED> + %b = tt.trans %tensor : (tensor<16x32xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED_T> +@@ -184,7 +184,7 @@ func @trans(%A : !tt.ptr) { + } + + // CHECK-LABEL: insert_slice_async +-func @insert_slice_async(%A : !tt.ptr, %i1 : i1) { ++func.func @insert_slice_async(%A : !tt.ptr, %i1 : i1) { + %a_ptr = tt.broadcast %A : (!tt.ptr) -> tensor<16x16x!tt.ptr, #AL> + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> +@@ -197,7 +197,7 @@ func @insert_slice_async(%A : !tt.ptr, %i1 : i1) { + } + + // CHECK-LABEL: extract_slice +-func @extract_slice(%A : !tt.ptr) { ++func.func @extract_slice(%A : !tt.ptr) { + // CHECK: offset = 0, size = 512 + %cst0 = arith.constant dense<0.000000e+00> : tensor<1x16x16xf16, #A_SHARED> + %index = arith.constant 0 : index +@@ -209,7 +209,7 @@ func @extract_slice(%A : !tt.ptr) { + // B0 -> (B1) -> B0 + // Memory used by B1 can be reused by B0. + // CHECK-LABEL: if +-func @if(%i1 : i1) { ++func.func @if(%i1 : i1) { + // CHECK: offset = 0, size = 512 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + // CHECK-NEXT: offset = 512, size = 512 +@@ -233,7 +233,7 @@ func @if(%i1 : i1) { + // B0 -> (B1) -> (B2) -> B0 + // Memory used by B0 cannot be reused by B1 or B2. + // CHECK-LABEL: if_else +-func @if_else(%i1 : i1) { ++func.func @if_else(%i1 : i1) { + // CHECK: offset = 0, size = 512 + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + // CHECK-NEXT: offset = 512, size = 512 +@@ -260,7 +260,7 @@ func @if_else(%i1 : i1) { + // Block arguments and yields are memory aliases that do not trigger a new + // allocation. + // CHECK-LABEL: for +-func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + // CHECK: offset = 0, size = 8192 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: offset = 8192, size = 8192 +@@ -275,7 +275,7 @@ func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.p + } + + // CHECK-LABEL: for_if_slice +-func @for_if_slice(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { ++func.func @for_if_slice(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { + // CHECK: offset = 0, size = 8192 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: offset = 8192, size = 8192 +@@ -296,7 +296,7 @@ func @for_if_slice(%lb : index, %ub : index, %step : index, %A : !tt.ptr, % + + // c0 cannot be released in the loop + // CHECK-LABEL: for_use_ancestor +-func @for_use_ancestor(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { ++func.func @for_use_ancestor(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { + // CHECK: offset = 0, size = 8192 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: offset = 8192, size = 8192 +@@ -316,7 +316,7 @@ func @for_use_ancestor(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { ++func.func @for_if_for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr, %i1 : i1) { + // CHECK: offset = 0, size = 8192 + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: offset = 8192, size = 8192 +diff --git a/test/Analysis/test-membar.mlir b/test/Analysis/test-membar.mlir +index 7199e5f53d..17880b2094 100644 +--- a/test/Analysis/test-membar.mlir ++++ b/test/Analysis/test-membar.mlir +@@ -14,7 +14,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + + // CHECK-LABEL: matmul_loop + // There shouldn't be any membar with the dot op encoding. +-func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + %a_ptr_init = tt.broadcast %A : (!tt.ptr) -> tensor<128x32x!tt.ptr, #AL> + %b_ptr_init = tt.broadcast %B : (!tt.ptr) -> tensor<32x128x!tt.ptr, #BL> + +@@ -42,7 +42,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B + } + + // CHECK-LABEL: raw_single_block +-func @raw_single_block(%A : !tt.ptr) { ++func.func @raw_single_block(%A : !tt.ptr) { + %cst1 = arith.constant dense : tensor<128x32xi1, #AL> + %cst2 = arith.constant dense<0.000000e+00> : tensor<128x32xf16, #AL> + %a_ptr = tt.broadcast %A : (!tt.ptr) -> tensor<128x32x!tt.ptr, #AL> +@@ -54,7 +54,7 @@ func @raw_single_block(%A : !tt.ptr) { + } + + // CHECK-LABEL: war_single_block +-func @war_single_block(%A : !tt.ptr) { ++func.func @war_single_block(%A : !tt.ptr) { + %cst1 = arith.constant dense : tensor<128x32xi1, #AL> + %cst2 = arith.constant dense<0.000000e+00> : tensor<128x32xf16, #AL> + %a_ptr = tt.broadcast %A : (!tt.ptr) -> tensor<128x32x!tt.ptr, #AL> +@@ -70,7 +70,7 @@ func @war_single_block(%A : !tt.ptr) { + } + + // CHECK-LABEL: scratch +-func @scratch() { ++func.func @scratch() { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + // CHECK: Membar 1 + %a = tt.cat %cst0, %cst0 {axis = 0} : (tensor<16x16xf16, #A_SHARED>, tensor<16x16xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED> +@@ -81,7 +81,7 @@ func @scratch() { + } + + // CHECK-LABEL: async_wait +-func @async_wait() { ++func.func @async_wait() { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + // CHECK: Membar 1 + %a = tt.cat %cst0, %cst0 {axis = 0} : (tensor<16x16xf16, #A_SHARED>, tensor<16x16xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED> +@@ -92,7 +92,7 @@ func @async_wait() { + } + + // CHECK-LABEL: alloc +-func @alloc() { ++func.func @alloc() { + %cst0 = triton_gpu.alloc_tensor : tensor<16x16xf16, #A_SHARED> + %a = tt.cat %cst0, %cst0 {axis = 0} : (tensor<16x16xf16, #A_SHARED>, tensor<16x16xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED> + // CHECK: Membar 2 +@@ -101,7 +101,7 @@ func @alloc() { + } + + // CHECK-LABEL: extract_slice +-func @extract_slice() { ++func.func @extract_slice() { + %cst0 = arith.constant dense<0.000000e+00> : tensor<1x16x16xf16, #A_SHARED> + %index = arith.constant 0 : index + %cst1 = tensor.extract_slice %cst0[%index, 0, 0][1, 16, 16][1, 1, 1] : tensor<1x16x16xf16, #A_SHARED> to tensor<16x16xf16, #A_SHARED> +@@ -113,14 +113,14 @@ func @extract_slice() { + } + + // CHECK-LABEL: trans +-func @trans() { ++func.func @trans() { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x32xf16, #A_SHARED> + %b = tt.trans %cst0 : (tensor<16x32xf16, #A_SHARED>) -> tensor<32x16xf16, #A_SHARED_T> + return + } + + // CHECK-LABEL: insert_slice_async +-func @insert_slice_async(%A : !tt.ptr, %i1 : i1) { ++func.func @insert_slice_async(%A : !tt.ptr, %i1 : i1) { + %a_ptr = tt.broadcast %A : (!tt.ptr) -> tensor<16x16x!tt.ptr, #AL> + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> +@@ -135,7 +135,7 @@ func @insert_slice_async(%A : !tt.ptr, %i1 : i1) { + } + + // CHECK-LABEL: insert_slice +-func @insert_slice(%A : !tt.ptr, %i1 : i1) { ++func.func @insert_slice(%A : !tt.ptr, %i1 : i1) { + %a_ptr = tt.broadcast %A : (!tt.ptr) -> tensor<16x16x!tt.ptr, #AL> + %mask = tt.splat %i1 : (i1) -> tensor<16x16xi1, #AL> + %other = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #AL> +@@ -153,7 +153,7 @@ func @insert_slice(%A : !tt.ptr, %i1 : i1) { + + // If branch inserted a barrier for %cst0 and %cst1, but else didn't, then the barrier should be inserted in the parent region + // CHECK-LABEL: multi_blocks +-func @multi_blocks(%i1 : i1) { ++func.func @multi_blocks(%i1 : i1) { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + scf.if %i1 { +@@ -174,7 +174,7 @@ func @multi_blocks(%i1 : i1) { + + // Both branches inserted a barrier for %cst0 and %cst1, then the barrier doesn't need to be inserted in the parent region + // CHECK-LABEL: multi_blocks_join_barrier +-func @multi_blocks_join_barrier(%i1 : i1) { ++func.func @multi_blocks_join_barrier(%i1 : i1) { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + scf.if %i1 { +@@ -192,7 +192,7 @@ func @multi_blocks_join_barrier(%i1 : i1) { + + // Read yielded tensor requires a barrier + // CHECK-LABEL: multi_blocks_yield +-func @multi_blocks_yield(%i1 : i1) { ++func.func @multi_blocks_yield(%i1 : i1) { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + %a = scf.if %i1 -> (tensor<32x16xf16, #A_SHARED>) { +@@ -212,7 +212,7 @@ func @multi_blocks_yield(%i1 : i1) { + + // Conservatively add a barrier as if the branch (%i1) is never taken + // CHECK-LABEL: multi_blocks_noelse +-func @multi_blocks_noelse(%i1 : i1) { ++func.func @multi_blocks_noelse(%i1 : i1) { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + scf.if %i1 { +@@ -226,7 +226,7 @@ func @multi_blocks_noelse(%i1 : i1) { + + // Conservatively add a barrier as if the branch (%i2) is never taken + // CHECK-LABEL: multi_blocks_nested_scf +-func @multi_blocks_nested_scf(%i1 : i1, %i2 : i1) { ++func.func @multi_blocks_nested_scf(%i1 : i1, %i2 : i1) { + %cst0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + %cst1 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #A_SHARED> + scf.if %i1 { +@@ -247,7 +247,7 @@ func @multi_blocks_nested_scf(%i1 : i1, %i2 : i1) { + } + + // CHECK-LABEL: for +-func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + %c_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> +@@ -262,7 +262,7 @@ func @for(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.p + // Although a_shared and b_shared are synced before entering the loop, + // they are reassociated with aliases (c_shared) and thus require a barrier. + // CHECK-LABEL: for_alias +-func @for_alias(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @for_alias(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: Membar 2 +@@ -282,7 +282,7 @@ func @for_alias(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : + // Although cst2 is not an argument of scf.yield, its memory is reused by cst1. + // So we need a barrier both before and after cst1 + // CHECK-LABEL: for_reuse +-func @for_reuse(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @for_reuse(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: Membar 2 +@@ -302,7 +302,7 @@ func @for_reuse(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : + + + // CHECK-LABEL: for_reuse_nested +-func @for_reuse_nested(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @for_reuse_nested(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + %a_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + %b_shared_init = arith.constant dense<0.00e+00> : tensor<128x32xf16, #A_SHARED> + // CHECK-NEXT: Membar 2 +diff --git a/test/Conversion/triton_ops.mlir b/test/Conversion/triton_ops.mlir +index e9ee502435..0e979b148d 100644 +--- a/test/Conversion/triton_ops.mlir ++++ b/test/Conversion/triton_ops.mlir +@@ -1,6 +1,6 @@ + // RUN: triton-opt %s | FileCheck %s + +-func @cast_ops(%scalar_ptr: !tt.ptr, %scalar_f32: f32, %scalar_i64: i64) { ++func.func @cast_ops(%scalar_ptr: !tt.ptr, %scalar_f32: f32, %scalar_i64: i64) { + // scalar -> scalar + // CHECK: i64 -> !tt.ptr + %0 = tt.int_to_ptr %scalar_i64 : i64 -> !tt.ptr +@@ -35,7 +35,7 @@ func @cast_ops(%scalar_ptr: !tt.ptr, %scalar_f32: f32, %scalar_i64: i64) { + return + } + +-func @addptr_ops(%scalar_ptr: !tt.ptr, %scalar_i32: i32) { ++func.func @addptr_ops(%scalar_ptr: !tt.ptr, %scalar_i32: i32) { + // scalar -> scalar + // CHECK: !tt.ptr + %0 = tt.addptr %scalar_ptr, %scalar_i32 : !tt.ptr, i32 +@@ -54,7 +54,7 @@ func @addptr_ops(%scalar_ptr: !tt.ptr, %scalar_i32: i32) { + return + } + +-func @load_store_ops_scalar(%ptr: !tt.ptr {tt.divisibility = 16 : i32}, %mask : i1) { ++func.func @load_store_ops_scalar(%ptr: !tt.ptr {tt.divisibility = 16 : i32}, %mask : i1) { + // Test if Load/Store ops can handle scalar values + %other = arith.constant 0.0e+0 : f32 + +@@ -76,7 +76,7 @@ func @load_store_ops_scalar(%ptr: !tt.ptr {tt.divisibility = 16 : i32}, %ma + return + } + +-func @reduce_ops_infer(%ptr: !tt.ptr, %v : tensor<1x2x4xf32>) { ++func.func @reduce_ops_infer(%ptr: !tt.ptr, %v : tensor<1x2x4xf32>) { + // Test if reduce ops infer types correctly + + // CHECK: %{{.*}} = tt.reduce %{{.*}} -> tensor<2x4xf32> +@@ -101,7 +101,7 @@ func @reduce_ops_infer(%ptr: !tt.ptr, %v : tensor<1x2x4xf32>) { + return + } + +-func @dot_ops_infer(%ptr: !tt.ptr, %v : f32) { ++func.func @dot_ops_infer(%ptr: !tt.ptr, %v : f32) { + // Test if reduce ops infer types correctly + %v128x32 = tt.splat %v : (f32) -> tensor<128x32xf32> + %v32x128 = tt.splat %v : (f32) -> tensor<32x128xf32> +diff --git a/test/Conversion/triton_to_tritongpu.mlir b/test/Conversion/triton_to_tritongpu.mlir +index a160bc8815..b461ca542f 100644 +--- a/test/Conversion/triton_to_tritongpu.mlir ++++ b/test/Conversion/triton_to_tritongpu.mlir +@@ -1,6 +1,6 @@ + // RUN: triton-opt %s -split-input-file -convert-triton-to-tritongpu=num-warps=2 | FileCheck %s + +-func @ops() { ++func.func @ops() { + // CHECK: module attributes {"triton_gpu.num-warps" = 2 : i32} {{.*}} + %a = arith.constant dense<1.00e+00> : tensor<128x32xf16> + %b = arith.constant dense<2.00e+00> : tensor<32x128xf16> +@@ -11,7 +11,7 @@ func @ops() { + + // ----- + +-func @load_ops(%ptr: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @load_ops(%ptr: !tt.ptr {tt.divisibility = 16 : i32}) { + // Test if LoadOp is lowered properly (see #771) + %ptrs = tt.splat %ptr : (!tt.ptr) -> tensor<128x!tt.ptr> + %mask = arith.constant dense : tensor<128xi1> +@@ -30,7 +30,7 @@ func @load_ops(%ptr: !tt.ptr {tt.divisibility = 16 : i32}) { + + // ----- + +-func @reduce_ops(%ptr: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @reduce_ops(%ptr: !tt.ptr {tt.divisibility = 16 : i32}) { + // Test if the total number of threadsPerWarp is 32 + // Test if the total number of warps is 2 + // CHECK: #blocked0 = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [4, 8], warpsPerCTA = [1, 2], order = [0, 1]}> +diff --git a/test/Conversion/tritongpu_to_llvm.mlir b/test/Conversion/tritongpu_to_llvm.mlir +index e9e7d5a340..507b362c99 100644 +--- a/test/Conversion/tritongpu_to_llvm.mlir ++++ b/test/Conversion/tritongpu_to_llvm.mlir +@@ -4,7 +4,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK: llvm.func @test_empty_kernel(%arg0: i32, %arg1: !llvm.ptr) + // Here the 128 comes from the 4 in module attribute multiples 32 + // CHECK: attributes {nvvm.kernel = 1 : ui1, nvvm.maxntid = 128 : i32} {{.*}} +- func @test_empty_kernel(%lb : index, %A : !tt.ptr) { ++ func.func @test_empty_kernel(%lb : index, %A : !tt.ptr) { + // CHECK: llvm.return + return + } +@@ -15,7 +15,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_load +- func @basic_load(%a_ptr_init : tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) { ++ func.func @basic_load(%a_ptr_init : tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) { + // CHECK: llvm.inline_asm + // CHECK: llvm.inline_asm + %1 = tt.load %a_ptr_init, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked0> +@@ -28,7 +28,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: vectorized_load +- func @vectorized_load(%a_ptr_init : tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) { ++ func.func @vectorized_load(%a_ptr_init : tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf32, #blocked0>) { + // CHECK: llvm.inline_asm + // CHECK-SAME: ld.global.b32 + // CHECK: llvm.inline_asm +@@ -43,7 +43,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK-LABEL: vectorized_load_f16 +- func @vectorized_load_f16(%a_ptr_init: tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf16, #blocked0>) { ++ func.func @vectorized_load_f16(%a_ptr_init: tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>, %cst_0 : tensor<256xf16, #blocked0>) { + // CHECK: llvm.inline_asm + // CHECK-SAME: ld.global.b16 + // CHECK: llvm.inline_asm +@@ -59,7 +59,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [8], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: masked_load_const_other +- func @masked_load_const_other(%a_ptr_init : tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>) { ++ func.func @masked_load_const_other(%a_ptr_init : tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>) { + %cst_0 = arith.constant dense<0.000000e+00> : tensor<256xf32, #blocked0> + %1 = tt.load %a_ptr_init, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked0> + return +@@ -72,7 +72,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [8], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: masked_load_const_other_vec +- func @masked_load_const_other_vec(%a_ptr_init : tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>) { ++ func.func @masked_load_const_other_vec(%a_ptr_init : tensor<256x!tt.ptr, #blocked0>, %cst : tensor<256xi1, #blocked0>) { + %cst_0 = arith.constant dense<0.000000e+00> : tensor<256xf32, #blocked0> + %1 = tt.load %a_ptr_init, %cst, %cst_0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked0> + return +@@ -84,7 +84,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [2], order = [0]}> + module attributes {"triton_gpu.num-warps" = 2 : i32} { + // CHECK-LABEL: global_load_store_no_vec +- func @global_load_store_no_vec(%arg0: !tt.ptr {tt.divisibility = 4 : i32}, %arg1: !tt.ptr {tt.divisibility = 4 : i32}, %arg2: !tt.ptr {tt.divisibility = 4 : i32}, %arg3: i32) { ++ func.func @global_load_store_no_vec(%arg0: !tt.ptr {tt.divisibility = 4 : i32}, %arg1: !tt.ptr {tt.divisibility = 4 : i32}, %arg2: !tt.ptr {tt.divisibility = 4 : i32}, %arg3: i32) { + %c256_i32 = arith.constant 256 : i32 + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = arith.muli %0, %c256_i32 : i32 +@@ -128,7 +128,7 @@ module attributes {"triton_gpu.num-warps" = 2 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [2], order = [0]}> + module attributes {"triton_gpu.num-warps" = 2 : i32} { + // CHECK-LABEL: global_load_store_vec4 +- func @global_load_store_vec4(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32) { ++ func.func @global_load_store_vec4(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32) { + %c256_i32 = arith.constant 256 : i32 + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = arith.muli %0, %c256_i32 : i32 +@@ -165,7 +165,7 @@ module attributes {"triton_gpu.num-warps" = 2 : i32} { + #blocked = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [2], order = [0]}> + // Note, the %n_elements doesn't have a "tt.divisibility" hint, so Triton assumes it's divisibility is 1, this should effect the mask's alignment and further restrict the load/store ops' vector width to be 1. + module attributes {"triton_gpu.num-warps" = 2 : i32} { +- func @vecadd_masked_vec1(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %n_elements: i32) { ++ func.func @vecadd_masked_vec1(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %n_elements: i32) { + %c64_i32 = arith.constant 64 : i32 + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = arith.muli %0, %c64_i32 : i32 +@@ -195,7 +195,7 @@ module attributes {"triton_gpu.num-warps" = 2 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK-LABEL: global_load_store_vec2 +- func @global_load_store_vec2(%arg0: !tt.ptr {tt.divisibility = 8 : i32}, %arg1: !tt.ptr {tt.divisibility = 8 : i32}, %arg2: !tt.ptr {tt.divisibility = 8 : i32}, %arg3: i32) { ++ func.func @global_load_store_vec2(%arg0: !tt.ptr {tt.divisibility = 8 : i32}, %arg1: !tt.ptr {tt.divisibility = 8 : i32}, %arg2: !tt.ptr {tt.divisibility = 8 : i32}, %arg3: i32) { + %c256_i32 = arith.constant 256 : i32 + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = arith.muli %0, %c256_i32 : i32 +@@ -240,7 +240,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK-LABEL: global_load_store_vec8 +- func @global_load_store_vec8(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32) { ++ func.func @global_load_store_vec8(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32) { + %c256_i32 = arith.constant 256 : i32 + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = arith.muli %0, %c256_i32 : i32 +@@ -283,7 +283,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + #blocked2 = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [4, 1], order = [0, 1]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_view_broadcast +- func @basic_view_broadcast(%arg : tensor<256xf32,#blocked0>) { ++ func.func @basic_view_broadcast(%arg : tensor<256xf32,#blocked0>) { + // CHECK: llvm.mlir.undef + // CHECK: %[[T0:.*]] = llvm.extractvalue + // CHECK: %[[T1:.*]] = llvm.extractvalue +@@ -307,7 +307,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_make_range +- func @basic_make_range() { ++ func.func @basic_make_range() { + // CHECK: nvvm.read.ptx.sreg.tid.x + // CHECK: llvm.mlir.undef + // CHECK: llvm.insertvalue +@@ -322,7 +322,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_addf +- func @basic_addf(%arg0 : tensor<256xf32,#blocked0>, %arg1 : tensor<256xf32,#blocked0>) { ++ func.func @basic_addf(%arg0 : tensor<256xf32,#blocked0>, %arg1 : tensor<256xf32,#blocked0>) { + // CHECK: llvm.fadd + // CHECK: llvm.fadd + %1 = arith.addf %arg0, %arg1 : tensor<256xf32,#blocked0> +@@ -335,7 +335,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_addi +- func @basic_addi(%arg0 : tensor<256xi32,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) { ++ func.func @basic_addi(%arg0 : tensor<256xi32,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) { + // CHECK: llvm.add + // CHECK: llvm.add + %1 = arith.addi %arg0, %arg1 : tensor<256xi32,#blocked0> +@@ -347,7 +347,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_program_id +- func @basic_program_id() { ++ func.func @basic_program_id() { + // CHECK: nvvm.read.ptx.sreg.ctaid.x : i32 + %0 = tt.get_program_id {axis = 0 : i32} : i32 + return +@@ -359,7 +359,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_addptr +- func @basic_addptr(%arg0 : tensor<256x!tt.ptr,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) { ++ func.func @basic_addptr(%arg0 : tensor<256x!tt.ptr,#blocked0>, %arg1 : tensor<256xi32,#blocked0>) { + // CHECK: llvm.getelementptr + // CHECK: llvm.getelementptr + %0 = tt.addptr %arg0, %arg1 : tensor<256x!tt.ptr, #blocked0>, tensor<256xi32, #blocked0> +@@ -373,7 +373,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK: llvm.mlir.global external @global_smem + // CHECK-LABEL: basic_alloc_tensor +- func @basic_alloc_tensor() { ++ func.func @basic_alloc_tensor() { + // CHECK: llvm.mlir.addressof @global_smem + // CHECK-NEXT: llvm.bitcast + // CHECK-NEXT: llvm.mlir.constant +@@ -390,7 +390,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK: llvm.mlir.global external @global_smem + // CHECK-LABEL: basic_extract_slice +- func @basic_extract_slice() { ++ func.func @basic_extract_slice() { + // CHECK: llvm.mlir.addressof @global_smem + // CHECK: llvm.extractvalue + // CHECK-NEXT: llvm.extractvalue +@@ -423,7 +423,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_async_wait +- func @basic_async_wait() { ++ func.func @basic_async_wait() { + // CHECK: cp.async.wait_group 0x4 + triton_gpu.async_wait {num = 4: i32} + return +@@ -442,7 +442,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #A = #triton_gpu.shared<{vec = 8, perPhase = 1, maxPhase = 4, order = [1, 0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_insert_slice_async_fallback +- func @basic_insert_slice_async_fallback(%arg0: !tt.ptr {tt.divisibility = 1 : i32}) { ++ func.func @basic_insert_slice_async_fallback(%arg0: !tt.ptr {tt.divisibility = 1 : i32}) { + %off0_ = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #slice2d1> + %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<64xi32, #slice3d0> + %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<16xi32, #slice2d1>) -> tensor<16x1xi32, #block2> +@@ -481,7 +481,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #A = #triton_gpu.shared<{vec = 8, perPhase = 1, maxPhase = 4, order = [1, 0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_insert_slice_async_v4 +- func @basic_insert_slice_async_v4(%arg0: !tt.ptr {tt.divisibility = 32 : i32}) { ++ func.func @basic_insert_slice_async_v4(%arg0: !tt.ptr {tt.divisibility = 32 : i32}) { + %off0_ = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #slice2d1> + %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<64xi32, #slice3d0> + %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<16xi32, #slice2d1>) -> tensor<16x1xi32, #block2> +@@ -523,7 +523,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #A = #triton_gpu.shared<{vec = 1, perPhase = 1, maxPhase = 4, order = [1, 0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_insert_slice_async_v1 +- func @basic_insert_slice_async_v1(%arg0: !tt.ptr {tt.divisibility = 4 : i32}) { ++ func.func @basic_insert_slice_async_v1(%arg0: !tt.ptr {tt.divisibility = 4 : i32}) { + %off0_ = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #slice2d1> + %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #slice3d0> + %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<16xi32, #slice2d1>) -> tensor<16x1xi32, #block2> +@@ -568,7 +568,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #A = #triton_gpu.shared<{vec = 1, perPhase = 1, maxPhase = 4, order = [1, 0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_insert_slice_async_v1_multictas +- func @basic_insert_slice_async_v1_multictas(%arg0: !tt.ptr {tt.divisibility = 4 : i32}) { ++ func.func @basic_insert_slice_async_v1_multictas(%arg0: !tt.ptr {tt.divisibility = 4 : i32}) { + %off0_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #slice2d1> + %off1_ = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #slice3d0> + %off0 = tt.expand_dims %off0_ {axis = 1 : i32} : (tensor<32xi32, #slice2d1>) -> tensor<32x1xi32, #block2> +@@ -619,7 +619,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK: basic_splat +- func @basic_splat(%ptr: !tt.ptr) { ++ func.func @basic_splat(%ptr: !tt.ptr) { + // CHECK: llvm.mlir.undef + // CHECK: llvm.insertvalue + // CHECK: llvm.insertvalue +@@ -633,7 +633,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: basic_store +- func @basic_store(%ptrs: tensor<256x!tt.ptr, #blocked0>, %vals: tensor<256xf32, #blocked0>, %mask: tensor<256xi1, #blocked0>) { ++ func.func @basic_store(%ptrs: tensor<256x!tt.ptr, #blocked0>, %vals: tensor<256xf32, #blocked0>, %mask: tensor<256xi1, #blocked0>) { + // CHECK: llvm.inline_asm + // CHECK-SAME: st.global.b32 [ ${{.*}} + 0 ], { ${{.*}} }; + // CHECK: llvm.inline_asm +@@ -650,7 +650,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> + // CHECK-LABEL: convert_layout_blocked_blocked +- func @convert_layout_blocked_blocked(%arg0: tensor<16x16xf32, #blocked0>) { ++ func.func @convert_layout_blocked_blocked(%arg0: tensor<16x16xf32, #blocked0>) { + // CHECK: llvm.mlir.addressof @global_smem + // CHECK: llvm.store + // CHECK-SAME: !llvm.ptr, 3> +@@ -697,7 +697,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> + // CHECK-LABEL: convert_layout_blocked_blocked_vec +- func @convert_layout_blocked_blocked_vec(%arg0: tensor<16x16xf32, #blocked0>) { ++ func.func @convert_layout_blocked_blocked_vec(%arg0: tensor<16x16xf32, #blocked0>) { + // CHECK: llvm.mlir.addressof @global_smem + // CHECK: llvm.store + // CHECK-SAME: !llvm.ptr, 3> +@@ -720,7 +720,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> + // CHECK-LABEL: convert_layout_blocked_blocked_multi_rep +- func @convert_layout_blocked_blocked_multi_rep(%arg0: tensor<16x16xf32, #blocked0>) { ++ func.func @convert_layout_blocked_blocked_multi_rep(%arg0: tensor<16x16xf32, #blocked0>) { + // CHECK: llvm.mlir.addressof @global_smem + // CHECK: llvm.store + // CHECK-SAME: !llvm.ptr, 3> +@@ -751,7 +751,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma0}> + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK-LABEL: convert_dot +- func @convert_dot(%A: tensor<16x16xf16, #blocked0>, %B: tensor<16x16xf16, #blocked0>) { ++ func.func @convert_dot(%A: tensor<16x16xf16, #blocked0>, %B: tensor<16x16xf16, #blocked0>) { + %AA = triton_gpu.convert_layout %A : (tensor<16x16xf16, #blocked0>) -> tensor<16x16xf16, #shared0> + %BB = triton_gpu.convert_layout %B : (tensor<16x16xf16, #blocked0>) -> tensor<16x16xf16, #shared0> + // CHECK: llvm.inline_asm +@@ -775,7 +775,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + // TODO: problems in MLIR's parser on slice layout + // #blocked0 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [8, 4], warpsPerCTA = [1, 1], order = [1, 0]}> + // module attributes {"triton_gpu.num-warps" = 1 : i32} { +-// func @make_range_sliced_layout() { ++// func.func @make_range_sliced_layout() { + // %0 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #triton_gpu.slice<{dim = 0, parent = #blocked0}>> + // return + // } +@@ -788,7 +788,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> + // CHECK-LABEL: convert_layout_mmav2_block +- func @convert_layout_mmav2_blocked(%arg0: tensor<32x16xf32, #mma>) { ++ func.func @convert_layout_mmav2_blocked(%arg0: tensor<32x16xf32, #mma>) { + // CHECK: llvm.store + // CHECK-SAME: !llvm.ptr, 3> + // CHECK: llvm.store +@@ -808,7 +808,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> + // CHECK-LABEL: convert_layout_mmav1_block +- func @convert_layout_mmav1_blocked(%arg0: tensor<32x64xf32, #mma>) { ++ func.func @convert_layout_mmav1_blocked(%arg0: tensor<32x64xf32, #mma>) { + // CHECK: llvm.store + // CHECK-SAME: !llvm.ptr, 3> + // CHECK: llvm.store +@@ -831,7 +831,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK: llvm.mlir.global external @global_smem() {addr_space = 3 : i32} : !llvm.array<0 x i8> + // CHECK-LABEL: convert_layout_blocked_shared +- func @convert_layout_blocked_shared(%arg0: tensor<128x32xf32, #blocked0>) { ++ func.func @convert_layout_blocked_shared(%arg0: tensor<128x32xf32, #blocked0>) { + // CHECK: llvm.store + // CHECK-SAME: !llvm.ptr, 3> + // CHECK: llvm.store +@@ -847,7 +847,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + #blocked1 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [4, 8], warpsPerCTA = [1, 1], order = [1, 0]}> + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK-LABEL: convert_blocked1d_to_slice0 +- func @convert_blocked1d_to_slice0(%src:tensor<32xi32, #blocked0>) { ++ func.func @convert_blocked1d_to_slice0(%src:tensor<32xi32, #blocked0>) { + // CHECK-COUNT-4: llvm.load {{.*}} : !llvm.ptr, 3> + %cvt = triton_gpu.convert_layout %src : (tensor<32xi32, #blocked0>) -> tensor<32xi32, #triton_gpu.slice<{dim = 0, parent = #blocked1}>> + return +@@ -860,7 +860,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + #blocked1 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [4, 8], warpsPerCTA = [1, 1], order = [1, 0]}> + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK-LABEL: convert_blocked1d_to_slice1 +- func @convert_blocked1d_to_slice1(%src:tensor<32xi32, #blocked0>) { ++ func.func @convert_blocked1d_to_slice1(%src:tensor<32xi32, #blocked0>) { + // CHECK-COUNT-32: llvm.load {{.*}} : !llvm.ptr, 3> + %cvt = triton_gpu.convert_layout %src : (tensor<32xi32, #blocked0>) -> tensor<32xi32, #triton_gpu.slice<{dim = 1, parent = #blocked1}>> + return +@@ -873,7 +873,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + #blocked1 = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK-LABEL: convert_blocked_to_blocked_ptr +- func @convert_blocked_to_blocked_ptr(%src:tensor<32x!tt.ptr, #blocked0>) { ++ func.func @convert_blocked_to_blocked_ptr(%src:tensor<32x!tt.ptr, #blocked0>) { + // CHECK: llvm.ptrtoint + // CHECK: llvm.store + // CHECK: nvvm.barrier0 +@@ -892,7 +892,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + #dot_operand_a = #triton_gpu.dot_op<{opIdx=0, parent=#mma}> + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { +- func @matmul_kernel_dot_operand_layout(%ptr:!tt.ptr {tt.divisibility = 16 : i32}, ++ func.func @matmul_kernel_dot_operand_layout(%ptr:!tt.ptr {tt.divisibility = 16 : i32}, + %a:tensor<128x32xf16, #shared>, %b:tensor<32x256xf16, #shared>) { + %cst = arith.constant dense<0.000000e+00> : tensor<128x256xf32, #mma> + // CHECK: ldmatrix.sync.aligned.m8n8.x4.shared.b16 +@@ -918,7 +918,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #dot_operand_a = #triton_gpu.dot_op<{opIdx=0, parent=#mma, isMMAv1Row=true}> + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma, isMMAv1Row=true}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { +- func @matmul884_kernel_dot_operand_layout(%ptr:!tt.ptr {tt.divisibility = 16 : i32}, ++ func.func @matmul884_kernel_dot_operand_layout(%ptr:!tt.ptr {tt.divisibility = 16 : i32}, + %a:tensor<32x64xf16, #shared0>, %b:tensor<64x64xf16, #shared1>) { + %cst = arith.constant dense<0.000000e+00> : tensor<32x64xf32, #mma> + // CHECK: ldmatrix.sync.aligned.m8n8.x4.shared.b16 +@@ -941,7 +941,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #dot_operand_a = #triton_gpu.dot_op<{opIdx=0, parent=#blocked}> + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#blocked}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { +- func @matmul_fmadot(%ptr:!tt.ptr {tt.divisibility = 16 : i32}, ++ func.func @matmul_fmadot(%ptr:!tt.ptr {tt.divisibility = 16 : i32}, + %a:tensor<32x16xf32, #shared>, %b:tensor<16x32xf32, #shared>) { + %cst = arith.constant dense<0.000000e+00> : tensor<32x32xf32, #blocked> + // CHECK: llvm.intr.fmuladd +@@ -965,7 +965,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #dot_operand_b = #triton_gpu.dot_op<{opIdx=1, parent=#mma}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: matmul_tf32dot +- func @matmul_tf32dot(%ptr:!tt.ptr {tt.divisibility = 16 : i32}, ++ func.func @matmul_tf32dot(%ptr:!tt.ptr {tt.divisibility = 16 : i32}, + %a:tensor<32x16xf32, #shared>, %b:tensor<16x32xf32, #shared>) { + %cst = arith.constant dense<0.000000e+00> : tensor<32x32xf32, #mma> + // CHECK: llvm.inline_asm +@@ -1000,7 +1000,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: atomic_add_f32 +- func @atomic_add_f32(%arg0 : tensor<256x!tt.ptr, #blocked0>, %arg1 : tensor<256xi1, #blocked0>, %arg2 : tensor<256xf32, #blocked0>) { ++ func.func @atomic_add_f32(%arg0 : tensor<256x!tt.ptr, #blocked0>, %arg1 : tensor<256xi1, #blocked0>, %arg2 : tensor<256xf32, #blocked0>) { + // CHECK: llvm.inline_asm + // CHECK-SAME: atom.global.gpu.add.f32 + %0 = "tt.atomic_rmw" (%arg0, %arg2, %arg1) {atomic_rmw_op = 5 : i32} : (tensor<256x!tt.ptr, #blocked0>, tensor<256xf32, #blocked0>, tensor<256xi1, #blocked0>) -> tensor<256xf32, #blocked0> +@@ -1012,7 +1012,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + +-func @test_get_program_id(%a: tensor<32x!tt.ptr, #blocked0>) { ++func.func @test_get_program_id(%a: tensor<32x!tt.ptr, #blocked0>) { + %blockidx = tt.get_program_id {axis=0:i32} : i32 + %blockidy = tt.get_program_id {axis=1:i32} : i32 + %blockidz = tt.get_program_id {axis=2:i32} : i32 +@@ -1032,7 +1032,7 @@ func @test_get_program_id(%a: tensor<32x!tt.ptr, #blocked0>) { + // ----- + #blocked0 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { +- func @test_get_num_program(%a: tensor<32x!tt.ptr, #blocked0>) { ++ func.func @test_get_num_program(%a: tensor<32x!tt.ptr, #blocked0>) { + // CHECK: nvvm.read.ptx.sreg.nctaid.x + // CHECK: nvvm.read.ptx.sreg.nctaid.y + // CHECK: nvvm.read.ptx.sreg.nctaid.z +@@ -1052,7 +1052,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #blocked0 = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [1], order = [0]}> + module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK-LABEL: test_index_cache +- func @test_index_cache() { ++ func.func @test_index_cache() { + // CHECK: nvvm.read.ptx.sreg.tid.x + %0 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32, #blocked0> + // CHECK-NOT: nvvm.read.ptx.sreg.tid.x +@@ -1066,7 +1066,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + #shared0 = #triton_gpu.shared<{vec = 8, perPhase = 2, maxPhase = 4, order = [1, 0]}> + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK-LABEL: test_base_index_cache +- func @test_base_index_cache(%arg0: tensor<128x32xf32, #blocked0>) { ++ func.func @test_base_index_cache(%arg0: tensor<128x32xf32, #blocked0>) { + // CHECK: nvvm.read.ptx.sreg.tid.x + %0 = triton_gpu.convert_layout %arg0 : (tensor<128x32xf32, #blocked0>) -> tensor<128x32xf32, #shared0> + // CHECK-NOT: nvvm.read.ptx.sreg.tid.x +@@ -1080,7 +1080,7 @@ module attributes {"triton_gpu.num-warps" = 1 : i32} { + #shared0 = #triton_gpu.shared<{vec = 8, perPhase = 2, maxPhase = 4, order = [1, 0]}> + module attributes {"triton_gpu.num-warps" = 1 : i32} { + // CHECK-LABEL: test_index_cache_different_block +- func @test_index_cache_different_block(%arg0: tensor<128x32xf32, #blocked0>, %arg1: i1) { ++ func.func @test_index_cache_different_block(%arg0: tensor<128x32xf32, #blocked0>, %arg1: i1) { + // CHECK: nvvm.read.ptx.sreg.tid.x + %0 = triton_gpu.convert_layout %arg0 : (tensor<128x32xf32, #blocked0>) -> tensor<128x32xf32, #shared0> + scf.if %arg1 { +diff --git a/test/Target/tritongpu_to_llvmir.mlir b/test/Target/tritongpu_to_llvmir.mlir +index cafff3ca60..114d3a9eb2 100644 +--- a/test/Target/tritongpu_to_llvmir.mlir ++++ b/test/Target/tritongpu_to_llvmir.mlir +@@ -4,11 +4,11 @@ + // CHECK-LABEL: ; ModuleID = 'LLVMDialectModule' + // CHECK: define void @test_empty_kernel + // CHECK: !nvvm.annotations +-// CHECK: !{void (i32, half addrspace(1)*)* @test_empty_kernel, !"maxntidx", i32 128} ++// CHECK: !{ptr @test_empty_kernel, !"maxntidx", i32 128} + + module attributes {"triton_gpu.num-warps" = 4 : i32} { + +-func @test_empty_kernel(%lb : index, %A : !tt.ptr) { ++func.func @test_empty_kernel(%lb : index, %A : !tt.ptr) { + + return + } +diff --git a/test/Target/tritongpu_to_ptx.mlir b/test/Target/tritongpu_to_ptx.mlir +index 404e970a29..12742ad9e2 100644 +--- a/test/Target/tritongpu_to_ptx.mlir ++++ b/test/Target/tritongpu_to_ptx.mlir +@@ -6,7 +6,7 @@ + + module attributes {"triton_gpu.num-warps" = 4 : i32} { + +-func @test_empty_kernel(%lb : index, %A : !tt.ptr) { ++func.func @test_empty_kernel(%lb : index, %A : !tt.ptr) { + + return + } +diff --git a/test/Triton/combine.mlir b/test/Triton/combine.mlir +index 050a3f7565..5ef6790e69 100644 +--- a/test/Triton/combine.mlir ++++ b/test/Triton/combine.mlir +@@ -2,10 +2,10 @@ + // RUN: triton-opt %s -split-input-file -canonicalize -triton-combine | FileCheck %s + + // CHECK-LABEL: @test_combine_dot_add_pattern +-func @test_combine_dot_add_pattern() -> (tensor<128x128xf32>, tensor<128x128xf32>) { +- // CHECK: %[[d:.*]] = arith.constant dense<3.000000e+00> : tensor<128x128xf32> +- // CHECK: %[[b:.*]] = arith.constant dense<2.000000e+00> : tensor<128x128xf32> +- // CHECK: %[[a:.*]] = arith.constant dense<1.000000e+00> : tensor<128x128xf32> ++func.func @test_combine_dot_add_pattern() -> (tensor<128x128xf32>, tensor<128x128xf32>) { ++ // CHECK-DAG: %[[d:.*]] = arith.constant dense<3.000000e+00> : tensor<128x128xf32> ++ // CHECK-DAG: %[[b:.*]] = arith.constant dense<2.000000e+00> : tensor<128x128xf32> ++ // CHECK-DAG: %[[a:.*]] = arith.constant dense<1.000000e+00> : tensor<128x128xf32> + %a = arith.constant dense<1.0> : tensor<128x128xf32> + %b = arith.constant dense<2.0> : tensor<128x128xf32> + %zero = arith.constant dense<0.0> : tensor<128x128xf32> +@@ -24,7 +24,7 @@ func @test_combine_dot_add_pattern() -> (tensor<128x128xf32>, tensor<128x128xf32 + + + // COM: CHECK-LABEL: @test_combine_addptr_pattern +-func @test_combine_addptr_pattern(%base: !tt.ptr) -> tensor<8x!tt.ptr> { ++func.func @test_combine_addptr_pattern(%base: !tt.ptr) -> tensor<8x!tt.ptr> { + %off0 = arith.constant 10 : i32 + %off1 = arith.constant 15 : i32 + +@@ -47,46 +47,46 @@ func @test_combine_addptr_pattern(%base: !tt.ptr) -> tensor<8x!tt.ptr> + + + // CHECK-LABEL: @test_combine_select_masked_load_pattern +-func @test_combine_select_masked_load_pattern(%ptr: tensor<8x!tt.ptr>, %cond: i1) -> (tensor<8xf32>, tensor<8xf32>) { ++func.func @test_combine_select_masked_load_pattern(%ptr: tensor<8x!tt.ptr>, %cond: i1) -> (tensor<8xf32>, tensor<8xf32>) { + %mask = tt.broadcast %cond : (i1) -> tensor<8xi1> + %false_val = arith.constant dense<0.0> : tensor<8xf32> + + // CHECK: %[[res1:.*]] = tt.load %{{.*}}, %{{.*}}, %{{.*}} {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> + %x = tt.load %ptr, %mask, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> +- %0 = select %cond, %x, %false_val : tensor<8xf32> ++ %0 = arith.select %cond, %x, %false_val : tensor<8xf32> + + // CHECK: %[[res2:.*]] = tt.load %{{.*}}, %{{.*}}, %{{.*}} {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> + %y = tt.load %ptr, %mask, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> +- %1 = select %cond, %y, %false_val : tensor<8xf32> ++ %1 = arith.select %cond, %y, %false_val : tensor<8xf32> + + // CHECK: return %[[res1]], %[[res2]] : tensor<8xf32>, tensor<8xf32> + return %0, %1 : tensor<8xf32>, tensor<8xf32> + } + + // CHECK-LABEL: @test_combine_select_masked_load_fail_pattern +-func @test_combine_select_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr>, %dummy_load: tensor<8xf32>, %dummy_broadcast: tensor<8xi1>, %cond0: i1, %cond1: i1) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) { ++func.func @test_combine_select_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr>, %dummy_load: tensor<8xf32>, %dummy_broadcast: tensor<8xi1>, %cond0: i1, %cond1: i1) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) { + %false_val = arith.constant dense<0.0> : tensor<8xf32> + + // Case 1: value at the "load" position is not an "op". Select should not be canonicalized. +- // CHECK: %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> +- %0 = select %cond0, %dummy_load, %false_val : tensor<8xf32> ++ // CHECK: %{{.*}} = arith.select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> ++ %0 = arith.select %cond0, %dummy_load, %false_val : tensor<8xf32> + + // Case 2: value at the "broadcast" position is not an "op". Select should not be canonicalized. + %real_load0 = tt.load %ptr, %dummy_broadcast, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> +- // CHECK: %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> +- %1 = select %cond0, %real_load0, %false_val : tensor<8xf32> ++ // CHECK: %{{.*}} = arith.select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> ++ %1 = arith.select %cond0, %real_load0, %false_val : tensor<8xf32> + + // Case 3: condition of "broadcast" is not the same as the condition of "select". Select should not be canonicalized. + %cond0_ = tt.broadcast %cond0 : (i1) -> tensor<8xi1> + %real_load1 = tt.load %ptr, %cond0_, %false_val {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<8xf32> +- // CHECK: %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> +- %2 = select %cond1, %real_load1, %false_val : tensor<8xf32> ++ // CHECK: %{{.*}} = arith.select %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> ++ %2 = arith.select %cond1, %real_load1, %false_val : tensor<8xf32> + + return %0, %1, %2 : tensor<8xf32>, tensor<8xf32>, tensor<8xf32> + } + + // CHECK-LABEL: @test_combine_broadcast_constant_pattern +-func @test_combine_broadcast_constant_pattern(%cst : f32) -> tensor<8x2xf32> { ++func.func @test_combine_broadcast_constant_pattern(%cst : f32) -> tensor<8x2xf32> { + // CHECK: %[[cst:.*]] = arith.constant dense<1.000000e+00> : tensor<8x2xf32> + %const = arith.constant dense<1.0> : tensor<8xf32> + %bst_out = tt.broadcast %const : (tensor<8xf32>) -> tensor<8x2xf32> +@@ -96,7 +96,7 @@ func @test_combine_broadcast_constant_pattern(%cst : f32) -> tensor<8x2xf32> { + } + + // CHECK-LABEL: @test_canonicalize_masked_load_pattern +-func @test_canonicalize_masked_load_pattern(%ptr: tensor<8x!tt.ptr>) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) { ++func.func @test_canonicalize_masked_load_pattern(%ptr: tensor<8x!tt.ptr>) -> (tensor<8xf32>, tensor<8xf32>, tensor<8xf32>) { + %true_mask = arith.constant dense : tensor<8xi1> + %false_mask = arith.constant dense : tensor<8xi1> + %other_val = arith.constant dense<0.0> : tensor<8xf32> +@@ -117,7 +117,7 @@ func @test_canonicalize_masked_load_pattern(%ptr: tensor<8x!tt.ptr>) -> (te + } + + // CHECK-LABEL: @test_canonicalize_masked_load_fail_pattern +-func @test_canonicalize_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr>, %mask: tensor<8xi1>) -> (tensor<8xf32>, tensor<8xf32>) { ++func.func @test_canonicalize_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr>, %mask: tensor<8xi1>) -> (tensor<8xf32>, tensor<8xf32>) { + %other_val = arith.constant dense<0.0> : tensor<8xf32> + + // Case: value at the "mask" position is not an "op". Load should not be canonicalized. +@@ -130,7 +130,7 @@ func @test_canonicalize_masked_load_fail_pattern(%ptr: tensor<8x!tt.ptr>, % + } + + // CHECK-LABEL: @test_canonicalize_masked_store_pattern +-func @test_canonicalize_masked_store_pattern(%ptr: tensor<8x!tt.ptr>, %val: tensor<8xf32>) { ++func.func @test_canonicalize_masked_store_pattern(%ptr: tensor<8x!tt.ptr>, %val: tensor<8xf32>) { + %true_mask = arith.constant dense : tensor<8xi1> + %false_mask = arith.constant dense : tensor<8xi1> + +@@ -144,7 +144,7 @@ func @test_canonicalize_masked_store_pattern(%ptr: tensor<8x!tt.ptr>, %val: + } + + // CHECK-LABEL: @test_canonicalize_masked_store_fail_pattern +-func @test_canonicalize_masked_store_fail_pattern(%ptr: tensor<8x!tt.ptr>, %val: tensor<8xf32>, %mask: tensor<8xi1>) { ++func.func @test_canonicalize_masked_store_fail_pattern(%ptr: tensor<8x!tt.ptr>, %val: tensor<8xf32>, %mask: tensor<8xi1>) { + // Case: value at the "mask" position is not an "op". Store should not be canonicalized. + // CHECK: tt.store %{{.*}}, %{{.*}}, %{{.*}} : tensor<8xf32> + tt.store %ptr, %val, %mask : tensor<8xf32> +diff --git a/test/Triton/vecadd.mlir b/test/Triton/vecadd.mlir +index 0b69ef3054..f5019b1cdd 100644 +--- a/test/Triton/vecadd.mlir ++++ b/test/Triton/vecadd.mlir +@@ -1,7 +1,7 @@ + // RUN: triton-opt %s -verify-diagnostics + + module { +- func @add_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32__(%arg0: !tt.ptr, %arg1: !tt.ptr, %arg2: !tt.ptr, %arg3: i32, %arg4: i32, %arg5: i32) { ++ func.func @add_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32__(%arg0: !tt.ptr, %arg1: !tt.ptr, %arg2: !tt.ptr, %arg3: i32, %arg4: i32, %arg5: i32) { + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %c256_i32 = arith.constant 256 : i32 + %1 = arith.muli %0, %c256_i32 : i32 +@@ -43,7 +43,7 @@ module { + } + } + // module { +-// func @add_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32__(%arg0: !tt.ptr, %arg1: !tt.ptr, %arg2: !tt.ptr, %arg3: i32, %arg4: i32, %arg5: i32) { ++// func.func @add_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32__(%arg0: !tt.ptr, %arg1: !tt.ptr, %arg2: !tt.ptr, %arg3: i32, %arg4: i32, %arg5: i32) { + // %c64 = arith.constant 64 : index + // %c32 = arith.constant 32 : index + // %c0 = arith.constant 0 : index +diff --git a/test/TritonGPU/coalesce.mlir b/test/TritonGPU/coalesce.mlir +index 60e359f527..51cccccfbd 100644 +--- a/test/TritonGPU/coalesce.mlir ++++ b/test/TritonGPU/coalesce.mlir +@@ -19,7 +19,7 @@ module attributes {"triton_gpu.num-warps" = 4 : i32} { + // CHECK: [[store_val:%.*]] = triton_gpu.convert_layout {{.*}} -> tensor<64x64xf32, [[col_layout]]> + // CHECK: [[store_mask:%.*]] = triton_gpu.convert_layout {{.*}} -> tensor<64x64xi1, [[col_layout]]> + // CHECK: tt.store [[store_ptr]], [[store_val]], [[store_mask]] +-func @transpose(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, ++func.func @transpose(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, + %arg1: i32 {tt.divisibility = 16 : i32}, + %arg2: !tt.ptr {tt.divisibility = 16 : i32}, + %arg3: i32 {tt.divisibility = 16 : i32}) { +diff --git a/test/TritonGPU/combine.mlir b/test/TritonGPU/combine.mlir +index 2c009ffa48..7e9cb9d504 100644 +--- a/test/TritonGPU/combine.mlir ++++ b/test/TritonGPU/combine.mlir +@@ -9,7 +9,7 @@ + // CHECK: [[col_layout:#.*]] = #triton_gpu.blocked<{sizePerThread = [4, 1], threadsPerWarp = [16, 2], warpsPerCTA = [4, 1], order = [0, 1]}> + // CHECK: [[col_layout_novec:#.*]] = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [4, 1], order = [0, 1]}> + // CHECK-LABEL: cst +-func @cst() -> tensor<1024xi32, #layout1> { ++func.func @cst() -> tensor<1024xi32, #layout1> { + %cst = arith.constant dense<0> : tensor<1024xi32, #layout0> + %1 = triton_gpu.convert_layout %cst : (tensor<1024xi32, #layout0>) -> tensor<1024xi32, #layout1> + // CHECK-NOT: triton_gpu.convert_layout +@@ -18,7 +18,7 @@ func @cst() -> tensor<1024xi32, #layout1> { + } + + // CHECK-LABEL: range +-func @range() -> tensor<1024xi32, #layout1> { ++func.func @range() -> tensor<1024xi32, #layout1> { + %0 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #layout0> + %1 = triton_gpu.convert_layout %0 : (tensor<1024xi32, #layout0>) -> tensor<1024xi32, #layout1> + // CHECK-NOT: triton_gpu.convert_layout +@@ -27,7 +27,7 @@ func @range() -> tensor<1024xi32, #layout1> { + } + + // CHECK-LABEL: splat +-func @splat(%arg0: i32) -> tensor<1024xi32, #layout1> { ++func.func @splat(%arg0: i32) -> tensor<1024xi32, #layout1> { + %0 = tt.splat %arg0 : (i32) -> tensor<1024xi32, #layout0> + %1 = triton_gpu.convert_layout %0 : (tensor<1024xi32, #layout0>) -> tensor<1024xi32, #layout1> + // CHECK-NOT: triton_gpu.convert_layout +@@ -36,7 +36,7 @@ func @splat(%arg0: i32) -> tensor<1024xi32, #layout1> { + } + + // CHECK-LABEL: remat +-func @remat(%arg0: i32) -> tensor<1024xi32, #layout1> { ++func.func @remat(%arg0: i32) -> tensor<1024xi32, #layout1> { + %0 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #layout0> + %1 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #layout0> + %2 = arith.muli %0, %1 : tensor<1024xi32, #layout0> +@@ -56,7 +56,7 @@ func @remat(%arg0: i32) -> tensor<1024xi32, #layout1> { + } + + // CHECK-LABEL: remat_load_store +-func @remat_load_store(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @remat_load_store(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { + %0 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #layout0> + %1 = tt.splat %arg : (!tt.ptr) -> tensor<64x!tt.ptr, #layout0> + %2 = tt.addptr %1, %0 : tensor<64x!tt.ptr, #layout0>, tensor<64xi32, #layout0> +@@ -70,7 +70,7 @@ func @remat_load_store(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { + + // Don't rematerialize vectorized loads + // CHECK-LABEL: remat_expensive +-func @remat_expensive(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @remat_expensive(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { + %0 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #layout1> + %1 = tt.splat %arg : (!tt.ptr) -> tensor<64x!tt.ptr, #layout1> + %2 = tt.addptr %1, %0 : tensor<64x!tt.ptr, #layout1>, tensor<64xi32, #layout1> +@@ -85,7 +85,7 @@ func @remat_expensive(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { + + // Don't rematerialize loads when original and target layouts are different + // CHECK-LABEL: remat_multi_layout +-func @remat_multi_layout(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @remat_multi_layout(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { + %0 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #layout0> + %1 = tt.splat %arg : (!tt.ptr) -> tensor<64x!tt.ptr, #layout0> + %2 = tt.addptr %1, %0 : tensor<64x!tt.ptr, #layout0>, tensor<64xi32, #layout0> +@@ -100,7 +100,7 @@ func @remat_multi_layout(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { + + // Always rematerialize single value loads + // CHECK-LABEL: remat_single_value +-func @remat_single_value(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @remat_single_value(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { + %0 = tt.splat %arg : (!tt.ptr) -> tensor<1x!tt.ptr, #layout1> + %1 = tt.load %0 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<1xi32, #layout1> + // CHECK-NOT: triton_gpu.convert_layout +@@ -111,7 +111,7 @@ func @remat_single_value(%arg: !tt.ptr {tt.divisibility = 16 : i32}) { + } + + // CHECK-LABEL: if +-func @if(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @if(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { + // CHECK-NOT: triton_gpu.convert_layout + %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout1> + %0 = tt.get_program_id {axis = 0 : i32} : i32 +@@ -128,7 +128,7 @@ func @if(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { + } + + // CHECK-LABEL: if_convert_else_not +-func @if_convert_else_not(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @if_convert_else_not(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { + %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout0> + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = tt.splat %0 : (i32) -> tensor<1024xi32, #layout0> +@@ -149,7 +149,7 @@ func @if_convert_else_not(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 + } + + // CHECK-LABEL: if_not_else_convert +-func @if_not_else_convert(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @if_not_else_convert(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { + %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout0> + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = tt.splat %0 : (i32) -> tensor<1024xi32, #layout0> +@@ -170,7 +170,7 @@ func @if_not_else_convert(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 + } + + // CHECK-LABEL: if_else_both_convert +-func @if_else_both_convert(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { ++func.func @if_else_both_convert(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 : i32}) { + %c32_i32 = arith.constant dense<32> : tensor<1024xi32, #layout0> + %0 = tt.get_program_id {axis = 0 : i32} : i32 + %1 = tt.splat %0 : (i32) -> tensor<1024xi32, #layout0> +@@ -200,7 +200,7 @@ func @if_else_both_convert(%arg0: i32, %arg1: !tt.ptr {tt.divisibility = 16 + #blocked4 = #triton_gpu.blocked<{sizePerThread = [4, 1], threadsPerWarp = [16, 2], warpsPerCTA = [4, 1], order = [0, 1]}> + + // CHECK-LABEL: transpose +-func @transpose(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32}) { ++func.func @transpose(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32}) { + // CHECK-NOT: triton_gpu.convert_layout + // CHECK: [[loaded_val:%.*]] = tt.load {{.*}}, {{%cst.*}}, {{%cst.*}} {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<64x64xf32, [[row_layout]]> + // CHECK: [[cvt_val:%.*]] = triton_gpu.convert_layout [[loaded_val]] : (tensor<64x64xf32, [[row_layout]]>) -> tensor<64x64xf32, [[col_layout]]> +@@ -241,7 +241,7 @@ func @transpose(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: i32 {tt + } + + // CHECK-LABEL: loop +-func @loop(%arg0: !tt.ptr, %arg1: i32, %arg2: !tt.ptr, %arg3: i32, %arg4: i32) { ++func.func @loop(%arg0: !tt.ptr, %arg1: i32, %arg2: !tt.ptr, %arg3: i32, %arg4: i32) { + // CHECK-NOT: triton_gpu.convert_layout + // CHECK: [[loop_ret:%.*]]:2 = scf.for {{.*}} -> (tensor<64x64xf32, [[row_layout]]>, tensor<64x64x!tt.ptr, [[row_layout]]>) + // CHECK-NEXT: {{.*}} = tt.load {{.*}} : tensor<64x64xf32, [[row_layout]]> +@@ -295,7 +295,7 @@ func @loop(%arg0: !tt.ptr, %arg1: i32, %arg2: !tt.ptr, %arg3: i32, %ar + } + + // CHECK-LABEL: vecadd +-func @vecadd(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32) { ++func.func @vecadd(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32) { + // CHECK-NOT: triton_gpu.convert_layout + %c256_i32 = arith.constant 256 : i32 + %0 = tt.get_program_id {axis = 0 : i32} : i32 +@@ -327,7 +327,7 @@ func @vecadd(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}) { ++func.func @select(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}) { + // CHECK-NOT: triton_gpu.convert_layout + %cst = arith.constant dense<30000> : tensor<1x1xi32, #blocked2> + %cst_0 = arith.constant dense<30000> : tensor<1x512xi32, #blocked2> +@@ -378,7 +378,7 @@ func @select(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: !tt.ptr {tt.divisibility = 16 : i32}, %arg4: !tt.ptr {tt.divisibility = 16 : i32}, %arg5: !tt.ptr {tt.divisibility = 16 : i32}, %arg6: !tt.ptr {tt.divisibility = 16 : i32}, %arg7: !tt.ptr {tt.divisibility = 16 : i32}, %arg8: !tt.ptr {tt.divisibility = 16 : i32}, %arg9: !tt.ptr {tt.divisibility = 16 : i32}, %arg10: !tt.ptr {tt.divisibility = 16 : i32}, %arg11: !tt.ptr {tt.divisibility = 16 : i32}, %arg12: !tt.ptr {tt.divisibility = 16 : i32}, %arg13: !tt.ptr {tt.divisibility = 16 : i32}, %arg14: !tt.ptr {tt.divisibility = 16 : i32}, %arg15: !tt.ptr {tt.divisibility = 16 : i32}, %arg16: i32 {tt.divisibility = 16 : i32}) { ++func.func public @long_func(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: !tt.ptr {tt.divisibility = 16 : i32}, %arg4: !tt.ptr {tt.divisibility = 16 : i32}, %arg5: !tt.ptr {tt.divisibility = 16 : i32}, %arg6: !tt.ptr {tt.divisibility = 16 : i32}, %arg7: !tt.ptr {tt.divisibility = 16 : i32}, %arg8: !tt.ptr {tt.divisibility = 16 : i32}, %arg9: !tt.ptr {tt.divisibility = 16 : i32}, %arg10: !tt.ptr {tt.divisibility = 16 : i32}, %arg11: !tt.ptr {tt.divisibility = 16 : i32}, %arg12: !tt.ptr {tt.divisibility = 16 : i32}, %arg13: !tt.ptr {tt.divisibility = 16 : i32}, %arg14: !tt.ptr {tt.divisibility = 16 : i32}, %arg15: !tt.ptr {tt.divisibility = 16 : i32}, %arg16: i32 {tt.divisibility = 16 : i32}) { + %cst = arith.constant dense<1.000000e+00> : tensor<1024xf32, #blocked0> + %cst_0 = arith.constant dense<5.000000e-04> : tensor<1024xf32, #blocked0> + %cst_1 = arith.constant dense<0.999499976> : tensor<1024xf32, #blocked0> +@@ -775,7 +775,7 @@ func public @long_func(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: + // A mnist model from torch inductor. + // Check if topological sort is working correct and there's no unnecessary convert + // CHECK-LABEL: mnist +-func public @mnist(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}, %arg3: i32) { ++func.func public @mnist(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}, %arg3: i32) { + // CHECK-NOT: triton_gpu.convert_layout + %cst = arith.constant dense<10> : tensor<16x1xi32, #blocked2> + %cst_0 = arith.constant dense<10> : tensor<1x16xi32, #blocked3> +@@ -862,7 +862,7 @@ func public @mnist(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt. + #blocked5 = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [2, 16], warpsPerCTA = [8, 1], order = [1, 0]}> + // cmpf and cmpi have different operands and result types + // CHECK-LABEL: cmp +-func public @cmp(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: !tt.ptr {tt.divisibility = 16 : i32}, %arg4: i32 {tt.divisibility = 16 : i32}, %arg5: i32 {tt.divisibility = 16 : i32}) { ++func.func public @cmp(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: !tt.ptr {tt.divisibility = 16 : i32}, %arg4: i32 {tt.divisibility = 16 : i32}, %arg5: i32 {tt.divisibility = 16 : i32}) { + %c64 = arith.constant 64 : index + %c2048 = arith.constant 2048 : index + %c0 = arith.constant 0 : index +diff --git a/test/TritonGPU/loop-pipeline.mlir b/test/TritonGPU/loop-pipeline.mlir +index 6ee3b15fbc..663f2da7b0 100644 +--- a/test/TritonGPU/loop-pipeline.mlir ++++ b/test/TritonGPU/loop-pipeline.mlir +@@ -10,7 +10,7 @@ + #A = #triton_gpu.dot_op<{opIdx = 0, parent = #C}> + #B = #triton_gpu.dot_op<{opIdx = 1, parent = #C}> + +-// CHECK: func @matmul_loop ++// CHECK: func.func @matmul_loop + // CHECK-DAG: %[[CONSTANT_0:.*]] = arith.constant 0 : i32 + // CHECK-DAG: %[[CONSTANT_1:.*]] = arith.constant 1 : i32 + // CHECK-DAG: %[[CONSTANT_2:.*]] = arith.constant 2 : i32 +@@ -46,8 +46,8 @@ + // CHECK-DAG: %[[NEXT_PIPELINE_IDX:.*]] = arith.addi %[[PIPELINE_IDX]], %[[CONSTANT_1]] + // CHECK-DAG: %[[NEXT_LOOP_IDX:.*]] = arith.addi %[[LOOP_IDX]], %[[CONSTANT_1]] + // CHECK: scf.yield {{.*}}, {{.*}}, {{.*}}, %[[NEXT_A_BUFFER]], %[[NEXT_B_BUFFER]], %[[NEXT_A]], %[[NEXT_B]], {{.*}}, {{.*}}, {{.*}}, %[[NEXT_PIPELINE_IDX]], %[[NEXT_LOOP_IDX]] +-func @matmul_loop(%lb : index, %ub : index, %step : index, +- %A : !tt.ptr {tt.divisibility = 16 : i32}, ++func.func @matmul_loop(%lb : index, %ub : index, %step : index, ++ %A : !tt.ptr {tt.divisibility = 16 : i32}, + %B : !tt.ptr {tt.divisibility = 16 : i32}) { + // A ptrs + %a_ptr_splat = tt.splat %A : (!tt.ptr) -> tensor<128x32x!tt.ptr, #AL> +@@ -61,7 +61,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, + %b_tmp1 = tt.expand_dims %b_tmp0 {axis = 0 : i32} : (tensor<128xi32, #BLs0>) -> tensor<1x128xi32, #BL> + %b_offs = tt.broadcast %b_tmp1 : (tensor<1x128xi32, #BL>) -> tensor<32x128xi32, #BL> + %b_ptr_init = tt.addptr %b_ptr_splat, %b_offs : tensor<32x128x!tt.ptr, #BL>, tensor<32x128xi32, #BL> +- ++ + + %a_mask = arith.constant dense : tensor<128x32xi1, #AL> + %a_other = arith.constant dense<0.00e+00> : tensor<128x32xf16, #AL> +@@ -88,7 +88,7 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, + } + + +-// CHECK: func @matmul_loop_nested ++// CHECK: func.func @matmul_loop_nested + // CHECK-DAG: %[[CONSTANT_0:.*]] = arith.constant 0 : i32 + // CHECK-DAG: %[[CONSTANT_1:.*]] = arith.constant 1 : i32 + // CHECK-DAG: %[[CONSTANT_2:.*]] = arith.constant 2 : i32 +@@ -118,8 +118,8 @@ func @matmul_loop(%lb : index, %ub : index, %step : index, + // CHECK-DAG: %[[NEXT_PIPELINE_IDX:.*]] = arith.addi %[[PIPELINE_IDX]], %[[CONSTANT_1]] + // CHECK-DAG: %[[NEXT_LOOP_IDX:.*]] = arith.addi %[[LOOP_IDX]], %[[CONSTANT_1]] + // CHECK: scf.yield {{.*}}, {{.*}}, {{.*}}, %[[NEXT_A_BUFFER]], %[[NEXT_B_BUFFER]], %[[NEXT_A]], %[[NEXT_B]], {{.*}}, {{.*}}, {{.*}}, %[[NEXT_PIPELINE_IDX]], %[[NEXT_LOOP_IDX]] +-func @matmul_loop_nested(%lb : index, %ub : index, %step : index, +- %A : !tt.ptr {tt.divisibility = 16 : i32}, ++func.func @matmul_loop_nested(%lb : index, %ub : index, %step : index, ++ %A : !tt.ptr {tt.divisibility = 16 : i32}, + %B : !tt.ptr {tt.divisibility = 16 : i32}) { + scf.for %iv0 = %lb to %ub step %step { + // A ptrs +@@ -134,7 +134,7 @@ func @matmul_loop_nested(%lb : index, %ub : index, %step : index, + %b_tmp1 = tt.expand_dims %b_tmp0 {axis = 0 : i32} : (tensor<128xi32, #BLs0>) -> tensor<1x128xi32, #BL> + %b_offs = tt.broadcast %b_tmp1 : (tensor<1x128xi32, #BL>) -> tensor<32x128xi32, #BL> + %b_ptr_init = tt.addptr %b_ptr_splat, %b_offs : tensor<32x128x!tt.ptr, #BL>, tensor<32x128xi32, #BL> +- ++ + %a_mask = arith.constant dense : tensor<128x32xi1, #AL> + %a_other = arith.constant dense<0.00e+00> : tensor<128x32xf16, #AL> + %b_mask = arith.constant dense : tensor<32x128xi1, #BL> +@@ -161,7 +161,7 @@ func @matmul_loop_nested(%lb : index, %ub : index, %step : index, + } + + +-// CHECK: func @matmul_loop_single_pipeline ++// CHECK: func.func @matmul_loop_single_pipeline + // CHECK-DAG: %[[CONSTANT_0:.*]] = arith.constant 0 : i32 + // CHECK-DAG: %[[CONSTANT_1:.*]] = arith.constant 1 : i32 + // CHECK-DAG: %[[CONSTANT_2:.*]] = arith.constant 2 : i32 +@@ -183,8 +183,8 @@ func @matmul_loop_nested(%lb : index, %ub : index, %step : index, + // CHECK-DAG: %[[NEXT_PIPELINE_IDX:.*]] = arith.addi %[[PIPELINE_IDX]], %[[CONSTANT_1]] + // CHECK-DAG: %[[NEXT_LOOP_IDX:.*]] = arith.addi %[[LOOP_IDX]], %[[CONSTANT_1]] + // CHECK: scf.yield {{.*}}, {{.*}}, %[[NEXT_B_BUFFER]], %[[NEXT_B]], {{.*}}, {{.*}}, %[[NEXT_PIPELINE_IDX]], %[[NEXT_LOOP_IDX]] +-func @matmul_loop_single_pipeline(%lb : index, %ub : index, %step : index, +- %A : !tt.ptr {tt.divisibility = 16 : i32}, ++func.func @matmul_loop_single_pipeline(%lb : index, %ub : index, %step : index, ++ %A : !tt.ptr {tt.divisibility = 16 : i32}, + %B : !tt.ptr {tt.divisibility = 16 : i32}) { + // A ptrs + %a_ptr_splat = tt.splat %A : (!tt.ptr) -> tensor<128x32x!tt.ptr, #AL> +diff --git a/test/TritonGPU/matmul.mlir b/test/TritonGPU/matmul.mlir +index 9bd5318e1e..01dc3f0ab1 100644 +--- a/test/TritonGPU/matmul.mlir ++++ b/test/TritonGPU/matmul.mlir +@@ -4,7 +4,7 @@ + // CHECK: offset = 49152, size = 49152 + // CHECK: size = 98304 + module { +-func @matmul_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32_i32_i32_i32_i32_i32_i32__12c64_13c64_14c64_15c8(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32, %arg4: i32, %arg5: i32, %arg6: i32 {tt.divisibility = 16 : i32}, %arg7: i32, %arg8: i32 {tt.divisibility = 16 : i32}, %arg9: i32, %arg10: i32 {tt.divisibility = 16 : i32}, %arg11: i32) { ++func.func @matmul_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32_i32_i32_i32_i32_i32_i32__12c64_13c64_14c64_15c8(%arg0: !tt.ptr {tt.divisibility = 16 : i32}, %arg1: !tt.ptr {tt.divisibility = 16 : i32}, %arg2: !tt.ptr {tt.divisibility = 16 : i32}, %arg3: i32, %arg4: i32, %arg5: i32, %arg6: i32 {tt.divisibility = 16 : i32}, %arg7: i32, %arg8: i32 {tt.divisibility = 16 : i32}, %arg9: i32, %arg10: i32 {tt.divisibility = 16 : i32}, %arg11: i32) { + %cst = arith.constant dense : tensor<64x64xi1> + %c64 = arith.constant 64 : index + %c0 = arith.constant 0 : index +@@ -22,7 +22,7 @@ func @matmul_kernel__Pfp32_Pfp32_Pfp32_i32_i32_i32_i32_i32_i32_i32_i32_i32__12c6 + %7 = arith.muli %6, %c8_i32 : i32 + %8 = arith.subi %2, %7 : i32 + %9 = arith.cmpi slt, %8, %c8_i32 : i32 +- %10 = select %9, %8, %c8_i32 : i32 ++ %10 = arith.select %9, %8, %c8_i32 : i32 + %11 = arith.remsi %0, %10 : i32 + %12 = arith.addi %7, %11 : i32 + %13 = arith.remsi %0, %5 : i32 +diff --git a/test/TritonGPU/prefetch.mlir b/test/TritonGPU/prefetch.mlir +index 52b4dddec1..b427547890 100644 +--- a/test/TritonGPU/prefetch.mlir ++++ b/test/TritonGPU/prefetch.mlir +@@ -11,7 +11,7 @@ + #B_OP = #triton_gpu.dot_op<{opIdx = 1, parent = #C}> + + +-// CHECK: func @matmul_loop ++// CHECK: func.func @matmul_loop + // CHECK-DAG: %[[A0_PREFETCH_SMEM:.*]] = tensor.extract_slice %[[A0:.*]][0, 0] [128, 16] + // CHECK-DAG: %[[A0_PREFETCH:.*]] = triton_gpu.convert_layout %[[A0_PREFETCH_SMEM]] + // CHECK-DAG: %[[B0_PREFETCH_SMEM:.*]] = tensor.extract_slice %[[B0:.*]][0, 0] [16, 128] +@@ -28,7 +28,7 @@ + // CHECK-DAG: %[[NEXT_B_PREFETCH_SMEM:.*]] = tensor.extract_slice {{.*}}[0, 0] [16, 128] + // CHECK-DAG: %[[NEXT_B_PREFETCH:.*]] = triton_gpu.convert_layout %[[NEXT_B_PREFETCH_SMEM]] + // CHECK: scf.yield {{.*}}, {{.*}}, {{.*}}, {{.*}}, {{.*}}, %[[NEXT_A_PREFETCH]], %[[NEXT_B_PREFETCH]] +-func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { ++func.func @matmul_loop(%lb : index, %ub : index, %step : index, %A : !tt.ptr, %B : !tt.ptr) { + %a_ptr_init = tt.broadcast %A : (!tt.ptr) -> tensor<128x32x!tt.ptr, #AL> + %b_ptr_init = tt.broadcast %B : (!tt.ptr) -> tensor<32x128x!tt.ptr, #BL> + +diff --git a/test/TritonGPU/update-mma-for-volta.mlir b/test/TritonGPU/update-mma-for-volta.mlir +index d587fffcca..7571ec6185 100644 +--- a/test/TritonGPU/update-mma-for-volta.mlir ++++ b/test/TritonGPU/update-mma-for-volta.mlir +@@ -15,7 +15,7 @@ + // CHECK: [[new_mma:#mma.*]] = #triton_gpu.mma<{versionMajor = 1, versionMinor = 3, warpsPerCTA = [4, 2]}> + module attributes {"triton_gpu.num-warps" = 16 : i32} { + // CHECK-LABEL: dot_mmav1 +- func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> { ++ func.func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> { + %C = arith.constant dense<0.000000e+00> : tensor<64x64xf32, #blocked0> + %AA = triton_gpu.convert_layout %A : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_a> + %BB = triton_gpu.convert_layout %B : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_b> +@@ -50,7 +50,7 @@ module attributes {"triton_gpu.num-warps" = 16 : i32} { + + module attributes {"triton_gpu.num-warps" = 16 : i32} { + // CHECK-LABEL: dot_mmav1 +- func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> { ++ func.func @dot_mmav1(%A: tensor<64x64xf16, #blocked0>, %B: tensor<64x64xf16, #blocked0>) -> tensor<64x64xf32, #blocked0> { + %C = arith.constant dense<0.000000e+00> : tensor<64x64xf32, #blocked0> + %AA = triton_gpu.convert_layout %A : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_a> + %BB = triton_gpu.convert_layout %B : (tensor<64x64xf16, #blocked0>) -> tensor<64x64xf16, #dot_operand_b> +diff --git a/test/lib/Analysis/TestAlias.cpp b/test/lib/Analysis/TestAlias.cpp +index 88a4118fe9..3fd0cfd0d3 100644 +--- a/test/lib/Analysis/TestAlias.cpp ++++ b/test/lib/Analysis/TestAlias.cpp +@@ -9,10 +9,10 @@ using namespace mlir; + namespace { + + struct TestAliasPass +- : public PassWrapper> { ++ : public PassWrapper> { ++ ++ MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAliasPass); + +- // LLVM15+ +- // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAliasPass); + static void print(StringRef name, SmallVector &vals, + raw_ostream &os) { + if (vals.empty()) +@@ -39,23 +39,24 @@ struct TestAliasPass + auto opName = SymbolTable::getSymbolName(operation).getValue().str(); + os << opName << "\n"; + +- SharedMemoryAliasAnalysis analysis(&getContext()); +- analysis.run(operation); ++ std::unique_ptr solver = createDataFlowSolver(); ++ SharedMemoryAliasAnalysis *analysis = ++ solver->load(); ++ if (failed(solver->initializeAndRun(operation))) ++ return signalPassFailure(); + + AsmState state(operation->getParentOfType()); + // Get operation ids of value's aliases + auto getAllocOpNames = [&](Value value) { +- LatticeElement *latticeElement = +- analysis.lookupLatticeElement(value); ++ dataflow::Lattice *latticeElement = ++ analysis->getLatticeElement(value); + SmallVector opNames; +- if (latticeElement) { ++ if (latticeElement && !latticeElement->isUninitialized()) { + auto &info = latticeElement->getValue(); +- if (!info.getAllocs().empty()) { +- for (auto &alias : info.getAllocs()) { +- auto opName = +- getValueOperandName(alias.getDefiningOp()->getResult(0), state); +- opNames.push_back(std::move(opName)); +- } ++ for (auto &alias : info.getAllocs()) { ++ auto opName = ++ getValueOperandName(alias.getDefiningOp()->getResult(0), state); ++ opNames.push_back(std::move(opName)); + } + } + // Ensure deterministic output +diff --git a/test/lib/Analysis/TestAllocation.cpp b/test/lib/Analysis/TestAllocation.cpp +index 84108c4d36..35e42242bd 100644 +--- a/test/lib/Analysis/TestAllocation.cpp ++++ b/test/lib/Analysis/TestAllocation.cpp +@@ -6,10 +6,9 @@ using namespace mlir; + namespace { + + struct TestAllocationPass +- : public PassWrapper> { ++ : public PassWrapper> { + +- // LLVM15+ +- // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAllocationPass); ++ MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAllocationPass); + + StringRef getArgument() const final { return "test-print-allocation"; } + StringRef getDescription() const final { +diff --git a/test/lib/Analysis/TestAxisInfo.cpp b/test/lib/Analysis/TestAxisInfo.cpp +index a5205bb0a0..22347c32f0 100644 +--- a/test/lib/Analysis/TestAxisInfo.cpp ++++ b/test/lib/Analysis/TestAxisInfo.cpp +@@ -1,25 +1,15 @@ + #include "mlir/Pass/Pass.h" + #include "triton/Analysis/AxisInfo.h" ++#include "triton/Analysis/Utility.h" + + using namespace mlir; + + namespace { + + struct TestAxisInfoPass +- : public PassWrapper> { ++ : public PassWrapper> { + +- // LLVM15+ +- // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAlignmentPass); +- +- void print(const std::string &name, raw_ostream &os, ArrayRef vals) { +- os << name << ": ["; +- for (size_t d = 0; d < vals.size(); d++) { +- if (d != 0) +- os << ", "; +- os << vals[d]; +- } +- os << "]"; +- } ++ MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestAxisInfoPass); + + StringRef getArgument() const final { return "test-print-alignment"; } + StringRef getDescription() const final { +@@ -30,38 +20,19 @@ struct TestAxisInfoPass + Operation *operation = getOperation(); + auto &os = llvm::errs(); + auto opName = SymbolTable::getSymbolName(operation).getValue().str(); +- os << opName << "\n"; +- AxisInfoAnalysis analysis(&getContext()); +- analysis.run(operation); ++ os << "@" << opName << "\n"; ++ ++ std::unique_ptr solver = createDataFlowSolver(); ++ AxisInfoAnalysis *analysis = solver->load(); ++ if (failed(solver->initializeAndRun(operation))) ++ return signalPassFailure(); + operation->walk([&](Operation *op) { + if (op->getNumResults() < 1) + return; + for (Value result : op->getResults()) { +- // std::ostringstream oss; +- // result.print(oss); +- // os << " => "; +- LatticeElement *latticeElement = +- analysis.lookupLatticeElement(result); +- if (!latticeElement) { +- os << "None\n"; +- return; +- } +- AxisInfo &info = latticeElement->getValue(); +- print("Contiguity", os, info.getContiguity()); +- os << " ; "; +- print("Divisibility", os, info.getDivisibility()); +- os << " ; "; +- print("Constancy", os, info.getConstancy()); +- os << " ; "; +- auto constantValue = info.getConstantValue(); +- os << "ConstantValue: ["; +- if (constantValue.has_value()) +- os << constantValue.value(); +- else +- os << "None"; +- os << "] ( "; + result.print(os); +- os << " ) "; ++ os << " => "; ++ analysis->getLatticeElement(result)->getValue().print(os); + os << "\n"; + } + }); +diff --git a/test/lib/Analysis/TestMembar.cpp b/test/lib/Analysis/TestMembar.cpp +index df4279fe24..ab9b9f3fb7 100644 +--- a/test/lib/Analysis/TestMembar.cpp ++++ b/test/lib/Analysis/TestMembar.cpp +@@ -1,4 +1,4 @@ +-#include "mlir/Dialect/GPU/GPUDialect.h" ++#include "mlir/Dialect/GPU/IR/GPUDialect.h" + #include "mlir/IR/Dialect.h" + #include "mlir/Pass/Pass.h" + #include "triton/Analysis/Allocation.h" +@@ -9,10 +9,9 @@ using namespace mlir; + namespace { + + struct TestMembarPass +- : public PassWrapper> { ++ : public PassWrapper> { + +- // LLVM15+ +- // MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestMembarPass); ++ MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestMembarPass); + + StringRef getArgument() const final { return "test-print-membar"; } + StringRef getDescription() const final { diff --git a/pkgs/development/python-modules/torch/default.nix b/pkgs/development/python-modules/torch/default.nix index 9f701493f07..95cd7955053 100644 --- a/pkgs/development/python-modules/torch/default.nix +++ b/pkgs/development/python-modules/torch/default.nix @@ -6,6 +6,7 @@ # Native build inputs cmake, util-linux, linkFarm, symlinkJoin, which, pybind11, removeReferencesTo, + pythonRelaxDepsHook, # Build inputs numactl, @@ -13,9 +14,10 @@ # Propagated build inputs filelock, - sympy, - networkx, jinja2, + networkx, + openai-triton, + sympy, numpy, pyyaml, cffi, click, typing-extensions, # Unit tests @@ -271,6 +273,7 @@ in buildPythonPackage rec { which ninja pybind11 + pythonRelaxDepsHook removeReferencesTo ] ++ lib.optionals cudaSupport [ cudatoolkit_joined ] ++ lib.optionals rocmSupport [ rocmtoolkit_joined ]; @@ -298,8 +301,17 @@ in buildPythonPackage rec { # the following are required for tensorboard support pillow six future tensorboard protobuf - ] ++ lib.optionals MPISupport [ mpi ] - ++ lib.optionals rocmSupport [ rocmtoolkit_joined ]; + ] + ++ lib.optionals MPISupport [ mpi ] + ++ lib.optionals rocmSupport [ rocmtoolkit_joined ] + # rocm build requires openai-triton; + # openai-triton currently requires cuda_nvcc, + # so not including it in the cpu-only build; + # torch.compile relies on openai-triton, + # so we include it for the cuda build as well + ++ lib.optionals (rocmSupport || cudaSupport) [ + openai-triton + ]; # Tests take a long time and may be flaky, so just sanity-check imports doCheck = false; @@ -327,6 +339,11 @@ in buildPythonPackage rec { "runHook postCheck" ]; + pythonRemoveDeps = [ + # In our dist-info the name is just "triton" + "pytorch-triton-rocm" + ]; + postInstall = '' find "$out/${python.sitePackages}/torch/include" "$out/${python.sitePackages}/torch/lib" -type f -exec remove-references-to -t ${stdenv.cc} '{}' + diff --git a/pkgs/tools/audio/tts/default.nix b/pkgs/tools/audio/tts/default.nix index 1c3cc91616c..9887353665e 100644 --- a/pkgs/tools/audio/tts/default.nix +++ b/pkgs/tools/audio/tts/default.nix @@ -184,6 +184,8 @@ python.pkgs.buildPythonApplication rec { "tests/vocoder_tests/test_multiband_melgan_train.py" "tests/vocoder_tests/test_melgan_train.py" "tests/vocoder_tests/test_wavernn_train.py" + # only a feed forward test, but still takes too long + "tests/tts_tests/test_overflow.py" ]; passthru = { diff --git a/pkgs/top-level/python-packages.nix b/pkgs/top-level/python-packages.nix index 5349c922a4e..cc8d2da989e 100644 --- a/pkgs/top-level/python-packages.nix +++ b/pkgs/top-level/python-packages.nix @@ -6801,6 +6801,8 @@ self: super: with self; { open-meteo = callPackage ../development/python-modules/open-meteo { }; + openai-triton = callPackage ../development/python-modules/openai-triton { llvmPackages = pkgs.llvmPackages_rocm; }; + openai-whisper = callPackage ../development/python-modules/openai-whisper { }; openant = callPackage ../development/python-modules/openant { }; From 5e8008a536aeb17f3b90a9926e59d217591697e7 Mon Sep 17 00:00:00 2001 From: Someone Serge Date: Wed, 22 Mar 2023 16:20:14 +0200 Subject: [PATCH 16/32] python3Packages.torchWithCuda: avoid "unknown-warning" when building with cuda-compatible stdenv --- pkgs/development/python-modules/torch/default.nix | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pkgs/development/python-modules/torch/default.nix b/pkgs/development/python-modules/torch/default.nix index 95cd7955053..d9cf4e5760d 100644 --- a/pkgs/development/python-modules/torch/default.nix +++ b/pkgs/development/python-modules/torch/default.nix @@ -256,7 +256,7 @@ in buildPythonPackage rec { # Suppress gcc regression: avx512 math function raises uninitialized variable warning # https://gcc.gnu.org/bugzilla/show_bug.cgi?id=105593 # See also: Fails to compile with GCC 12.1.0 https://github.com/pytorch/pytorch/issues/77939 - ++ lib.optionals stdenv.cc.isGNU [ + ++ lib.optionals (stdenv.cc.isGNU && lib.versionAtLeast stdenv.cc.version "12.0.0") [ "-Wno-error=maybe-uninitialized" "-Wno-error=uninitialized" ] From 455d23b95d67f1774a59eabdbaafa9650ca2f61f Mon Sep 17 00:00:00 2001 From: Someone Serge Date: Thu, 23 Mar 2023 20:35:01 +0200 Subject: [PATCH 17/32] python3Packages.torchinfo: 1.64 -> 1.7.2 catch up with pytorch 2.0.0 and updated interfaces --- pkgs/development/python-modules/torchinfo/default.nix | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/pkgs/development/python-modules/torchinfo/default.nix b/pkgs/development/python-modules/torchinfo/default.nix index c36372235e7..67945d9bd2f 100644 --- a/pkgs/development/python-modules/torchinfo/default.nix +++ b/pkgs/development/python-modules/torchinfo/default.nix @@ -9,7 +9,7 @@ buildPythonPackage rec { pname = "torchinfo"; - version = "1.64"; + version = "1.7.2"; format = "setuptools"; disabled = pythonOlder "3.7"; @@ -18,7 +18,7 @@ buildPythonPackage rec { owner = "TylerYep"; repo = pname; rev = "refs/tags/v${version}"; - hash = "sha256-gcl8RxCD017FP4LtB60WVtOh7jg2Otv/vNd9hKneEAU="; + hash = "sha256-O+I7BNQ5moV/ZcbbuP/IFoi0LO0WsGHBbSfgPmFu1Ec="; }; propagatedBuildInputs = [ @@ -37,6 +37,11 @@ buildPythonPackage rec { "test_google" ]; + disabledTestPaths = [ + # Wants "compressai", which we don't package (2023-03-23) + "tests/torchinfo_xl_test.py" + ]; + pythonImportsCheck = [ "torchvision" ]; From 9b5fb1838d85f74a4c7f5b2362cec1570ba682e3 Mon Sep 17 00:00:00 2001 From: Someone Serge Date: Thu, 23 Mar 2023 20:35:15 +0200 Subject: [PATCH 18/32] python3Packages.torchinfo: fix pythonImportsCheck --- pkgs/development/python-modules/torchinfo/default.nix | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pkgs/development/python-modules/torchinfo/default.nix b/pkgs/development/python-modules/torchinfo/default.nix index 67945d9bd2f..420a5fd8dfc 100644 --- a/pkgs/development/python-modules/torchinfo/default.nix +++ b/pkgs/development/python-modules/torchinfo/default.nix @@ -43,7 +43,7 @@ buildPythonPackage rec { ]; pythonImportsCheck = [ - "torchvision" + "torchinfo" ]; meta = with lib; { From 91f24957262e76a6cff959ce49867703349d839d Mon Sep 17 00:00:00 2001 From: Someone Serge Date: Thu, 23 Mar 2023 22:27:56 +0200 Subject: [PATCH 19/32] ocamlPackages.torch: patch for pytorch 2.0.0 compatibility --- pkgs/development/ocaml-modules/torch/default.nix | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/pkgs/development/ocaml-modules/torch/default.nix b/pkgs/development/ocaml-modules/torch/default.nix index e46374ee1c2..9ba356fa939 100644 --- a/pkgs/development/ocaml-modules/torch/default.nix +++ b/pkgs/development/ocaml-modules/torch/default.nix @@ -2,6 +2,7 @@ , stdenv , buildDunePackage , fetchFromGitHub +, fetchpatch , cmdliner , ctypes , dune-configurator @@ -24,11 +25,19 @@ buildDunePackage rec { src = fetchFromGitHub { owner = "LaurentMazare"; - repo = "ocaml-${pname}"; - rev = version; + repo = "ocaml-${pname}"; + rev = version; hash = "sha256-z/9NUBjeFWE63Z/e8OyzDiy8hrn6qzjaiBH8G9MPeos="; }; + patches = [ + # Pytorch 2.0 support. Drop when it reaches a release + (fetchpatch { + url = "https://github.com/LaurentMazare/ocaml-torch/commit/ef7ef30cafecb09e45ec1ed8ce4bedae5947cfa5.patch"; + hash = "sha256-smdwKy40iIISp/25L2J4az6KmqFS1soeChBElUyhl5A="; + }) + ]; + buildInputs = [ dune-configurator ]; propagatedBuildInputs = [ From 24d20fefbf967ce83126098a1c788ca9866c4570 Mon Sep 17 00:00:00 2001 From: Someone Date: Mon, 3 Apr 2023 16:21:15 +0000 Subject: [PATCH 20/32] python3Packages.openai-triton: inline bash comments The drawback of this is that the comments now affect outPath's. Hopefully, though, we'll remove this preFixup soon anyway Co-authored-by: Sandro --- .../python-modules/openai-triton/default.nix | 37 ++++++++----------- 1 file changed, 16 insertions(+), 21 deletions(-) diff --git a/pkgs/development/python-modules/openai-triton/default.nix b/pkgs/development/python-modules/openai-triton/default.nix index 9340aad3a95..90c973d63b6 100644 --- a/pkgs/development/python-modules/openai-triton/default.nix +++ b/pkgs/development/python-modules/openai-triton/default.nix @@ -163,31 +163,26 @@ buildPythonPackage { ]; # Avoid GLIBCXX mismatch with other cuda-enabled python packages - preConfigure = - '' - export CC="${backendStdenv.cc}/bin/cc"; - export CXX="${backendStdenv.cc}/bin/c++"; - '' + preConfigure = '' + export CC="${backendStdenv.cc}/bin/cc"; + export CXX="${backendStdenv.cc}/bin/c++"; + # Upstream's setup.py tries to write cache somewhere in ~/ - + '' - export HOME=$TMPDIR - '' + export HOME=$TMPDIR + # Upstream's github actions patch setup.cfg to write base-dir. May be redundant - + '' - echo "" >> python/setup.cfg - echo "[build_ext]" >> python/setup.cfg - echo "base-dir=$PWD" >> python/setup.cfg - '' + echo " + [build_ext] + base-dir=$PWD" >> python/setup.cfg + # The rest (including buildPhase) is relative to ./python/ - + '' - cd python/ - '' + cd python/ + # Work around download_and_copy_ptxas() - + '' - dst_cuda="$PWD/triton/third_party/cuda/bin" - mkdir -p "$dst_cuda" - ln -s "${ptxas}" "$dst_cuda/" - ''; + dst_cuda="$PWD/triton/third_party/cuda/bin" + mkdir -p "$dst_cuda" + ln -s "${ptxas}" "$dst_cuda/" + ''; # CMake is run by setup.py instead dontUseCmakeConfigure = true; From 632cff6fcef0895202515d875dc89cc41fc14a8c Mon Sep 17 00:00:00 2001 From: Someone Serge Date: Tue, 4 Apr 2023 14:41:31 +0300 Subject: [PATCH 21/32] python3Packages.openai-triton: justify the use of pkgsTargetTarget --- .../python-modules/openai-triton/default.nix | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/pkgs/development/python-modules/openai-triton/default.nix b/pkgs/development/python-modules/openai-triton/default.nix index 90c973d63b6..0e10642f069 100644 --- a/pkgs/development/python-modules/openai-triton/default.nix +++ b/pkgs/development/python-modules/openai-triton/default.nix @@ -26,6 +26,18 @@ let version = "2.0.0"; inherit (cudaPackages) cuda_cudart backendStdenv; + + # A time may come we'll want to be cross-friendly + # + # Short explanation: we need pkgsTargetTarget, because we use string + # interpolation instead of buildInputs. + # + # Long explanation: OpenAI/triton downloads and vendors a copy of NVidia's + # ptxas compiler. We're not running this ptxas on the build machine, but on + # the user's machine, i.e. our Target platform. The second "Target" in + # pkgsTargetTarget maybe doesn't matter, because ptxas compiles programs to + # be executed on the GPU. + # Cf. https://nixos.org/manual/nixpkgs/unstable/#sec-cross-infra ptxas = "${pkgsTargetTarget.cudaPackages.cuda_nvcc}/bin/ptxas"; llvm = (llvmPackages.llvm.override { From abebf789a3131d9527a497b326f09e0f12cf827d Mon Sep 17 00:00:00 2001 From: Colin Date: Sat, 8 Apr 2023 00:26:41 +0000 Subject: [PATCH 22/32] rapidfuzz-cpp: inherit default `doCheck` instead of setting it explicitly --- pkgs/development/libraries/rapidfuzz-cpp/default.nix | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) diff --git a/pkgs/development/libraries/rapidfuzz-cpp/default.nix b/pkgs/development/libraries/rapidfuzz-cpp/default.nix index 9cde99b69f0..3050d64a4fd 100644 --- a/pkgs/development/libraries/rapidfuzz-cpp/default.nix +++ b/pkgs/development/libraries/rapidfuzz-cpp/default.nix @@ -5,14 +5,14 @@ , catch2_3 }: -stdenv.mkDerivation rec { +stdenv.mkDerivation (finalAttrs: { pname = "rapidfuzz-cpp"; version = "1.10.4"; src = fetchFromGitHub { owner = "maxbachmann"; repo = "rapidfuzz-cpp"; - rev = "v${version}"; + rev = "v${finalAttrs.version}"; hash = "sha256-I7MdeLs+J5a57ypgUJIW0/pSFPzK4nZA4ZrVRdKX7J4="; }; @@ -20,7 +20,7 @@ stdenv.mkDerivation rec { cmake ]; - cmakeFlags = lib.optionals doCheck [ + cmakeFlags = lib.optionals finalAttrs.finalPackage.doCheck [ "-DRAPIDFUZZ_BUILD_TESTING=ON" ]; @@ -33,14 +33,12 @@ stdenv.mkDerivation rec { catch2_3 ]; - doCheck = stdenv.buildPlatform.canExecute stdenv.hostPlatform; - meta = { description = "Rapid fuzzy string matching in C++ using the Levenshtein Distance"; homepage = "https://github.com/maxbachmann/rapidfuzz-cpp"; - changelog = "https://github.com/maxbachmann/rapidfuzz-cpp/blob/${src.rev}/CHANGELOG.md"; + changelog = "https://github.com/maxbachmann/rapidfuzz-cpp/blob/${finalAttrs.src.rev}/CHANGELOG.md"; license = lib.licenses.mit; maintainers = with lib.maintainers; [ dotlambda ]; platforms = lib.platforms.unix; }; -} +}) From 09462c7f66f51d172955e6cee70a58f697979f41 Mon Sep 17 00:00:00 2001 From: ocfox Date: Sat, 8 Apr 2023 15:06:12 +0800 Subject: [PATCH 23/32] rustdesk: fix build --- .../networking/remote/rustdesk/Cargo.lock | 5193 +++++++++++++++++ .../networking/remote/rustdesk/cargo.patch | 35 - .../networking/remote/rustdesk/default.nix | 98 +- 3 files changed, 5252 insertions(+), 74 deletions(-) create mode 100644 pkgs/applications/networking/remote/rustdesk/Cargo.lock diff --git a/pkgs/applications/networking/remote/rustdesk/Cargo.lock b/pkgs/applications/networking/remote/rustdesk/Cargo.lock new file mode 100644 index 00000000000..656b0fd9b06 --- /dev/null +++ b/pkgs/applications/networking/remote/rustdesk/Cargo.lock @@ -0,0 +1,5193 @@ +# This file is automatically @generated by Cargo. +# It is not intended for manual editing. +version = 3 + +[[package]] +name = "addr2line" +version = "0.19.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a76fd60b23679b7d19bd066031410fb7e458ccc5e958eb5c325888ce4baedc97" +dependencies = [ + "gimli", +] + +[[package]] +name = "adler" +version = "1.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f26201604c87b1e01bd3d98f8d5d9a8fcbb815e8cedb41ffccbeb4bf593a35fe" + +[[package]] +name = "adler32" +version = "1.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "aae1277d39aeec15cb388266ecc24b11c80469deae6067e17a1a7aa9e5c1f234" + +[[package]] +name = "aho-corasick" +version = "0.7.20" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cc936419f96fa211c1b9166887b38e5e40b19958e5b895be7c1f93adec7071ac" +dependencies = [ + "memchr", +] + +[[package]] +name = "alsa" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5915f52fe2cf65e83924d037b6c5290b7cee097c6b5c8700746e6168a343fd6b" +dependencies = [ + "alsa-sys", + "bitflags", + "libc", + "nix 0.23.2", +] + +[[package]] +name = "alsa-sys" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "db8fee663d06c4e303404ef5f40488a53e062f89ba8bfed81f42325aafad1527" +dependencies = [ + "libc", + "pkg-config", +] + +[[package]] +name = "android_log-sys" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "85965b6739a430150bdd138e2374a98af0c3ee0d030b3bb7fc3bddff58d0102e" + +[[package]] +name = "android_logger" +version = "0.11.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8619b80c242aa7bd638b5c7ddd952addeecb71f69c75e33f1d47b2804f8f883a" +dependencies = [ + "android_log-sys", + "env_logger 0.10.0", + "log", + "once_cell", +] + +[[package]] +name = "android_system_properties" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "819e7219dbd41043ac279b19830f2efc897156490d7fd6ea916720117ee66311" +dependencies = [ + "libc", +] + +[[package]] +name = "ansi_term" +version = "0.12.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d52a9bb7ec0cf484c551830a7ce27bd20d67eac647e1befb56b0be4ee39a55d2" +dependencies = [ + "winapi 0.3.9", +] + +[[package]] +name = "anyhow" +version = "1.0.70" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7de8ce5e0f9f8d88245311066a578d72b7af3e7088f32783804676302df237e4" + +[[package]] +name = "arboard" +version = "2.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dc120354d1b5ec6d7aaf4876b602def75595937b5e15d356eb554ab5177e08bb" +dependencies = [ + "clipboard-win", + "core-graphics", + "image", + "log", + "objc", + "objc-foundation", + "objc_id", + "parking_lot 0.12.1", + "thiserror", + "winapi 0.3.9", + "x11rb", +] + +[[package]] +name = "async-trait" +version = "0.1.68" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b9ccdd8f2a161be9bd5c023df56f1b2a0bd1d83872ae53b71a84a12c9bf6e842" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.13", +] + +[[package]] +name = "atk" +version = "0.9.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "812b4911e210bd51b24596244523c856ca749e6223c50a7fbbba3f89ee37c426" +dependencies = [ + "atk-sys 0.10.0", + "bitflags", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", +] + +[[package]] +name = "atk" +version = "0.15.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2c3d816ce6f0e2909a96830d6911c2aff044370b1ef92d7f267b43bae5addedd" +dependencies = [ + "atk-sys 0.15.1", + "bitflags", + "glib 0.15.12", + "libc", +] + +[[package]] +name = "atk-sys" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f530e4af131d94cc4fa15c5c9d0348f0ef28bac64ba660b6b2a1cf2605dedfce" +dependencies = [ + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "atk-sys" +version = "0.15.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "58aeb089fb698e06db8089971c7ee317ab9644bade33383f63631437b03aafb6" +dependencies = [ + "glib-sys 0.15.10", + "gobject-sys 0.15.10", + "libc", + "system-deps 6.0.4", +] + +[[package]] +name = "atty" +version = "0.2.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d9b39be18770d11421cdb1b9947a45dd3f37e93092cbf377614828a319d5fee8" +dependencies = [ + "hermit-abi 0.1.19", + "libc", + "winapi 0.3.9", +] + +[[package]] +name = "autocfg" +version = "0.1.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0dde43e75fd43e8a1bf86103336bc699aa8d17ad1be60c76c0bdfd4828e19b78" +dependencies = [ + "autocfg 1.1.0", +] + +[[package]] +name = "autocfg" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d468802bab17cbc0cc575e9b053f41e72aa36bfa6b7f55e3529ffa43161b97fa" + +[[package]] +name = "backtrace" +version = "0.3.67" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "233d376d6d185f2a3093e58f283f60f880315b6c60075b01f36b3b85154564ca" +dependencies = [ + "addr2line", + "cc", + "cfg-if 1.0.0", + "libc", + "miniz_oxide 0.6.2", + "object", + "rustc-demangle", +] + +[[package]] +name = "base64" +version = "0.13.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9e1b586273c5702936fe7b7d6896644d8be71e6314cfe09d3167c95f712589e8" + +[[package]] +name = "base64" +version = "0.21.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a4a4ddaa51a5bc52a6948f74c06d20aaaddb71924eab79b8c97a8c556e942d6a" + +[[package]] +name = "bindgen" +version = "0.59.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2bd2a9a458e8f4304c52c43ebb0cfbd520289f8379a52e329a38afda99bf8eb8" +dependencies = [ + "bitflags", + "cexpr", + "clang-sys", + "clap 2.34.0", + "env_logger 0.9.3", + "lazy_static", + "lazycell", + "log", + "peeking_take_while", + "proc-macro2", + "quote", + "regex", + "rustc-hash", + "shlex", + "which 4.4.0", +] + +[[package]] +name = "bindgen" +version = "0.64.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c4243e6031260db77ede97ad86c27e501d646a27ab57b59a574f725d98ab1fb4" +dependencies = [ + "bitflags", + "cexpr", + "clang-sys", + "lazy_static", + "lazycell", + "peeking_take_while", + "proc-macro2", + "quote", + "regex", + "rustc-hash", + "shlex", + "syn 1.0.109", +] + +[[package]] +name = "bitflags" +version = "1.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bef38d45163c2f1dde094a7dfd33ccf595c92905c8f8f4fdc18d06fb1037718a" + +[[package]] +name = "block" +version = "0.1.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a" + +[[package]] +name = "block-buffer" +version = "0.10.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3078c7629b62d3f0439517fa394996acacc5cbc91c5a20d8c658e77abd503a71" +dependencies = [ + "generic-array", +] + +[[package]] +name = "bumpalo" +version = "3.12.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0d261e256854913907f67ed06efbc3338dfe6179796deefc1ff763fc1aee5535" + +[[package]] +name = "bytemuck" +version = "1.13.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "17febce684fd15d89027105661fec94afb475cb995fbc59d2865198446ba2eea" + +[[package]] +name = "byteorder" +version = "1.4.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "14c189c53d098945499cdfa7ecc63567cf3886b3332b312a5b4585d8d3a6a610" + +[[package]] +name = "bytes" +version = "1.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "89b2fd2a0dcf38d7971e2194b6b6eebab45ae01067456a7fd93d5547a61b70be" + +[[package]] +name = "cairo-rs" +version = "0.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c5c0f2e047e8ca53d0ff249c54ae047931d7a6ebe05d00af73e0ffeb6e34bdb8" +dependencies = [ + "bitflags", + "cairo-sys-rs 0.10.0", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "thiserror", +] + +[[package]] +name = "cairo-rs" +version = "0.15.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c76ee391b03d35510d9fa917357c7f1855bd9a6659c95a1b392e33f49b3369bc" +dependencies = [ + "bitflags", + "cairo-sys-rs 0.15.1", + "glib 0.15.12", + "libc", + "thiserror", +] + +[[package]] +name = "cairo-sys-rs" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2ed2639b9ad5f1d6efa76de95558e11339e7318426d84ac4890b86c03e828ca7" +dependencies = [ + "glib-sys 0.10.1", + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "cairo-sys-rs" +version = "0.15.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3c55d429bef56ac9172d25fecb85dc8068307d17acd74b377866b7a1ef25d3c8" +dependencies = [ + "glib-sys 0.15.10", + "libc", + "system-deps 6.0.4", +] + +[[package]] +name = "cc" +version = "1.0.79" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "50d30906286121d95be3d479533b458f87493b30a4b5f79a607db8f5d11aa91f" +dependencies = [ + "jobserver", +] + +[[package]] +name = "cesu8" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6d43a04d8753f35258c91f8ec639f792891f748a1edbd759cf1dcea3382ad83c" + +[[package]] +name = "cexpr" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6fac387a98bb7c37292057cffc56d62ecb629900026402633ae9160df93a8766" +dependencies = [ + "nom", +] + +[[package]] +name = "cfg-expr" +version = "0.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a35b255461940a32985c627ce82900867c61db1659764d3675ea81963f72a4c6" +dependencies = [ + "smallvec", +] + +[[package]] +name = "cfg-if" +version = "0.1.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4785bdd1c96b2a846b2bd7cc02e86b6b3dbf14e7e53446c4f54c92a361040822" + +[[package]] +name = "cfg-if" +version = "1.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "baf1de4339761588bc0619e3cbc0120ee582ebb74b53b4efbf79117bd2da40fd" + +[[package]] +name = "chrono" +version = "0.4.24" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4e3c5919066adf22df73762e50cffcde3a758f2a848b113b586d1f86728b673b" +dependencies = [ + "iana-time-zone", + "num-integer", + "num-traits 0.2.15", + "winapi 0.3.9", +] + +[[package]] +name = "clang-sys" +version = "1.6.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c688fc74432808e3eb684cae8830a86be1d66a2bd58e1f248ed0960a590baf6f" +dependencies = [ + "glob", + "libc", + "libloading", +] + +[[package]] +name = "clap" +version = "2.34.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a0610544180c38b88101fecf2dd634b174a62eef6946f84dfc6a7127512b381c" +dependencies = [ + "ansi_term", + "atty", + "bitflags", + "strsim 0.8.0", + "textwrap 0.11.0", + "unicode-width", + "vec_map", +] + +[[package]] +name = "clap" +version = "3.2.23" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "71655c45cb9845d3270c9d6df84ebe72b4dad3c2ba3f7023ad47c144e4e473a5" +dependencies = [ + "atty", + "bitflags", + "clap_lex", + "indexmap", + "strsim 0.10.0", + "termcolor", + "textwrap 0.16.0", +] + +[[package]] +name = "clap_lex" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2850f2f5a82cbf437dd5af4d49848fbdfc27c157c3d010345776f952765261c5" +dependencies = [ + "os_str_bytes", +] + +[[package]] +name = "clipboard" +version = "0.1.0" +dependencies = [ + "cc", + "hbb_common", + "lazy_static", + "serde 1.0.159", + "serde_derive", + "thiserror", +] + +[[package]] +name = "clipboard-master" +version = "3.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "459887701008d8ee21f8de7f45f0f0707417c7eea3311973f6e67222bd686b7a" +dependencies = [ + "objc", + "objc-foundation", + "objc_id", + "winapi 0.3.9", + "windows-win", + "x11-clipboard", +] + +[[package]] +name = "clipboard-win" +version = "4.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7191c27c2357d9b7ef96baac1773290d4ca63b24205b82a3fd8a0637afcf0362" +dependencies = [ + "error-code", + "str-buf", + "winapi 0.3.9", +] + +[[package]] +name = "cloudabi" +version = "0.0.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ddfc5b9aa5d4507acaf872de71051dfd0e309860e88966e1051e462a077aac4f" +dependencies = [ + "bitflags", +] + +[[package]] +name = "cmake" +version = "0.1.50" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a31c789563b815f77f4250caee12365734369f942439b7defd71e18a48197130" +dependencies = [ + "cc", +] + +[[package]] +name = "cocoa" +version = "0.24.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f425db7937052c684daec3bd6375c8abe2d146dca4b8b143d6db777c39138f3a" +dependencies = [ + "bitflags", + "block", + "cocoa-foundation", + "core-foundation", + "core-graphics", + "foreign-types", + "libc", + "objc", +] + +[[package]] +name = "cocoa-foundation" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "931d3837c286f56e3c58423ce4eba12d08db2374461a785c86f672b08b5650d6" +dependencies = [ + "bitflags", + "block", + "core-foundation", + "core-graphics-types", + "foreign-types", + "libc", + "objc", +] + +[[package]] +name = "codespan-reporting" +version = "0.11.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3538270d33cc669650c4b093848450d380def10c331d38c768e34cac80576e6e" +dependencies = [ + "termcolor", + "unicode-width", +] + +[[package]] +name = "color_quant" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3d7b894f5411737b7867f4827955924d7c254fc9f4d91a6aad6b097804b1018b" + +[[package]] +name = "combine" +version = "4.6.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "35ed6e9d84f0b51a7f52daf1c7d71dd136fd7a3f41a8462b8cdb8c78d920fad4" +dependencies = [ + "bytes", + "memchr", +] + +[[package]] +name = "confy" +version = "0.4.0" +source = "git+https://github.com/open-trade/confy#630cc28a396cb7d01eefdd9f3824486fe4d8554b" +dependencies = [ + "directories-next", + "serde 1.0.159", + "thiserror", + "toml 0.5.11", +] + +[[package]] +name = "core-foundation" +version = "0.9.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "194a7a9e6de53fa55116934067c844d9d749312f75c6f6d0980e8c252f8c2146" +dependencies = [ + "core-foundation-sys", + "libc", +] + +[[package]] +name = "core-foundation-sys" +version = "0.8.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e496a50fda8aacccc86d7529e2c1e0892dbd0f898a6b5645b5561b89c3210efa" + +[[package]] +name = "core-graphics" +version = "0.22.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2581bbab3b8ffc6fcbd550bf46c355135d16e9ff2a6ea032ad6b9bf1d7efe4fb" +dependencies = [ + "bitflags", + "core-foundation", + "core-graphics-types", + "foreign-types", + "libc", +] + +[[package]] +name = "core-graphics-types" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3a68b68b3446082644c91ac778bf50cd4104bfb002b5a6a7c44cca5a2c70788b" +dependencies = [ + "bitflags", + "core-foundation", + "foreign-types", + "libc", +] + +[[package]] +name = "coreaudio-rs" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "11894b20ebfe1ff903cbdc52259693389eea03b94918a2def2c30c3bf227ad88" +dependencies = [ + "bitflags", + "coreaudio-sys", +] + +[[package]] +name = "coreaudio-sys" +version = "0.2.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f034b2258e6c4ade2f73bf87b21047567fb913ee9550837c2316d139b0262b24" +dependencies = [ + "bindgen 0.64.0", +] + +[[package]] +name = "cpal" +version = "0.13.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "74117836a5124f3629e4b474eed03e479abaf98988b4bb317e29f08cfe0e4116" +dependencies = [ + "alsa", + "core-foundation-sys", + "coreaudio-rs", + "jni", + "js-sys", + "lazy_static", + "libc", + "mach", + "ndk", + "ndk-glue", + "nix 0.23.2", + "oboe", + "parking_lot 0.11.2", + "stdweb", + "thiserror", + "web-sys", + "winapi 0.3.9", +] + +[[package]] +name = "cpufeatures" +version = "0.2.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "280a9f2d8b3a38871a3c8a46fb80db65e5e5ed97da80c4d08bf27fb63e35e181" +dependencies = [ + "libc", +] + +[[package]] +name = "crc32fast" +version = "1.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b540bd8bc810d3885c6ea91e2018302f68baba2129ab3e88f32389ee9370880d" +dependencies = [ + "cfg-if 1.0.0", +] + +[[package]] +name = "crossbeam-channel" +version = "0.5.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cf2b3e8478797446514c91ef04bafcb59faba183e621ad488df88983cc14128c" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-deque" +version = "0.8.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ce6fd6f855243022dcecf8702fef0c297d4338e226845fe067f6341ad9fa0cef" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-epoch", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-epoch" +version = "0.9.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "46bd5f3f85273295a9d14aedfb86f6aadbff6d8f5295c4a9edb08e819dcf5695" +dependencies = [ + "autocfg 1.1.0", + "cfg-if 1.0.0", + "crossbeam-utils", + "memoffset 0.8.0", + "scopeguard", +] + +[[package]] +name = "crossbeam-queue" +version = "0.3.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d1cfb3ea8a53f37c40dea2c7bedcbd88bdfae54f5e2175d6ecaff1c988353add" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-utils" +version = "0.8.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3c063cd8cc95f5c377ed0d4b49a4b21f632396ff690e8470c29b3359b346984b" +dependencies = [ + "cfg-if 1.0.0", +] + +[[package]] +name = "crypto-common" +version = "0.1.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1bfb12502f3fc46cca1bb51ac28df9d618d813cdc3d2f25b9fe775a34af26bb3" +dependencies = [ + "generic-array", + "typenum", +] + +[[package]] +name = "ctrlc" +version = "3.2.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bbcf33c2a618cbe41ee43ae6e9f2e48368cd9f9db2896f10167d8d762679f639" +dependencies = [ + "nix 0.26.2", + "windows-sys 0.45.0", +] + +[[package]] +name = "cxx" +version = "1.0.94" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f61f1b6389c3fe1c316bf8a4dccc90a38208354b330925bce1f74a6c4756eb93" +dependencies = [ + "cc", + "cxxbridge-flags", + "cxxbridge-macro", + "link-cplusplus", +] + +[[package]] +name = "cxx-build" +version = "1.0.94" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "12cee708e8962df2aeb38f594aae5d827c022b6460ac71a7a3e2c3c2aae5a07b" +dependencies = [ + "cc", + "codespan-reporting", + "once_cell", + "proc-macro2", + "quote", + "scratch", + "syn 2.0.13", +] + +[[package]] +name = "cxxbridge-flags" +version = "1.0.94" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7944172ae7e4068c533afbb984114a56c46e9ccddda550499caa222902c7f7bb" + +[[package]] +name = "cxxbridge-macro" +version = "1.0.94" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2345488264226bf682893e25de0769f3360aac9957980ec49361b083ddaa5bc5" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.13", +] + +[[package]] +name = "darling" +version = "0.13.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a01d95850c592940db9b8194bc39f4bc0e89dee5c4265e4b1807c34a9aba453c" +dependencies = [ + "darling_core", + "darling_macro", +] + +[[package]] +name = "darling_core" +version = "0.13.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "859d65a907b6852c9361e3185c862aae7fafd2887876799fa55f5f99dc40d610" +dependencies = [ + "fnv", + "ident_case", + "proc-macro2", + "quote", + "strsim 0.10.0", + "syn 1.0.109", +] + +[[package]] +name = "darling_macro" +version = "0.13.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9c972679f83bdf9c42bd905396b6c3588a843a17f0f16dfcfa3e2c5d57441835" +dependencies = [ + "darling_core", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "dasp" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7381b67da416b639690ac77c73b86a7b5e64a29e31d1f75fb3b1102301ef355a" +dependencies = [ + "dasp_envelope", + "dasp_frame", + "dasp_interpolate", + "dasp_peak", + "dasp_ring_buffer", + "dasp_rms", + "dasp_sample", + "dasp_signal", + "dasp_slice", + "dasp_window", +] + +[[package]] +name = "dasp_envelope" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8ec617ce7016f101a87fe85ed44180839744265fae73bb4aa43e7ece1b7668b6" +dependencies = [ + "dasp_frame", + "dasp_peak", + "dasp_ring_buffer", + "dasp_rms", + "dasp_sample", +] + +[[package]] +name = "dasp_frame" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b2a3937f5fe2135702897535c8d4a5553f8b116f76c1529088797f2eee7c5cd6" +dependencies = [ + "dasp_sample", +] + +[[package]] +name = "dasp_interpolate" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7fc975a6563bb7ca7ec0a6c784ead49983a21c24835b0bc96eea11ee407c7486" +dependencies = [ + "dasp_frame", + "dasp_ring_buffer", + "dasp_sample", +] + +[[package]] +name = "dasp_peak" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5cf88559d79c21f3d8523d91250c397f9a15b5fc72fbb3f87fdb0a37b79915bf" +dependencies = [ + "dasp_frame", + "dasp_sample", +] + +[[package]] +name = "dasp_ring_buffer" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "07d79e19b89618a543c4adec9c5a347fe378a19041699b3278e616e387511ea1" + +[[package]] +name = "dasp_rms" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a6c5dcb30b7e5014486e2822537ea2beae50b19722ffe2ed7549ab03774575aa" +dependencies = [ + "dasp_frame", + "dasp_ring_buffer", + "dasp_sample", +] + +[[package]] +name = "dasp_sample" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0c87e182de0887fd5361989c677c4e8f5000cd9491d6d563161a8f3a5519fc7f" + +[[package]] +name = "dasp_signal" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "aa1ab7d01689c6ed4eae3d38fe1cea08cba761573fbd2d592528d55b421077e7" +dependencies = [ + "dasp_envelope", + "dasp_frame", + "dasp_interpolate", + "dasp_peak", + "dasp_ring_buffer", + "dasp_rms", + "dasp_sample", + "dasp_window", +] + +[[package]] +name = "dasp_slice" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4e1c7335d58e7baedafa516cb361360ff38d6f4d3f9d9d5ee2a2fc8e27178fa1" +dependencies = [ + "dasp_frame", + "dasp_sample", +] + +[[package]] +name = "dasp_window" +version = "0.11.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "99ded7b88821d2ce4e8b842c9f1c86ac911891ab89443cc1de750cae764c5076" +dependencies = [ + "dasp_sample", +] + +[[package]] +name = "dbus" +version = "0.9.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1bb21987b9fb1613058ba3843121dd18b163b254d8a6e797e144cbac14d96d1b" +dependencies = [ + "libc", + "libdbus-sys", + "winapi 0.3.9", +] + +[[package]] +name = "deflate" +version = "0.8.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "73770f8e1fe7d64df17ca66ad28994a0a623ea497fa69486e14984e715c5d174" +dependencies = [ + "adler32", + "byteorder", +] + +[[package]] +name = "digest" +version = "0.10.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8168378f4e5023e7218c89c891c0fd8ecdb5e5e4f18cb78f38cf245dd021e76f" +dependencies = [ + "block-buffer", + "crypto-common", +] + +[[package]] +name = "directories-next" +version = "2.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "339ee130d97a610ea5a5872d2bbb130fdf68884ff09d3028b81bec8a1ac23bbc" +dependencies = [ + "cfg-if 1.0.0", + "dirs-sys-next", +] + +[[package]] +name = "dirs-next" +version = "2.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b98cf8ebf19c3d1b223e151f99a4f9f0690dca41414773390fc824184ac833e1" +dependencies = [ + "cfg-if 1.0.0", + "dirs-sys-next", +] + +[[package]] +name = "dirs-sys-next" +version = "0.1.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4ebda144c4fe02d1f7ea1a7d9641b6fc6b580adcfa024ae48797ecdeb6825b4d" +dependencies = [ + "libc", + "redox_users", + "winapi 0.3.9", +] + +[[package]] +name = "dispatch" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bd0c93bb4b0c6d9b77f4435b0ae98c24d17f1c45b2ff844c6151a07256ca923b" + +[[package]] +name = "docopt" +version = "1.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7f3f119846c823f9eafcf953a8f6ffb6ed69bf6240883261a7f13b634579a51f" +dependencies = [ + "lazy_static", + "regex", + "serde 1.0.159", + "strsim 0.10.0", +] + +[[package]] +name = "dtoa" +version = "0.4.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "56899898ce76aaf4a0f24d914c97ea6ed976d42fec6ad33fcbb0a1103e07b2b0" + +[[package]] +name = "ed25519" +version = "1.5.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "91cff35c70bba8a626e3185d8cd48cc11b5437e1a5bcd15b9b5fa3c64b6dfee7" +dependencies = [ + "signature", +] + +[[package]] +name = "either" +version = "1.8.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7fcaabb2fef8c910e7f4c7ce9f67a1283a1715879a7c230ca9d6d1ae31f16d91" + +[[package]] +name = "enigo" +version = "0.0.14" +dependencies = [ + "core-graphics", + "libc", + "log", + "objc", + "pkg-config", + "serde 1.0.159", + "serde_derive", + "unicode-segmentation", + "winapi 0.3.9", +] + +[[package]] +name = "enum-map" +version = "2.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "988f0d17a0fa38291e5f41f71ea8d46a5d5497b9054d5a759fae2cbb819f2356" +dependencies = [ + "enum-map-derive", +] + +[[package]] +name = "enum-map-derive" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2a4da76b3b6116d758c7ba93f7ec6a35d2e2cf24feda76c6e38a375f4d5c59f2" +dependencies = [ + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "env_logger" +version = "0.9.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a12e6657c4c97ebab115a42dcee77225f7f482cdd841cf7088c657a42e9e00e7" +dependencies = [ + "atty", + "humantime", + "log", + "regex", + "termcolor", +] + +[[package]] +name = "env_logger" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "85cdab6a89accf66733ad5a1693a4dcced6aeff64602b634530dd73c1f3ee9f0" +dependencies = [ + "log", + "regex", +] + +[[package]] +name = "err-derive" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c34a887c8df3ed90498c1c437ce21f211c8e27672921a8ffa293cb8d6d4caa9e" +dependencies = [ + "proc-macro-error", + "proc-macro2", + "quote", + "rustversion", + "syn 1.0.109", + "synstructure", +] + +[[package]] +name = "errno" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4bcfec3a70f97c962c307b2d2c56e358cf1d00b558d74262b5f929ee8cc7e73a" +dependencies = [ + "errno-dragonfly", + "libc", + "windows-sys 0.48.0", +] + +[[package]] +name = "errno-dragonfly" +version = "0.1.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "aa68f1b12764fab894d2755d2518754e71b4fd80ecfb822714a1206c2aab39bf" +dependencies = [ + "cc", + "libc", +] + +[[package]] +name = "error-code" +version = "2.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "64f18991e7bf11e7ffee451b5318b5c1a73c52d0d0ada6e5a3017c8c1ced6a21" +dependencies = [ + "libc", + "str-buf", +] + +[[package]] +name = "failure" +version = "0.1.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d32e9bd16cc02eae7db7ef620b392808b89f6a5e16bb3497d159c6b92a0f4f86" +dependencies = [ + "backtrace", +] + +[[package]] +name = "fastrand" +version = "1.9.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e51093e27b0797c359783294ca4f0a911c270184cb10f85783b118614a1501be" +dependencies = [ + "instant", +] + +[[package]] +name = "field-offset" +version = "0.3.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a3cf3a800ff6e860c863ca6d4b16fd999db8b752819c1606884047b73e468535" +dependencies = [ + "memoffset 0.8.0", + "rustc_version", +] + +[[package]] +name = "filetime" +version = "0.2.21" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5cbc844cecaee9d4443931972e1289c8ff485cb4cc2767cb03ca139ed6885153" +dependencies = [ + "cfg-if 1.0.0", + "libc", + "redox_syscall 0.2.16", + "windows-sys 0.48.0", +] + +[[package]] +name = "flate2" +version = "1.0.25" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a8a2db397cb1c8772f31494cb8917e48cd1e64f0fa7efac59fbd741a0a8ce841" +dependencies = [ + "crc32fast", + "miniz_oxide 0.6.2", +] + +[[package]] +name = "flexi_logger" +version = "0.22.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0c76a80dd14a27fc3d8bc696502132cb52b3f227256fd8601166c3a35e45f409" +dependencies = [ + "ansi_term", + "atty", + "chrono", + "crossbeam-channel", + "crossbeam-queue", + "glob", + "lazy_static", + "log", + "regex", + "rustversion", + "thiserror", + "time", +] + +[[package]] +name = "fnv" +version = "1.0.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3f9eec918d3f24069decb9af1554cad7c880e2da24a9afd88aca000531ab82c1" + +[[package]] +name = "foreign-types" +version = "0.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f6f339eb8adc052cd2ca78910fda869aefa38d22d5cb648e6485e4d3fc06f3b1" +dependencies = [ + "foreign-types-shared", +] + +[[package]] +name = "foreign-types-shared" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "00b0228411908ca8685dba7fc2cdd70ec9990a6e753e89b6ac91a84c40fbaf4b" + +[[package]] +name = "fuchsia-cprng" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a06f77d526c1a601b7c4cdd98f54b5eaabffc14d5f2f0296febdc7f357c6d3ba" + +[[package]] +name = "fuchsia-zircon" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2e9763c69ebaae630ba35f74888db465e49e259ba1bc0eda7d06f4a067615d82" +dependencies = [ + "bitflags", + "fuchsia-zircon-sys", +] + +[[package]] +name = "fuchsia-zircon-sys" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3dcaa9ae7725d12cdb85b3ad99a434db70b468c09ded17e012d86b5c1010f7a7" + +[[package]] +name = "futures" +version = "0.3.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "23342abe12aba583913b2e62f22225ff9c950774065e4bfb61a19cd9770fec40" +dependencies = [ + "futures-channel", + "futures-core", + "futures-executor", + "futures-io", + "futures-sink", + "futures-task", + "futures-util", +] + +[[package]] +name = "futures-channel" +version = "0.3.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "955518d47e09b25bbebc7a18df10b81f0c766eaf4c4f1cccef2fca5f2a4fb5f2" +dependencies = [ + "futures-core", + "futures-sink", +] + +[[package]] +name = "futures-core" +version = "0.3.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4bca583b7e26f571124fe5b7561d49cb2868d79116cfa0eefce955557c6fee8c" + +[[package]] +name = "futures-executor" +version = "0.3.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ccecee823288125bd88b4d7f565c9e58e41858e47ab72e8ea2d64e93624386e0" +dependencies = [ + "futures-core", + "futures-task", + "futures-util", +] + +[[package]] +name = "futures-io" +version = "0.3.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4fff74096e71ed47f8e023204cfd0aa1289cd54ae5430a9523be060cdb849964" + +[[package]] +name = "futures-macro" +version = "0.3.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "89ca545a94061b6365f2c7355b4b32bd20df3ff95f02da9329b34ccc3bd6ee72" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.13", +] + +[[package]] +name = "futures-sink" +version = "0.3.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f43be4fe21a13b9781a69afa4985b0f6ee0e1afab2c6f454a8cf30e2b2237b6e" + +[[package]] +name = "futures-task" +version = "0.3.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "76d3d132be6c0e6aa1534069c705a74a5997a356c0dc2f86a47765e5617c5b65" + +[[package]] +name = "futures-util" +version = "0.3.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "26b01e40b772d54cf6c6d721c1d1abd0647a0106a12ecaa1c186273392a69533" +dependencies = [ + "futures-channel", + "futures-core", + "futures-io", + "futures-macro", + "futures-sink", + "futures-task", + "memchr", + "pin-project-lite", + "pin-utils", + "slab", +] + +[[package]] +name = "fxhash" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c31b6d751ae2c7f11320402d34e41349dd1016f8d5d45e48c4312bc8625af50c" +dependencies = [ + "byteorder", +] + +[[package]] +name = "gdk" +version = "0.13.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "db00839b2a68a7a10af3fa28dfb3febaba3a20c3a9ac2425a33b7df1f84a6b7d" +dependencies = [ + "bitflags", + "cairo-rs 0.9.1", + "cairo-sys-rs 0.10.0", + "gdk-pixbuf 0.9.0", + "gdk-sys 0.10.0", + "gio 0.9.1", + "gio-sys 0.10.1", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "pango 0.9.1", +] + +[[package]] +name = "gdk" +version = "0.15.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a6e05c1f572ab0e1f15be94217f0dc29088c248b14f792a5ff0af0d84bcda9e8" +dependencies = [ + "bitflags", + "cairo-rs 0.15.12", + "gdk-pixbuf 0.15.11", + "gdk-sys 0.15.1", + "gio 0.15.12", + "glib 0.15.12", + "libc", + "pango 0.15.10", +] + +[[package]] +name = "gdk-pixbuf" +version = "0.9.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8f6dae3cb99dd49b758b88f0132f8d401108e63ae8edd45f432d42cdff99998a" +dependencies = [ + "gdk-pixbuf-sys 0.10.0", + "gio 0.9.1", + "gio-sys 0.10.1", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", +] + +[[package]] +name = "gdk-pixbuf" +version = "0.15.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad38dd9cc8b099cceecdf41375bb6d481b1b5a7cd5cd603e10a69a9383f8619a" +dependencies = [ + "bitflags", + "gdk-pixbuf-sys 0.15.10", + "gio 0.15.12", + "glib 0.15.12", + "libc", +] + +[[package]] +name = "gdk-pixbuf-sys" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3bfe468a7f43e97b8d193a762b6c5cf67a7d36cacbc0b9291dbcae24bfea1e8f" +dependencies = [ + "gio-sys 0.10.1", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "gdk-pixbuf-sys" +version = "0.15.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "140b2f5378256527150350a8346dbdb08fadc13453a7a2d73aecd5fab3c402a7" +dependencies = [ + "gio-sys 0.15.10", + "glib-sys 0.15.10", + "gobject-sys 0.15.10", + "libc", + "system-deps 6.0.4", +] + +[[package]] +name = "gdk-sys" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0a9653cfc500fd268015b1ac055ddbc3df7a5c9ea3f4ccef147b3957bd140d69" +dependencies = [ + "cairo-sys-rs 0.10.0", + "gdk-pixbuf-sys 0.10.0", + "gio-sys 0.10.1", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "pango-sys 0.10.0", + "pkg-config", + "system-deps 1.3.2", +] + +[[package]] +name = "gdk-sys" +version = "0.15.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "32e7a08c1e8f06f4177fb7e51a777b8c1689f743a7bc11ea91d44d2226073a88" +dependencies = [ + "cairo-sys-rs 0.15.1", + "gdk-pixbuf-sys 0.15.10", + "gio-sys 0.15.10", + "glib-sys 0.15.10", + "gobject-sys 0.15.10", + "libc", + "pango-sys 0.15.10", + "pkg-config", + "system-deps 6.0.4", +] + +[[package]] +name = "generic-array" +version = "0.14.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "85649ca51fd72272d7821adaf274ad91c288277713d9c18820d8499a7ff69e9a" +dependencies = [ + "typenum", + "version_check", +] + +[[package]] +name = "gethostname" +version = "0.2.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c1ebd34e35c46e00bb73e81363248d627782724609fe1b6396f553f68fe3862e" +dependencies = [ + "libc", + "winapi 0.3.9", +] + +[[package]] +name = "getrandom" +version = "0.2.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c85e1d9ab2eadba7e5040d4e09cbd6d072b76a557ad64e797c2cb9d4da21d7e4" +dependencies = [ + "cfg-if 1.0.0", + "libc", + "wasi", +] + +[[package]] +name = "gimli" +version = "0.27.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad0a93d233ebf96623465aad4046a8d3aa4da22d4f4beba5388838c8a434bbb4" + +[[package]] +name = "gio" +version = "0.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1fb60242bfff700772dae5d9e3a1f7aa2e4ebccf18b89662a16acb2822568561" +dependencies = [ + "bitflags", + "futures", + "futures-channel", + "futures-core", + "futures-io", + "futures-util", + "gio-sys 0.10.1", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "once_cell", + "thiserror", +] + +[[package]] +name = "gio" +version = "0.15.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68fdbc90312d462781a395f7a16d96a2b379bb6ef8cd6310a2df272771c4283b" +dependencies = [ + "bitflags", + "futures-channel", + "futures-core", + "futures-io", + "gio-sys 0.15.10", + "glib 0.15.12", + "libc", + "once_cell", + "thiserror", +] + +[[package]] +name = "gio-sys" +version = "0.10.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5e24fb752f8f5d2cf6bbc2c606fd2bc989c81c5e2fe321ab974d54f8b6344eac" +dependencies = [ + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "system-deps 1.3.2", + "winapi 0.3.9", +] + +[[package]] +name = "gio-sys" +version = "0.15.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "32157a475271e2c4a023382e9cab31c4584ee30a97da41d3c4e9fdd605abcf8d" +dependencies = [ + "glib-sys 0.15.10", + "gobject-sys 0.15.10", + "libc", + "system-deps 6.0.4", + "winapi 0.3.9", +] + +[[package]] +name = "glib" +version = "0.10.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0c685013b7515e668f1b57a165b009d4d28cb139a8a989bbd699c10dad29d0c5" +dependencies = [ + "bitflags", + "futures-channel", + "futures-core", + "futures-executor", + "futures-task", + "futures-util", + "glib-macros 0.10.1", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "once_cell", +] + +[[package]] +name = "glib" +version = "0.15.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "edb0306fbad0ab5428b0ca674a23893db909a98582969c9b537be4ced78c505d" +dependencies = [ + "bitflags", + "futures-channel", + "futures-core", + "futures-executor", + "futures-task", + "glib-macros 0.15.13", + "glib-sys 0.15.10", + "gobject-sys 0.15.10", + "libc", + "once_cell", + "smallvec", + "thiserror", +] + +[[package]] +name = "glib-macros" +version = "0.10.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "41486a26d1366a8032b160b59065a59fb528530a46a49f627e7048fb8c064039" +dependencies = [ + "anyhow", + "heck 0.3.3", + "itertools", + "proc-macro-crate 0.1.5", + "proc-macro-error", + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "glib-macros" +version = "0.15.13" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "10c6ae9f6fa26f4fb2ac16b528d138d971ead56141de489f8111e259b9df3c4a" +dependencies = [ + "anyhow", + "heck 0.4.1", + "proc-macro-crate 1.3.1", + "proc-macro-error", + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "glib-sys" +version = "0.10.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c7e9b997a66e9a23d073f2b1abb4dbfc3925e0b8952f67efd8d9b6e168e4cdc1" +dependencies = [ + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "glib-sys" +version = "0.15.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ef4b192f8e65e9cf76cbf4ea71fa8e3be4a0e18ffe3d68b8da6836974cc5bad4" +dependencies = [ + "libc", + "system-deps 6.0.4", +] + +[[package]] +name = "glob" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d2fabcfbdc87f4758337ca535fb41a6d701b65693ce38287d856d1674551ec9b" + +[[package]] +name = "gobject-sys" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "952133b60c318a62bf82ee75b93acc7e84028a093e06b9e27981c2b6fe68218c" +dependencies = [ + "glib-sys 0.10.1", + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "gobject-sys" +version = "0.15.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0d57ce44246becd17153bd035ab4d32cfee096a657fc01f2231c9278378d1e0a" +dependencies = [ + "glib-sys 0.15.10", + "libc", + "system-deps 6.0.4", +] + +[[package]] +name = "gstreamer" +version = "0.16.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9ff5d0f7ff308ae37e6eb47b6ded17785bdea06e438a708cd09e0288c1862f33" +dependencies = [ + "bitflags", + "cfg-if 1.0.0", + "futures-channel", + "futures-core", + "futures-util", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "gstreamer-sys", + "libc", + "muldiv", + "num-rational", + "once_cell", + "paste", + "pretty-hex", + "thiserror", +] + +[[package]] +name = "gstreamer-app" +version = "0.16.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cc80888271338c3ede875d8cafc452eb207476ff5539dcbe0018a8f5b827af0e" +dependencies = [ + "bitflags", + "futures-core", + "futures-sink", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "gstreamer", + "gstreamer-app-sys", + "gstreamer-base", + "gstreamer-sys", + "libc", + "once_cell", +] + +[[package]] +name = "gstreamer-app-sys" +version = "0.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "813f64275c9e7b33b828b9efcf9dfa64b95996766d4de996e84363ac65b87e3d" +dependencies = [ + "glib-sys 0.10.1", + "gstreamer-base-sys", + "gstreamer-sys", + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "gstreamer-base" +version = "0.16.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bafd01c56f59cb10f4b5a10f97bb4bdf8c2b2784ae5b04da7e2d400cf6e6afcf" +dependencies = [ + "bitflags", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "gstreamer", + "gstreamer-base-sys", + "gstreamer-sys", + "libc", +] + +[[package]] +name = "gstreamer-base-sys" +version = "0.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a4b7b6dc2d6e160a1ae28612f602bd500b3fa474ce90bf6bb2f08072682beef5" +dependencies = [ + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "gstreamer-sys", + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "gstreamer-sys" +version = "0.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fc1f154082d01af5718c5f8a8eb4f565a4ea5586ad8833a8fc2c2aa6844b601d" +dependencies = [ + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "gstreamer-video" +version = "0.16.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f7bbb1485d87469849ec45c08e03c2f280d3ea20ff3c439d03185be54e3ce98e" +dependencies = [ + "bitflags", + "futures-channel", + "futures-util", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "gstreamer", + "gstreamer-base", + "gstreamer-base-sys", + "gstreamer-sys", + "gstreamer-video-sys", + "libc", + "once_cell", +] + +[[package]] +name = "gstreamer-video-sys" +version = "0.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "92347e46438007d6a2386302125f62cb9df6769cdacb931af5c0f12c1ee21de4" +dependencies = [ + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "gstreamer-base-sys", + "gstreamer-sys", + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "gtk" +version = "0.9.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2f022f2054072b3af07666341984562c8e626a79daa8be27b955d12d06a5ad6a" +dependencies = [ + "atk 0.9.0", + "bitflags", + "cairo-rs 0.9.1", + "cairo-sys-rs 0.10.0", + "cc", + "gdk 0.13.2", + "gdk-pixbuf 0.9.0", + "gdk-pixbuf-sys 0.10.0", + "gdk-sys 0.10.0", + "gio 0.9.1", + "gio-sys 0.10.1", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "gtk-sys 0.10.0", + "libc", + "once_cell", + "pango 0.9.1", + "pango-sys 0.10.0", + "pkg-config", +] + +[[package]] +name = "gtk" +version = "0.15.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "92e3004a2d5d6d8b5057d2b57b3712c9529b62e82c77f25c1fecde1fd5c23bd0" +dependencies = [ + "atk 0.15.1", + "bitflags", + "cairo-rs 0.15.12", + "field-offset", + "futures-channel", + "gdk 0.15.4", + "gdk-pixbuf 0.15.11", + "gio 0.15.12", + "glib 0.15.12", + "gtk-sys 0.15.3", + "gtk3-macros", + "libc", + "once_cell", + "pango 0.15.10", + "pkg-config", +] + +[[package]] +name = "gtk-sys" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "89acda6f084863307d948ba64a4b1ef674e8527dddab147ee4cdcc194c880457" +dependencies = [ + "atk-sys 0.10.0", + "cairo-sys-rs 0.10.0", + "gdk-pixbuf-sys 0.10.0", + "gdk-sys 0.10.0", + "gio-sys 0.10.1", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "pango-sys 0.10.0", + "system-deps 1.3.2", +] + +[[package]] +name = "gtk-sys" +version = "0.15.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d5bc2f0587cba247f60246a0ca11fe25fb733eabc3de12d1965fc07efab87c84" +dependencies = [ + "atk-sys 0.15.1", + "cairo-sys-rs 0.15.1", + "gdk-pixbuf-sys 0.15.10", + "gdk-sys 0.15.1", + "gio-sys 0.15.10", + "glib-sys 0.15.10", + "gobject-sys 0.15.10", + "libc", + "pango-sys 0.15.10", + "system-deps 6.0.4", +] + +[[package]] +name = "gtk3-macros" +version = "0.15.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "684c0456c086e8e7e9af73ec5b84e35938df394712054550e81558d21c44ab0d" +dependencies = [ + "anyhow", + "proc-macro-crate 1.3.1", + "proc-macro-error", + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "hashbrown" +version = "0.12.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8a9ee70c43aaf417c914396645a0fa852624801b24ebb7ae78fe8272889ac888" + +[[package]] +name = "hbb_common" +version = "0.1.0" +dependencies = [ + "anyhow", + "bytes", + "confy", + "directories-next", + "dirs-next", + "env_logger 0.9.3", + "filetime", + "futures", + "futures-util", + "lazy_static", + "log", + "mac_address", + "protobuf", + "protobuf-codegen-pure", + "quinn", + "rand 0.8.5", + "regex", + "serde 1.0.159", + "serde_derive", + "serde_json 1.0.95", + "socket2 0.3.19", + "sodiumoxide", + "tokio", + "tokio-socks", + "tokio-util 0.6.10", + "toml 0.5.11", + "winapi 0.3.9", + "zstd", +] + +[[package]] +name = "heck" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6d621efb26863f0e9924c6ac577e8275e5e6b77455db64ffa6c65c904e9e132c" +dependencies = [ + "unicode-segmentation", +] + +[[package]] +name = "heck" +version = "0.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "95505c38b4572b2d910cecb0281560f54b440a19336cbbcb27bf6ce6adc6f5a8" + +[[package]] +name = "hermit-abi" +version = "0.1.19" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "62b467343b94ba476dcb2500d242dadbb39557df889310ac77c5d99100aaac33" +dependencies = [ + "libc", +] + +[[package]] +name = "hermit-abi" +version = "0.2.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ee512640fe35acbfb4bb779db6f0d80704c2cacfa2e39b601ef3e3f47d1ae4c7" +dependencies = [ + "libc", +] + +[[package]] +name = "hermit-abi" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fed44880c466736ef9a5c5b5facefb5ed0785676d0c02d612db14e54f0d84286" + +[[package]] +name = "hound" +version = "3.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4d13cdbd5dbb29f9c88095bbdc2590c9cba0d0a1269b983fef6b2cdd7e9f4db1" + +[[package]] +name = "humantime" +version = "2.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9a3a5bfb195931eeb336b2a7b4d761daec841b97f947d34394601737a7bba5e4" + +[[package]] +name = "iana-time-zone" +version = "0.1.56" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0722cd7114b7de04316e7ea5456a0bbb20e4adb46fd27a3697adb812cff0f37c" +dependencies = [ + "android_system_properties", + "core-foundation-sys", + "iana-time-zone-haiku", + "js-sys", + "wasm-bindgen", + "windows", +] + +[[package]] +name = "iana-time-zone-haiku" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0703ae284fc167426161c2e3f1da3ea71d94b21bedbcc9494e92b28e334e3dca" +dependencies = [ + "cxx", + "cxx-build", +] + +[[package]] +name = "ident_case" +version = "1.0.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b9e0384b61958566e926dc50660321d12159025e767c18e043daf26b70104c39" + +[[package]] +name = "image" +version = "0.23.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "24ffcb7e7244a9bf19d35bf2883b9c080c4ced3c07a9895572178cdb8f13f6a1" +dependencies = [ + "bytemuck", + "byteorder", + "color_quant", + "num-iter", + "num-rational", + "num-traits 0.2.15", + "png", + "tiff", +] + +[[package]] +name = "include_dir" +version = "0.7.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "18762faeff7122e89e0857b02f7ce6fcc0d101d5e9ad2ad7846cc01d61b7f19e" +dependencies = [ + "include_dir_macros", +] + +[[package]] +name = "include_dir_macros" +version = "0.7.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b139284b5cf57ecfa712bcc66950bb635b31aff41c188e8a4cfc758eca374a3f" +dependencies = [ + "proc-macro2", + "quote", +] + +[[package]] +name = "indexmap" +version = "1.9.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bd070e393353796e801d209ad339e89596eb4c8d430d18ede6a1cced8fafbd99" +dependencies = [ + "autocfg 1.1.0", + "hashbrown", +] + +[[package]] +name = "instant" +version = "0.1.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7a5bbe824c507c5da5956355e86a746d82e0e1464f65d862cc5e71da70e94b2c" +dependencies = [ + "cfg-if 1.0.0", +] + +[[package]] +name = "io-lifetimes" +version = "1.0.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9c66c74d2ae7e79a5a8f7ac924adbe38ee42a859c6539ad869eb51f0b52dc220" +dependencies = [ + "hermit-abi 0.3.1", + "libc", + "windows-sys 0.48.0", +] + +[[package]] +name = "iovec" +version = "0.1.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b2b3ea6ff95e175473f8ffe6a7eb7c00d054240321b84c57051175fe3c1e075e" +dependencies = [ + "libc", +] + +[[package]] +name = "itertools" +version = "0.9.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "284f18f85651fe11e8a991b2adb42cb078325c996ed026d994719efcfca1d54b" +dependencies = [ + "either", +] + +[[package]] +name = "itoa" +version = "0.3.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8324a32baf01e2ae060e9de58ed0bc2320c9a2833491ee36cd3b4c414de4db8c" + +[[package]] +name = "itoa" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "453ad9f582a441959e5f0d088b02ce04cfe8d51a8eaf077f12ac6d3e94164ca6" + +[[package]] +name = "jni" +version = "0.19.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c6df18c2e3db7e453d3c6ac5b3e9d5182664d28788126d39b91f2d1e22b017ec" +dependencies = [ + "cesu8", + "combine", + "jni-sys", + "log", + "thiserror", + "walkdir", +] + +[[package]] +name = "jni-sys" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8eaf4bc02d17cbdd7ff4c7438cafcdf7fb9a4613313ad11b4f8fefe7d3fa0130" + +[[package]] +name = "jobserver" +version = "0.1.26" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "936cfd212a0155903bcbc060e316fb6cc7cbf2e1907329391ebadc1fe0ce77c2" +dependencies = [ + "libc", +] + +[[package]] +name = "jpeg-decoder" +version = "0.1.22" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "229d53d58899083193af11e15917b5640cd40b29ff475a1fe4ef725deb02d0f2" + +[[package]] +name = "js-sys" +version = "0.3.61" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "445dde2150c55e483f3d8416706b97ec8e8237c307e5b7b4b8dd15e6af2a0730" +dependencies = [ + "wasm-bindgen", +] + +[[package]] +name = "kernel32-sys" +version = "0.2.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7507624b29483431c0ba2d82aece8ca6cdba9382bff4ddd0f7490560c056098d" +dependencies = [ + "winapi 0.2.8", + "winapi-build", +] + +[[package]] +name = "lazy_static" +version = "1.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e2abad23fbc42b3700f2f279844dc832adb2b2eb069b2df918f455c4e18cc646" + +[[package]] +name = "lazycell" +version = "1.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "830d08ce1d1d941e6b30645f1a0eb5643013d835ce3779a5fc208261dbe10f55" + +[[package]] +name = "libappindicator" +version = "0.6.1" +source = "git+https://github.com/liyue201/libappindicator-rs#3763cfd629dd90050af1feafa643cbfca0bf487e" +dependencies = [ + "glib 0.10.3", + "gtk 0.9.2", + "gtk-sys 0.10.0", + "libappindicator-sys 0.6.1", + "log", +] + +[[package]] +name = "libappindicator" +version = "0.7.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "db2d3cb96d092b4824cb306c9e544c856a4cb6210c1081945187f7f1924b47e8" +dependencies = [ + "glib 0.15.12", + "gtk 0.15.5", + "gtk-sys 0.15.3", + "libappindicator-sys 0.7.3", + "log", +] + +[[package]] +name = "libappindicator-sys" +version = "0.6.1" +source = "git+https://github.com/liyue201/libappindicator-rs#3763cfd629dd90050af1feafa643cbfca0bf487e" +dependencies = [ + "gtk-sys 0.10.0", + "pkg-config", +] + +[[package]] +name = "libappindicator-sys" +version = "0.7.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f1b3b6681973cea8cc3bce7391e6d7d5502720b80a581c9a95c9cbaf592826aa" +dependencies = [ + "gtk-sys 0.15.3", + "libloading", + "once_cell", +] + +[[package]] +name = "libc" +version = "0.2.141" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3304a64d199bb964be99741b7a14d26972741915b3649639149b2479bb46f4b5" + +[[package]] +name = "libdbus-sys" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9f8d7ae751e1cb825c840ae5e682f59b098cdfd213c350ac268b61449a5f58a0" +dependencies = [ + "pkg-config", +] + +[[package]] +name = "libloading" +version = "0.7.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b67380fd3b2fbe7527a606e18729d21c6f3951633d0500574c4dc22d2d638b9f" +dependencies = [ + "cfg-if 1.0.0", + "winapi 0.3.9", +] + +[[package]] +name = "libpulse-binding" +version = "2.27.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1745b20bfc194ac12ef828f144f0ec2d4a7fe993281fa3567a0bd4969aee6890" +dependencies = [ + "bitflags", + "libc", + "libpulse-sys", + "num-derive", + "num-traits 0.2.15", + "winapi 0.3.9", +] + +[[package]] +name = "libpulse-simple-binding" +version = "2.27.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5ced94199e6e44133431374e4043f34e1f0697ebfb7b7d6c244a65bfaedf0e31" +dependencies = [ + "libpulse-binding", + "libpulse-simple-sys", + "libpulse-sys", +] + +[[package]] +name = "libpulse-simple-sys" +version = "1.20.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "84e423d9c619c908ce9b4916080e65ab586ca55b8c4939379f15e6e72fb43842" +dependencies = [ + "libpulse-sys", + "pkg-config", +] + +[[package]] +name = "libpulse-sys" +version = "1.20.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2191e6880818d1df4cf72eac8e91dce7a5a52ba0da4b2a5cdafabc22b937eadb" +dependencies = [ + "libc", + "num-derive", + "num-traits 0.2.15", + "pkg-config", + "winapi 0.3.9", +] + +[[package]] +name = "libsamplerate-sys" +version = "0.1.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "28853b399f78f8281cd88d333b54a63170c4275f6faea66726a2bea5cca72e0d" +dependencies = [ + "cmake", +] + +[[package]] +name = "libsodium-sys" +version = "0.2.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6b779387cd56adfbc02ea4a668e704f729be8d6a6abd2c27ca5ee537849a92fd" +dependencies = [ + "cc", + "libc", + "pkg-config", + "walkdir", +] + +[[package]] +name = "link-cplusplus" +version = "1.0.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ecd207c9c713c34f95a097a5b029ac2ce6010530c7b49d7fea24d977dede04f5" +dependencies = [ + "cc", +] + +[[package]] +name = "linux-raw-sys" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d59d8c75012853d2e872fb56bc8a2e53718e2cafe1a4c823143141c6d90c322f" + +[[package]] +name = "lock_api" +version = "0.4.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "435011366fe56583b16cf956f9df0095b405b82d76425bc8981c0e22e60ec4df" +dependencies = [ + "autocfg 1.1.0", + "scopeguard", +] + +[[package]] +name = "log" +version = "0.4.17" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "abb12e687cfb44aa40f41fc3978ef76448f9b6038cad6aef4259d3c095a2382e" +dependencies = [ + "cfg-if 1.0.0", +] + +[[package]] +name = "mac_address" +version = "1.1.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b238e3235c8382b7653c6408ed1b08dd379bdb9fdf990fb0bbae3db2cc0ae963" +dependencies = [ + "nix 0.23.2", + "winapi 0.3.9", +] + +[[package]] +name = "mach" +version = "0.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b823e83b2affd8f40a9ee8c29dbc56404c1e34cd2710921f2801e2cf29527afa" +dependencies = [ + "libc", +] + +[[package]] +name = "machine-uid" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1f1595709b0a7386bcd56ba34d250d626e5503917d05d32cdccddcd68603e212" +dependencies = [ + "winreg 0.6.2", +] + +[[package]] +name = "magnum-opus" +version = "0.4.0" +source = "git+https://github.com/TheRadioGuy/magnum-opus#171e1d021004626f7444d1e39b98f50bc3cb2604" +dependencies = [ + "libc", + "opusic-sys", +] + +[[package]] +name = "malloc_buf" +version = "0.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "62bb907fe88d54d8d9ce32a3cceab4218ed2f6b7d35617cafe9adf84e43919cb" +dependencies = [ + "libc", +] + +[[package]] +name = "memchr" +version = "2.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2dffe52ecf27772e601905b7522cb4ef790d2cc203488bbd0e2fe85fcb74566d" + +[[package]] +name = "memoffset" +version = "0.6.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5aa361d4faea93603064a027415f07bd8e1d5c88c9fbf68bf56a285428fd79ce" +dependencies = [ + "autocfg 1.1.0", +] + +[[package]] +name = "memoffset" +version = "0.8.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d61c719bcfbcf5d62b3a09efa6088de8c54bc0bfcd3ea7ae39fcc186108b8de1" +dependencies = [ + "autocfg 1.1.0", +] + +[[package]] +name = "minimal-lexical" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a" + +[[package]] +name = "miniz_oxide" +version = "0.3.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "791daaae1ed6889560f8c4359194f56648355540573244a5448a83ba1ecc7435" +dependencies = [ + "adler32", +] + +[[package]] +name = "miniz_oxide" +version = "0.4.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a92518e98c078586bc6c934028adcca4c92a53d6a958196de835170a01d84e4b" +dependencies = [ + "adler", + "autocfg 1.1.0", +] + +[[package]] +name = "miniz_oxide" +version = "0.6.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b275950c28b37e794e8c55d88aeb5e139d0ce23fdbbeda68f8d7174abdf9e8fa" +dependencies = [ + "adler", +] + +[[package]] +name = "mio" +version = "0.6.23" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4afd66f5b91bf2a3bc13fad0e21caedac168ca4c707504e75585648ae80e4cc4" +dependencies = [ + "cfg-if 0.1.10", + "fuchsia-zircon", + "fuchsia-zircon-sys", + "iovec", + "kernel32-sys", + "libc", + "log", + "miow 0.2.2", + "net2", + "slab", + "winapi 0.2.8", +] + +[[package]] +name = "mio" +version = "0.8.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5b9d9a46eff5b4ff64b45a9e316a6d1e0bc719ef429cbec4dc630684212bfdf9" +dependencies = [ + "libc", + "log", + "wasi", + "windows-sys 0.45.0", +] + +[[package]] +name = "mio-named-pipes" +version = "0.1.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0840c1c50fd55e521b247f949c241c9997709f23bd7f023b9762cd561e935656" +dependencies = [ + "log", + "mio 0.6.23", + "miow 0.3.7", + "winapi 0.3.9", +] + +[[package]] +name = "miow" +version = "0.2.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ebd808424166322d4a38da87083bfddd3ac4c131334ed55856112eb06d46944d" +dependencies = [ + "kernel32-sys", + "net2", + "winapi 0.2.8", + "ws2_32-sys", +] + +[[package]] +name = "miow" +version = "0.3.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b9f1c5b025cda876f66ef43a113f91ebc9f4ccef34843000e0adf6ebbab84e21" +dependencies = [ + "winapi 0.3.9", +] + +[[package]] +name = "miow" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a7377f7792b3afb6a3cba68daa54ca23c032137010460d667fda53a8d66be00e" +dependencies = [ + "windows-sys 0.28.0", +] + +[[package]] +name = "muldiv" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0419348c027fa7be448d2ae7ea0e4e04c2334c31dc4e74ab29f00a2a7ca69204" + +[[package]] +name = "ndk" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2032c77e030ddee34a6787a64166008da93f6a352b629261d0fee232b8742dd4" +dependencies = [ + "bitflags", + "jni-sys", + "ndk-sys", + "num_enum", + "thiserror", +] + +[[package]] +name = "ndk-context" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "27b02d87554356db9e9a873add8782d4ea6e3e58ea071a9adb9a2e8ddb884a8b" + +[[package]] +name = "ndk-glue" +version = "0.6.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0d0c4a7b83860226e6b4183edac21851f05d5a51756e97a1144b7f5a6b63e65f" +dependencies = [ + "lazy_static", + "libc", + "log", + "ndk", + "ndk-context", + "ndk-macro", + "ndk-sys", +] + +[[package]] +name = "ndk-macro" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0df7ac00c4672f9d5aece54ee3347520b7e20f158656c7db2e6de01902eb7a6c" +dependencies = [ + "darling", + "proc-macro-crate 1.3.1", + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "ndk-sys" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6e5a6ae77c8ee183dcbbba6150e2e6b9f3f4196a7666c02a715a95692ec1fa97" +dependencies = [ + "jni-sys", +] + +[[package]] +name = "net2" +version = "0.2.38" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "74d0df99cfcd2530b2e694f6e17e7f37b8e26bb23983ac530c0c97408837c631" +dependencies = [ + "cfg-if 0.1.10", + "libc", + "winapi 0.3.9", +] + +[[package]] +name = "nix" +version = "0.22.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e4916f159ed8e5de0082076562152a76b7a1f64a01fd9d1e0fea002c37624faf" +dependencies = [ + "bitflags", + "cc", + "cfg-if 1.0.0", + "libc", + "memoffset 0.6.5", +] + +[[package]] +name = "nix" +version = "0.23.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8f3790c00a0150112de0f4cd161e3d7fc4b2d8a5542ffc35f099a2562aecb35c" +dependencies = [ + "bitflags", + "cc", + "cfg-if 1.0.0", + "libc", + "memoffset 0.6.5", +] + +[[package]] +name = "nix" +version = "0.26.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bfdda3d196821d6af13126e40375cdf7da646a96114af134d5f417a9a1dc8e1a" +dependencies = [ + "bitflags", + "cfg-if 1.0.0", + "libc", + "static_assertions", +] + +[[package]] +name = "nom" +version = "7.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d273983c5a657a70a3e8f2a01329822f3b8c8172b73826411a55751e404a0a4a" +dependencies = [ + "memchr", + "minimal-lexical", +] + +[[package]] +name = "ntapi" +version = "0.3.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c28774a7fd2fbb4f0babd8237ce554b73af68021b5f695a3cebd6c59bac0980f" +dependencies = [ + "winapi 0.3.9", +] + +[[package]] +name = "num-complex" +version = "0.4.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "02e0d21255c828d6f128a1e41534206671e8c3ea0c62f32291e808dc82cff17d" +dependencies = [ + "num-traits 0.2.15", +] + +[[package]] +name = "num-derive" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "876a53fff98e03a936a674b29568b0e605f06b29372c2489ff4de23f1949743d" +dependencies = [ + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "num-integer" +version = "0.1.45" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "225d3389fb3509a24c93f5c29eb6bde2586b98d9f016636dff58d7c6f7569cd9" +dependencies = [ + "autocfg 1.1.0", + "num-traits 0.2.15", +] + +[[package]] +name = "num-iter" +version = "0.1.43" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7d03e6c028c5dc5cac6e2dec0efda81fc887605bb3d884578bb6d6bf7514e252" +dependencies = [ + "autocfg 1.1.0", + "num-integer", + "num-traits 0.2.15", +] + +[[package]] +name = "num-rational" +version = "0.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "12ac428b1cb17fce6f731001d307d351ec70a6d202fc2e60f7d4c5e42d8f4f07" +dependencies = [ + "autocfg 1.1.0", + "num-integer", + "num-traits 0.2.15", +] + +[[package]] +name = "num-traits" +version = "0.1.43" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "92e5113e9fd4cc14ded8e499429f396a20f98c772a47cc8622a736e1ec843c31" +dependencies = [ + "num-traits 0.2.15", +] + +[[package]] +name = "num-traits" +version = "0.2.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "578ede34cf02f8924ab9447f50c28075b4d3e5b269972345e7e0372b38c6cdcd" +dependencies = [ + "autocfg 1.1.0", +] + +[[package]] +name = "num_cpus" +version = "1.15.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0fac9e2da13b5eb447a6ce3d392f23a29d8694bff781bf03a16cd9ac8697593b" +dependencies = [ + "hermit-abi 0.2.6", + "libc", +] + +[[package]] +name = "num_enum" +version = "0.5.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1f646caf906c20226733ed5b1374287eb97e3c2a5c227ce668c1f2ce20ae57c9" +dependencies = [ + "num_enum_derive", +] + +[[package]] +name = "num_enum_derive" +version = "0.5.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dcbff9bc912032c62bf65ef1d5aea88983b420f4f839db1e9b0c281a25c9c799" +dependencies = [ + "proc-macro-crate 1.3.1", + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "num_threads" +version = "0.1.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2819ce041d2ee131036f4fc9d6ae7ae125a3a40e97ba64d04fe799ad9dabbb44" +dependencies = [ + "libc", +] + +[[package]] +name = "objc" +version = "0.2.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "915b1b472bc21c53464d6c8461c9d3af805ba1ef837e1cac254428f4a77177b1" +dependencies = [ + "malloc_buf", +] + +[[package]] +name = "objc-foundation" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1add1b659e36c9607c7aab864a76c7a4c2760cd0cd2e120f3fb8b952c7e22bf9" +dependencies = [ + "block", + "objc", + "objc_id", +] + +[[package]] +name = "objc_id" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c92d4ddb4bd7b50d730c215ff871754d0da6b2178849f8a2a2ab69712d0c073b" +dependencies = [ + "objc", +] + +[[package]] +name = "object" +version = "0.30.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ea86265d3d3dcb6a27fc51bd29a4bf387fae9d2986b823079d4986af253eb439" +dependencies = [ + "memchr", +] + +[[package]] +name = "oboe" +version = "0.4.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "27f63c358b4fa0fbcfefd7c8be5cfc39c08ce2389f5325687e7762a48d30a5c1" +dependencies = [ + "jni", + "ndk", + "ndk-context", + "num-derive", + "num-traits 0.2.15", + "oboe-sys", +] + +[[package]] +name = "oboe-sys" +version = "0.4.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3370abb7372ed744232c12954d920d1a40f1c4686de9e79e800021ef492294bd" +dependencies = [ + "cc", +] + +[[package]] +name = "once_cell" +version = "1.17.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b7e5500299e16ebb147ae15a00a942af264cf3688f47923b8fc2cd5858f23ad3" + +[[package]] +name = "openssl-probe" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ff011a302c396a5197692431fc1948019154afc178baf7d8e37367442a4601cf" + +[[package]] +name = "opusic-sys" +version = "0.3.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5eace752ce07a037241dba8f02c654799f051e431b27028056bcb480e83b54f5" +dependencies = [ + "cmake", + "libc", +] + +[[package]] +name = "os_str_bytes" +version = "6.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ceedf44fb00f2d1984b0bc98102627ce622e083e49a5bacdb3e514fa4238e267" + +[[package]] +name = "padlock" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c10569378a1dacd9f30dbe7ae49e054d2c45dc2f8ee49899903e09c3924e8b6f" + +[[package]] +name = "pango" +version = "0.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9937068580bebd8ced19975938573803273ccbcbd598c58d4906efd4ac87c438" +dependencies = [ + "bitflags", + "glib 0.10.3", + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "once_cell", + "pango-sys 0.10.0", +] + +[[package]] +name = "pango" +version = "0.15.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "22e4045548659aee5313bde6c582b0d83a627b7904dd20dc2d9ef0895d414e4f" +dependencies = [ + "bitflags", + "glib 0.15.12", + "libc", + "once_cell", + "pango-sys 0.15.10", +] + +[[package]] +name = "pango-sys" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "24d2650c8b62d116c020abd0cea26a4ed96526afda89b1c4ea567131fdefc890" +dependencies = [ + "glib-sys 0.10.1", + "gobject-sys 0.10.0", + "libc", + "system-deps 1.3.2", +] + +[[package]] +name = "pango-sys" +version = "0.15.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d2a00081cde4661982ed91d80ef437c20eacaf6aa1a5962c0279ae194662c3aa" +dependencies = [ + "glib-sys 0.15.10", + "gobject-sys 0.15.10", + "libc", + "system-deps 6.0.4", +] + +[[package]] +name = "parity-tokio-ipc" +version = "0.7.3-1" +source = "git+https://github.com/open-trade/parity-tokio-ipc#20b2895910161605210657f3e751edd55321f698" +dependencies = [ + "futures", + "libc", + "log", + "mio-named-pipes", + "miow 0.4.0", + "rand 0.8.5", + "tokio", + "winapi 0.3.9", +] + +[[package]] +name = "parking_lot" +version = "0.11.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7d17b78036a60663b797adeaee46f5c9dfebb86948d1255007a1d6be0271ff99" +dependencies = [ + "instant", + "lock_api", + "parking_lot_core 0.8.6", +] + +[[package]] +name = "parking_lot" +version = "0.12.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3742b2c103b9f06bc9fff0a37ff4912935851bee6d36f3c02bcc755bcfec228f" +dependencies = [ + "lock_api", + "parking_lot_core 0.9.7", +] + +[[package]] +name = "parking_lot_core" +version = "0.8.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "60a2cfe6f0ad2bfc16aefa463b497d5c7a5ecd44a23efa72aa342d90177356dc" +dependencies = [ + "cfg-if 1.0.0", + "instant", + "libc", + "redox_syscall 0.2.16", + "smallvec", + "winapi 0.3.9", +] + +[[package]] +name = "parking_lot_core" +version = "0.9.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9069cbb9f99e3a5083476ccb29ceb1de18b9118cafa53e90c9551235de2b9521" +dependencies = [ + "cfg-if 1.0.0", + "libc", + "redox_syscall 0.2.16", + "smallvec", + "windows-sys 0.45.0", +] + +[[package]] +name = "paste" +version = "1.0.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9f746c4065a8fa3fe23974dd82f15431cc8d40779821001404d10d2e79ca7d79" + +[[package]] +name = "peeking_take_while" +version = "0.1.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "19b17cddbe7ec3f8bc800887bab5e717348c95ea2ca0b1bf0837fb964dc67099" + +[[package]] +name = "phf" +version = "0.7.24" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b3da44b85f8e8dfaec21adae67f95d93244b2ecf6ad2a692320598dcc8e6dd18" +dependencies = [ + "phf_shared", +] + +[[package]] +name = "phf_codegen" +version = "0.7.24" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b03e85129e324ad4166b06b2c7491ae27fe3ec353af72e72cd1654c7225d517e" +dependencies = [ + "phf_generator", + "phf_shared", +] + +[[package]] +name = "phf_generator" +version = "0.7.24" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "09364cc93c159b8b06b1f4dd8a4398984503483891b0c26b867cf431fb132662" +dependencies = [ + "phf_shared", + "rand 0.6.5", +] + +[[package]] +name = "phf_shared" +version = "0.7.24" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "234f71a15de2288bcb7e3b6515828d22af7ec8598ee6d24c3b526fa0a80b67a0" +dependencies = [ + "siphasher", +] + +[[package]] +name = "pin-project" +version = "1.0.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad29a609b6bcd67fee905812e544992d216af9d755757c05ed2d0e15a74c6ecc" +dependencies = [ + "pin-project-internal", +] + +[[package]] +name = "pin-project-internal" +version = "1.0.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "069bdb1e05adc7a8990dce9cc75370895fbe4e3d58b9b73bf1aee56359344a55" +dependencies = [ + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "pin-project-lite" +version = "0.2.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e0a7ae3ac2f1173085d398531c705756c94a4c56843785df85a60c1a0afac116" + +[[package]] +name = "pin-utils" +version = "0.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8b870d8c151b6f2fb93e84a13146138f05d02ed11c7e7c54f8826aaaf7c9f184" + +[[package]] +name = "pkg-config" +version = "0.3.26" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6ac9a59f73473f1b8d852421e59e64809f025994837ef743615c6d0c5b305160" + +[[package]] +name = "png" +version = "0.16.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3c3287920cb847dee3de33d301c463fba14dda99db24214ddf93f83d3021f4c6" +dependencies = [ + "bitflags", + "crc32fast", + "deflate", + "miniz_oxide 0.3.7", +] + +[[package]] +name = "ppv-lite86" +version = "0.2.17" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5b40af805b3121feab8a3c29f04d8ad262fa8e0561883e7653e024ae4479e6de" + +[[package]] +name = "pretty-hex" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bc5c99d529f0d30937f6f4b8a86d988047327bb88d04d2c4afc356de74722131" + +[[package]] +name = "primal-check" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9df7f93fd637f083201473dab4fee2db4c429d32e55e3299980ab3957ab916a0" +dependencies = [ + "num-integer", +] + +[[package]] +name = "proc-macro-crate" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1d6ea3c4595b96363c13943497db34af4460fb474a95c43f4446ad341b8c9785" +dependencies = [ + "toml 0.5.11", +] + +[[package]] +name = "proc-macro-crate" +version = "1.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7f4c021e1093a56626774e81216a4ce732a735e5bad4868a03f3ed65ca0c3919" +dependencies = [ + "once_cell", + "toml_edit", +] + +[[package]] +name = "proc-macro-error" +version = "1.0.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "da25490ff9892aab3fcf7c36f08cfb902dd3e71ca0f9f9517bea02a73a5ce38c" +dependencies = [ + "proc-macro-error-attr", + "proc-macro2", + "quote", + "syn 1.0.109", + "version_check", +] + +[[package]] +name = "proc-macro-error-attr" +version = "1.0.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a1be40180e52ecc98ad80b184934baf3d0d29f979574e439af5a55274b35f869" +dependencies = [ + "proc-macro2", + "quote", + "version_check", +] + +[[package]] +name = "proc-macro2" +version = "1.0.56" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2b63bdb0cd06f1f4dedf69b254734f9b45af66e4a031e42a7480257d9898b435" +dependencies = [ + "unicode-ident", +] + +[[package]] +name = "protobuf" +version = "3.0.0-alpha.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9d5ef59c35c7472ce5e1b6c5924b87585143d1fc2cf39eae0009bba6c4df62f1" + +[[package]] +name = "protobuf-codegen" +version = "3.0.0-alpha.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "89100ee819f69b77a4cab389fec9dd155a305af4c615e6413ec1ef9341f333ef" +dependencies = [ + "anyhow", + "protobuf", + "protobuf-parse", + "thiserror", +] + +[[package]] +name = "protobuf-codegen-pure" +version = "3.0.0-alpha.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "79453e74d08190551e821533ee42c447f9e21ca26f83520e120e6e8af27f6879" +dependencies = [ + "anyhow", + "protobuf", + "protobuf-codegen", + "protobuf-parse", + "thiserror", +] + +[[package]] +name = "protobuf-parse" +version = "3.0.0-alpha.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c265ffc69976efc3056955b881641add3186ad0be893ef10622482d80d1d2b68" +dependencies = [ + "anyhow", + "protobuf", + "protoc", + "tempfile", + "thiserror", +] + +[[package]] +name = "protoc" +version = "3.0.0-alpha.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1f1f8b318a54d18fbe542513331e058f4f8ce6502e542e057c50c7e5e803fdab" +dependencies = [ + "anyhow", + "log", + "thiserror", + "which 4.4.0", +] + +[[package]] +name = "quest" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "556af5f5c953a2ee13f45753e581a38f9778e6551bc3ccc56d90b14628fe59d8" +dependencies = [ + "cfg-if 0.1.10", + "rpassword 2.1.0", + "tempfile", + "termios", + "winapi 0.3.9", +] + +[[package]] +name = "quick-xml" +version = "0.22.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8533f14c8382aaad0d592c812ac3b826162128b65662331e1127b45c3d18536b" +dependencies = [ + "memchr", +] + +[[package]] +name = "quinn" +version = "0.8.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5b435e71d9bfa0d8889927231970c51fb89c58fa63bffcab117c9c7a41e5ef8f" +dependencies = [ + "bytes", + "futures-channel", + "futures-util", + "fxhash", + "quinn-proto", + "quinn-udp", + "rustls", + "thiserror", + "tokio", + "tracing", + "webpki", +] + +[[package]] +name = "quinn-proto" +version = "0.8.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3fce546b9688f767a57530652488420d419a8b1f44a478b451c3d1ab6d992a55" +dependencies = [ + "bytes", + "fxhash", + "rand 0.8.5", + "ring", + "rustls", + "rustls-native-certs", + "rustls-pemfile 0.2.1", + "slab", + "thiserror", + "tinyvec", + "tracing", + "webpki", +] + +[[package]] +name = "quinn-udp" +version = "0.1.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b07946277141531aea269befd949ed16b2c85a780ba1043244eda0969e538e54" +dependencies = [ + "futures-util", + "libc", + "quinn-proto", + "socket2 0.4.9", + "tokio", + "tracing", +] + +[[package]] +name = "quote" +version = "1.0.26" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4424af4bf778aae2051a77b60283332f386554255d722233d09fbfc7e30da2fc" +dependencies = [ + "proc-macro2", +] + +[[package]] +name = "rand" +version = "0.6.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6d71dacdc3c88c1fde3885a3be3fbab9f35724e6ce99467f7d9c5026132184ca" +dependencies = [ + "autocfg 0.1.8", + "libc", + "rand_chacha 0.1.1", + "rand_core 0.4.2", + "rand_hc", + "rand_isaac", + "rand_jitter", + "rand_os", + "rand_pcg", + "rand_xorshift", + "winapi 0.3.9", +] + +[[package]] +name = "rand" +version = "0.8.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "34af8d1a0e25924bc5b7c43c079c942339d8f0a8b57c39049bef581b46327404" +dependencies = [ + "libc", + "rand_chacha 0.3.1", + "rand_core 0.6.4", +] + +[[package]] +name = "rand_chacha" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "556d3a1ca6600bfcbab7c7c91ccb085ac7fbbcd70e008a98742e7847f4f7bcef" +dependencies = [ + "autocfg 0.1.8", + "rand_core 0.3.1", +] + +[[package]] +name = "rand_chacha" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e6c10a63a0fa32252be49d21e7709d4d4baf8d231c2dbce1eaa8141b9b127d88" +dependencies = [ + "ppv-lite86", + "rand_core 0.6.4", +] + +[[package]] +name = "rand_core" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7a6fdeb83b075e8266dcc8762c22776f6877a63111121f5f8c7411e5be7eed4b" +dependencies = [ + "rand_core 0.4.2", +] + +[[package]] +name = "rand_core" +version = "0.4.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9c33a3c44ca05fa6f1807d8e6743f3824e8509beca625669633be0acbdf509dc" + +[[package]] +name = "rand_core" +version = "0.6.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ec0be4795e2f6a28069bec0b5ff3e2ac9bafc99e6a9a7dc3547996c5c816922c" +dependencies = [ + "getrandom", +] + +[[package]] +name = "rand_hc" +version = "0.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7b40677c7be09ae76218dc623efbf7b18e34bced3f38883af07bb75630a21bc4" +dependencies = [ + "rand_core 0.3.1", +] + +[[package]] +name = "rand_isaac" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ded997c9d5f13925be2a6fd7e66bf1872597f759fd9dd93513dd7e92e5a5ee08" +dependencies = [ + "rand_core 0.3.1", +] + +[[package]] +name = "rand_jitter" +version = "0.1.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1166d5c91dc97b88d1decc3285bb0a99ed84b05cfd0bc2341bdf2d43fc41e39b" +dependencies = [ + "libc", + "rand_core 0.4.2", + "winapi 0.3.9", +] + +[[package]] +name = "rand_os" +version = "0.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7b75f676a1e053fc562eafbb47838d67c84801e38fc1ba459e8f180deabd5071" +dependencies = [ + "cloudabi", + "fuchsia-cprng", + "libc", + "rand_core 0.4.2", + "rdrand", + "winapi 0.3.9", +] + +[[package]] +name = "rand_pcg" +version = "0.1.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "abf9b09b01790cfe0364f52bf32995ea3c39f4d2dd011eac241d2914146d0b44" +dependencies = [ + "autocfg 0.1.8", + "rand_core 0.4.2", +] + +[[package]] +name = "rand_xorshift" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cbf7e9e623549b0e21f6e97cf8ecf247c1a8fd2e8a992ae265314300b2455d5c" +dependencies = [ + "rand_core 0.3.1", +] + +[[package]] +name = "rayon" +version = "1.7.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1d2df5196e37bcc87abebc0053e20787d73847bb33134a69841207dd0a47f03b" +dependencies = [ + "either", + "rayon-core", +] + +[[package]] +name = "rayon-core" +version = "1.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4b8f95bd6966f5c87776639160a66bd8ab9895d9d4ab01ddba9fc60661aebe8d" +dependencies = [ + "crossbeam-channel", + "crossbeam-deque", + "crossbeam-utils", + "num_cpus", +] + +[[package]] +name = "rdev" +version = "0.5.0-2" +source = "git+https://github.com/open-trade/rdev#0ad53987fa6f0e37a7bc000358f71c3802de4e7c" +dependencies = [ + "cocoa", + "core-foundation", + "core-foundation-sys", + "core-graphics", + "enum-map", + "lazy_static", + "libc", + "widestring 1.0.2", + "winapi 0.3.9", + "x11", +] + +[[package]] +name = "rdrand" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "678054eb77286b51581ba43620cc911abf02758c91f93f479767aed0f90458b2" +dependencies = [ + "rand_core 0.3.1", +] + +[[package]] +name = "realfft" +version = "3.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "93d6b8e8f0c6d2234aa58048d7290c60bf92cd36fd2888cd8331c66ad4f2e1d2" +dependencies = [ + "rustfft", +] + +[[package]] +name = "redox_syscall" +version = "0.2.16" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fb5a58c1855b4b6819d59012155603f0b22ad30cad752600aadfcb695265519a" +dependencies = [ + "bitflags", +] + +[[package]] +name = "redox_syscall" +version = "0.3.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "567664f262709473930a4bf9e51bf2ebf3348f2e748ccc50dea20646858f8f29" +dependencies = [ + "bitflags", +] + +[[package]] +name = "redox_users" +version = "0.4.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b033d837a7cf162d7993aded9304e30a83213c648b6e389db233191f891e5c2b" +dependencies = [ + "getrandom", + "redox_syscall 0.2.16", + "thiserror", +] + +[[package]] +name = "regex" +version = "1.7.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8b1f693b24f6ac912f4893ef08244d70b6067480d2f1a46e950c9691e6749d1d" +dependencies = [ + "aho-corasick", + "memchr", + "regex-syntax", +] + +[[package]] +name = "regex-syntax" +version = "0.6.29" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f162c6dd7b008981e4d40210aca20b4bd0f9b60ca9271061b07f78537722f2e1" + +[[package]] +name = "repng" +version = "0.2.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0dd57cd2cb5cc699b3eb4824d654e5a32f3bc013766da4966f71fe94805abbda" +dependencies = [ + "byteorder", + "flate2", +] + +[[package]] +name = "ring" +version = "0.16.20" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3053cf52e236a3ed746dfc745aa9cacf1b791d846bdaf412f60a8d7d6e17c8fc" +dependencies = [ + "cc", + "libc", + "once_cell", + "spin", + "untrusted", + "web-sys", + "winapi 0.3.9", +] + +[[package]] +name = "rpassword" +version = "2.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d37473170aedbe66ffa3ad3726939ba677d83c646ad4fd99e5b4bc38712f45ec" +dependencies = [ + "kernel32-sys", + "libc", + "winapi 0.2.8", +] + +[[package]] +name = "rpassword" +version = "6.0.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2bf099a1888612545b683d2661a1940089f6c2e5a8e38979b2159da876bfd956" +dependencies = [ + "libc", + "serde 1.0.159", + "serde_json 1.0.95", + "winapi 0.3.9", +] + +[[package]] +name = "rubato" +version = "0.12.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cd70209c27d5b08f5528bdc779ea3ffb418954e28987f9f9775c6eac41003f9c" +dependencies = [ + "num-complex", + "num-integer", + "num-traits 0.2.15", + "realfft", +] + +[[package]] +name = "runas" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a620b0994a180cdfa25c0439e6d58c0628272571501880d626ffff58e96a0799" +dependencies = [ + "cc", + "which 3.1.1", +] + +[[package]] +name = "rust-pulsectl" +version = "0.2.12" +source = "git+https://github.com/open-trade/pulsectl#5e68f4c2b7c644fa321984688602d71e8ad0bba3" +dependencies = [ + "libpulse-binding", +] + +[[package]] +name = "rustc-demangle" +version = "0.1.22" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d4a36c42d1873f9a77c53bde094f9664d9891bc604a45b4798fd2c389ed12e5b" + +[[package]] +name = "rustc-hash" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2" + +[[package]] +name = "rustc_version" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bfa0f585226d2e68097d4f95d113b15b83a82e819ab25717ec0590d9584ef366" +dependencies = [ + "semver", +] + +[[package]] +name = "rustdesk" +version = "1.1.9" +dependencies = [ + "android_logger", + "arboard", + "async-trait", + "base64 0.13.1", + "cc", + "cfg-if 1.0.0", + "clap 3.2.23", + "clipboard", + "clipboard-master", + "cocoa", + "core-foundation", + "core-graphics", + "cpal", + "ctrlc", + "dasp", + "dispatch", + "enigo", + "flexi_logger", + "hbb_common", + "hound", + "include_dir", + "lazy_static", + "libc", + "libpulse-binding", + "libpulse-simple-binding", + "mac_address", + "machine-uid", + "magnum-opus", + "num_cpus", + "objc", + "parity-tokio-ipc", + "rdev", + "repng", + "rpassword 6.0.1", + "rubato", + "runas", + "rust-pulsectl", + "samplerate", + "sciter-rs", + "scrap", + "serde 1.0.159", + "serde_derive", + "serde_json 1.0.95", + "sha2", + "sys-locale", + "sysinfo", + "systray", + "tray-item", + "uuid", + "virtual_display", + "whoami", + "winapi 0.3.9", + "windows-service", + "winreg 0.10.1", + "winres", +] + +[[package]] +name = "rustfft" +version = "6.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e17d4f6cbdb180c9f4b2a26bbf01c4e647f1e1dea22fe8eb9db54198b32f9434" +dependencies = [ + "num-complex", + "num-integer", + "num-traits 0.2.15", + "primal-check", + "strength_reduce", + "transpose", + "version_check", +] + +[[package]] +name = "rustix" +version = "0.37.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "85597d61f83914ddeba6a47b3b8ffe7365107221c2e557ed94426489fefb5f77" +dependencies = [ + "bitflags", + "errno", + "io-lifetimes", + "libc", + "linux-raw-sys", + "windows-sys 0.48.0", +] + +[[package]] +name = "rustls" +version = "0.20.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fff78fc74d175294f4e83b28343315ffcfb114b156f0185e9741cb5570f50e2f" +dependencies = [ + "ring", + "sct", + "webpki", +] + +[[package]] +name = "rustls-native-certs" +version = "0.6.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0167bac7a9f490495f3c33013e7722b53cb087ecbe082fb0c6387c96f634ea50" +dependencies = [ + "openssl-probe", + "rustls-pemfile 1.0.2", + "schannel", + "security-framework", +] + +[[package]] +name = "rustls-pemfile" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5eebeaeb360c87bfb72e84abdb3447159c0eaececf1bef2aecd65a8be949d1c9" +dependencies = [ + "base64 0.13.1", +] + +[[package]] +name = "rustls-pemfile" +version = "1.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d194b56d58803a43635bdc398cd17e383d6f71f9182b9a192c127ca42494a59b" +dependencies = [ + "base64 0.21.0", +] + +[[package]] +name = "rustversion" +version = "1.0.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4f3208ce4d8448b3f3e7d168a73f5e0c43a61e32930de3bceeccedb388b6bf06" + +[[package]] +name = "ryu" +version = "1.0.13" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f91339c0467de62360649f8d3e185ca8de4224ff281f66000de5eb2a77a79041" + +[[package]] +name = "same-file" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "93fc1dc3aaa9bfed95e02e6eadabb4baf7e3078b0bd1b4d7b6b0b68378900502" +dependencies = [ + "winapi-util", +] + +[[package]] +name = "samplerate" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e032b2b24715c4f982f483ea3abdb3c9ba444d9f63e87b2843d6f998f5ba2698" +dependencies = [ + "libsamplerate-sys", +] + +[[package]] +name = "schannel" +version = "0.1.21" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "713cfb06c7059f3588fb8044c0fad1d09e3c01d225e25b9220dbfdcf16dbb1b3" +dependencies = [ + "windows-sys 0.42.0", +] + +[[package]] +name = "sciter-rs" +version = "0.5.57" +source = "git+https://github.com/open-trade/rust-sciter?branch=dyn#82025b9ba77d5ae14543009444033036dbe25917" +dependencies = [ + "lazy_static", + "libc", + "objc", + "objc-foundation", +] + +[[package]] +name = "scopeguard" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d29ab0c6d3fc0ee92fe66e2d99f700eab17a8d57d1c1d3b748380fb20baa78cd" + +[[package]] +name = "scrap" +version = "0.5.0" +dependencies = [ + "bindgen 0.59.2", + "block", + "cfg-if 1.0.0", + "dbus", + "docopt", + "gstreamer", + "gstreamer-app", + "gstreamer-video", + "libc", + "num_cpus", + "quest", + "repng", + "serde 1.0.159", + "target_build_utils", + "tracing", + "vcpkg", + "webm", + "winapi 0.3.9", +] + +[[package]] +name = "scratch" +version = "1.0.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1792db035ce95be60c3f8853017b3999209281c24e2ba5bc8e59bf97a0c590c1" + +[[package]] +name = "sct" +version = "0.7.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d53dcdb7c9f8158937a7981b48accfd39a43af418591a5d008c7b22b5e1b7ca4" +dependencies = [ + "ring", + "untrusted", +] + +[[package]] +name = "security-framework" +version = "2.8.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a332be01508d814fed64bf28f798a146d73792121129962fdf335bb3c49a4254" +dependencies = [ + "bitflags", + "core-foundation", + "core-foundation-sys", + "libc", + "security-framework-sys", +] + +[[package]] +name = "security-framework-sys" +version = "2.8.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "31c9bb296072e961fcbd8853511dd39c2d8be2deb1e17c6860b1d30732b323b4" +dependencies = [ + "core-foundation-sys", + "libc", +] + +[[package]] +name = "semver" +version = "1.0.17" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bebd363326d05ec3e2f532ab7660680f3b02130d780c299bca73469d521bc0ed" + +[[package]] +name = "serde" +version = "0.9.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "34b623917345a631dc9608d5194cc206b3fe6c3554cd1c75b937e55e285254af" + +[[package]] +name = "serde" +version = "1.0.159" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3c04e8343c3daeec41f58990b9d77068df31209f2af111e059e9fe9646693065" +dependencies = [ + "serde_derive", +] + +[[package]] +name = "serde_derive" +version = "1.0.159" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4c614d17805b093df4b147b51339e7e44bf05ef59fba1e45d83500bcfb4d8585" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.13", +] + +[[package]] +name = "serde_json" +version = "0.9.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad8bcf487be7d2e15d3d543f04312de991d631cfe1b43ea0ade69e6a8a5b16a1" +dependencies = [ + "dtoa", + "itoa 0.3.4", + "num-traits 0.1.43", + "serde 0.9.15", +] + +[[package]] +name = "serde_json" +version = "1.0.95" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d721eca97ac802aa7777b701877c8004d950fc142651367300d21c1cc0194744" +dependencies = [ + "itoa 1.0.6", + "ryu", + "serde 1.0.159", +] + +[[package]] +name = "serde_spanned" +version = "0.6.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0efd8caf556a6cebd3b285caf480045fcc1ac04f6bd786b09a6f11af30c4fcf4" +dependencies = [ + "serde 1.0.159", +] + +[[package]] +name = "sha2" +version = "0.10.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "82e6b795fe2e3b1e845bafcb27aa35405c4d47cdfc92af5fc8d3002f76cebdc0" +dependencies = [ + "cfg-if 1.0.0", + "cpufeatures", + "digest", +] + +[[package]] +name = "shlex" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "43b2853a4d09f215c24cc5489c992ce46052d359b5109343cbafbf26bc62f8a3" + +[[package]] +name = "signal-hook-registry" +version = "1.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d8229b473baa5980ac72ef434c4415e70c4b5e71b423043adb4ba059f89c99a1" +dependencies = [ + "libc", +] + +[[package]] +name = "signature" +version = "1.6.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "74233d3b3b2f6d4b006dc19dee745e73e2a6bfb6f93607cd3b02bd5b00797d7c" + +[[package]] +name = "siphasher" +version = "0.2.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0b8de496cf83d4ed58b6be86c3a275b8602f6ffe98d3024a869e124147a9a3ac" + +[[package]] +name = "slab" +version = "0.4.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6528351c9bc8ab22353f9d776db39a20288e8d6c37ef8cfe3317cf875eecfc2d" +dependencies = [ + "autocfg 1.1.0", +] + +[[package]] +name = "smallvec" +version = "1.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a507befe795404456341dfab10cef66ead4c041f62b8b11bbb92bffe5d0953e0" + +[[package]] +name = "socket2" +version = "0.3.19" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "122e570113d28d773067fab24266b66753f6ea915758651696b6e35e49f88d6e" +dependencies = [ + "cfg-if 1.0.0", + "libc", + "winapi 0.3.9", +] + +[[package]] +name = "socket2" +version = "0.4.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "64a4a911eed85daf18834cfaa86a79b7d266ff93ff5ba14005426219480ed662" +dependencies = [ + "libc", + "winapi 0.3.9", +] + +[[package]] +name = "sodiumoxide" +version = "0.2.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e26be3acb6c2d9a7aac28482586a7856436af4cfe7100031d219de2d2ecb0028" +dependencies = [ + "ed25519", + "libc", + "libsodium-sys", + "serde 1.0.159", +] + +[[package]] +name = "spin" +version = "0.5.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6e63cff320ae2c57904679ba7cb63280a3dc4613885beafb148ee7bf9aa9042d" + +[[package]] +name = "static_assertions" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f" + +[[package]] +name = "stdweb" +version = "0.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ef5430c8e36b713e13b48a9f709cc21e046723fe44ce34587b73a830203b533e" + +[[package]] +name = "str-buf" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9e08d8363704e6c71fc928674353e6b7c23dcea9d82d7012c8faf2a3a025f8d0" + +[[package]] +name = "strength_reduce" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fe895eb47f22e2ddd4dabc02bce419d2e643c8e3b585c78158b349195bc24d82" + +[[package]] +name = "strsim" +version = "0.8.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8ea5119cdb4c55b55d432abb513a0429384878c15dde60cc77b1c99de1a95a6a" + +[[package]] +name = "strsim" +version = "0.10.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "73473c0e59e6d5812c5dfe2a064a6444949f089e20eec9a2e5506596494e4623" + +[[package]] +name = "strum" +version = "0.18.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "57bd81eb48f4c437cadc685403cad539345bf703d78e63707418431cecd4522b" + +[[package]] +name = "strum_macros" +version = "0.18.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "87c85aa3f8ea653bfd3ddf25f7ee357ee4d204731f6aa9ad04002306f6e2774c" +dependencies = [ + "heck 0.3.3", + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "syn" +version = "1.0.109" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "72b64191b275b66ffe2469e8af2c1cfe3bafa67b529ead792a6d0160888b4237" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "syn" +version = "2.0.13" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4c9da457c5285ac1f936ebd076af6dac17a61cfe7826f2076b4d015cf47bc8ec" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "synstructure" +version = "0.12.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f36bdaa60a83aca3921b5259d5400cbf5e90fc51931376a9bd4a0eb79aa7210f" +dependencies = [ + "proc-macro2", + "quote", + "syn 1.0.109", + "unicode-xid", +] + +[[package]] +name = "sys-locale" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f8a11bd9c338fdba09f7881ab41551932ad42e405f61d01e8406baea71c07aee" +dependencies = [ + "js-sys", + "libc", + "wasm-bindgen", + "web-sys", + "windows-sys 0.45.0", +] + +[[package]] +name = "sysinfo" +version = "0.23.13" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3977ec2e0520829be45c8a2df70db2bf364714d8a748316a10c3c35d4d2b01c9" +dependencies = [ + "cfg-if 1.0.0", + "core-foundation-sys", + "libc", + "ntapi", + "once_cell", + "rayon", + "winapi 0.3.9", +] + +[[package]] +name = "system-deps" +version = "1.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0f3ecc17269a19353b3558b313bba738b25d82993e30d62a18406a24aba4649b" +dependencies = [ + "heck 0.3.3", + "pkg-config", + "strum", + "strum_macros", + "thiserror", + "toml 0.5.11", + "version-compare 0.0.10", +] + +[[package]] +name = "system-deps" +version = "6.0.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "555fc8147af6256f3931a36bb83ad0023240ce9cf2b319dec8236fd1f220b05f" +dependencies = [ + "cfg-expr", + "heck 0.4.1", + "pkg-config", + "toml 0.7.3", + "version-compare 0.1.1", +] + +[[package]] +name = "systray" +version = "0.4.1" +source = "git+https://github.com/liyue201/systray-rs#84cca4b4171661bc6c4d1ba5aaa2320ff8e085aa" +dependencies = [ + "glib 0.10.3", + "gtk 0.9.2", + "libappindicator 0.6.1", + "libc", + "log", + "winapi 0.3.9", +] + +[[package]] +name = "target_build_utils" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "013d134ae4a25ee744ad6129db589018558f620ddfa44043887cdd45fa08e75c" +dependencies = [ + "phf", + "phf_codegen", + "serde_json 0.9.10", +] + +[[package]] +name = "tempfile" +version = "3.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b9fbec84f381d5795b08656e4912bec604d162bff9291d6189a78f4c8ab87998" +dependencies = [ + "cfg-if 1.0.0", + "fastrand", + "redox_syscall 0.3.5", + "rustix", + "windows-sys 0.45.0", +] + +[[package]] +name = "termcolor" +version = "1.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "be55cf8942feac5c765c2c993422806843c9a9a45d4d5c407ad6dd2ea95eb9b6" +dependencies = [ + "winapi-util", +] + +[[package]] +name = "termios" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "411c5bf740737c7918b8b1fe232dca4dc9f8e754b8ad5e20966814001ed0ac6b" +dependencies = [ + "libc", +] + +[[package]] +name = "textwrap" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d326610f408c7a4eb6f51c37c330e496b08506c9457c9d34287ecc38809fb060" +dependencies = [ + "unicode-width", +] + +[[package]] +name = "textwrap" +version = "0.16.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "222a222a5bfe1bba4a77b45ec488a741b3cb8872e5e499451fd7d0129c9c7c3d" + +[[package]] +name = "thiserror" +version = "1.0.40" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "978c9a314bd8dc99be594bc3c175faaa9794be04a5a5e153caba6915336cebac" +dependencies = [ + "thiserror-impl", +] + +[[package]] +name = "thiserror-impl" +version = "1.0.40" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f9456a42c5b0d803c8cd86e73dd7cc9edd429499f37a3550d286d5e86720569f" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.13", +] + +[[package]] +name = "tiff" +version = "0.6.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9a53f4706d65497df0c4349241deddf35f84cee19c87ed86ea8ca590f4464437" +dependencies = [ + "jpeg-decoder", + "miniz_oxide 0.4.4", + "weezl", +] + +[[package]] +name = "time" +version = "0.3.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c2702e08a7a860f005826c6815dcac101b19b5eb330c27fe4a5928fec1d20ddd" +dependencies = [ + "itoa 1.0.6", + "libc", + "num_threads", + "time-macros", +] + +[[package]] +name = "time-macros" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "42657b1a6f4d817cda8e7a0ace261fe0cc946cf3a80314390b22cc61ae080792" + +[[package]] +name = "tinyvec" +version = "1.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "87cc5ceb3875bb20c2890005a4e226a4651264a5c75edb2421b52861a0a0cb50" +dependencies = [ + "tinyvec_macros", +] + +[[package]] +name = "tinyvec_macros" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1f3ccbac311fea05f86f61904b462b55fb3df8837a366dfc601a0161d0532f20" + +[[package]] +name = "tokio" +version = "1.27.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d0de47a4eecbe11f498978a9b29d792f0d2692d1dd003650c24c76510e3bc001" +dependencies = [ + "autocfg 1.1.0", + "bytes", + "libc", + "mio 0.8.6", + "num_cpus", + "parking_lot 0.12.1", + "pin-project-lite", + "signal-hook-registry", + "socket2 0.4.9", + "tokio-macros", + "windows-sys 0.45.0", +] + +[[package]] +name = "tokio-macros" +version = "2.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "61a573bdc87985e9d6ddeed1b3d864e8a302c847e40d647746df2f1de209d1ce" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.13", +] + +[[package]] +name = "tokio-socks" +version = "0.5.1-1" +source = "git+https://github.com/open-trade/tokio-socks#7034e79263ce25c348be072808d7601d82cd892d" +dependencies = [ + "bytes", + "either", + "futures-core", + "futures-sink", + "futures-util", + "pin-project", + "thiserror", + "tokio", + "tokio-util 0.7.7", +] + +[[package]] +name = "tokio-util" +version = "0.6.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "36943ee01a6d67977dd3f84a5a1d2efeb4ada3a1ae771cadfaa535d9d9fc6507" +dependencies = [ + "bytes", + "futures-core", + "futures-io", + "futures-sink", + "log", + "pin-project-lite", + "slab", + "tokio", +] + +[[package]] +name = "tokio-util" +version = "0.7.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5427d89453009325de0d8f342c9490009f76e999cb7672d77e46267448f7e6b2" +dependencies = [ + "bytes", + "futures-core", + "futures-sink", + "pin-project-lite", + "tokio", + "tracing", +] + +[[package]] +name = "toml" +version = "0.5.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f4f7f0dd8d50a853a531c426359045b1998f04219d88799810762cd4ad314234" +dependencies = [ + "serde 1.0.159", +] + +[[package]] +name = "toml" +version = "0.7.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b403acf6f2bb0859c93c7f0d967cb4a75a7ac552100f9322faf64dc047669b21" +dependencies = [ + "serde 1.0.159", + "serde_spanned", + "toml_datetime", + "toml_edit", +] + +[[package]] +name = "toml_datetime" +version = "0.6.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3ab8ed2edee10b50132aed5f331333428b011c99402b5a534154ed15746f9622" +dependencies = [ + "serde 1.0.159", +] + +[[package]] +name = "toml_edit" +version = "0.19.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "239410c8609e8125456927e6707163a3b1fdb40561e4b803bc041f466ccfdc13" +dependencies = [ + "indexmap", + "serde 1.0.159", + "serde_spanned", + "toml_datetime", + "winnow", +] + +[[package]] +name = "tracing" +version = "0.1.37" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8ce8c33a8d48bd45d624a6e523445fd21ec13d3653cd51f681abf67418f54eb8" +dependencies = [ + "cfg-if 1.0.0", + "pin-project-lite", + "tracing-attributes", + "tracing-core", +] + +[[package]] +name = "tracing-attributes" +version = "0.1.23" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4017f8f45139870ca7e672686113917c71c7a6e02d4924eda67186083c03081a" +dependencies = [ + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "tracing-core" +version = "0.1.30" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "24eb03ba0eab1fd845050058ce5e616558e8f8d8fca633e6b163fe25c797213a" +dependencies = [ + "once_cell", +] + +[[package]] +name = "transpose" +version = "0.2.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e6522d49d03727ffb138ae4cbc1283d3774f0d10aa7f9bf52e6784c45daf9b23" +dependencies = [ + "num-integer", + "strength_reduce", +] + +[[package]] +name = "tray-item" +version = "0.7.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0914b62e00e8f51241806cb9f9c4ea6b10c75d94cae02c89278de6f4b98c7d0f" +dependencies = [ + "cocoa", + "core-graphics", + "gtk 0.15.5", + "libappindicator 0.7.1", + "libc", + "objc", + "objc-foundation", + "objc_id", + "padlock", + "winapi 0.3.9", +] + +[[package]] +name = "typenum" +version = "1.16.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "497961ef93d974e23eb6f433eb5fe1b7930b659f06d12dec6fc44a8f554c0bba" + +[[package]] +name = "unicode-ident" +version = "1.0.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e5464a87b239f13a63a501f2701565754bae92d243d4bb7eb12f6d57d2269bf4" + +[[package]] +name = "unicode-segmentation" +version = "1.10.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1dd624098567895118886609431a7c3b8f516e41d30e0643f03d94592a147e36" + +[[package]] +name = "unicode-width" +version = "0.1.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c0edd1e5b14653f783770bce4a4dabb4a5108a5370a5f5d8cfe8710c361f6c8b" + +[[package]] +name = "unicode-xid" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f962df74c8c05a667b5ee8bcf162993134c104e96440b663c8daa176dc772d8c" + +[[package]] +name = "untrusted" +version = "0.7.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a156c684c91ea7d62626509bce3cb4e1d9ed5c4d978f7b4352658f96a4c26b4a" + +[[package]] +name = "uuid" +version = "1.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1674845326ee10d37ca60470760d4288a6f80f304007d92e5c53bab78c9cfd79" +dependencies = [ + "getrandom", +] + +[[package]] +name = "vcpkg" +version = "0.2.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "accd4ea62f7bb7a82fe23066fb0957d48ef677f6eeb8215f372f52e48bb32426" + +[[package]] +name = "vec_map" +version = "0.8.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f1bddf1187be692e79c5ffeab891132dfb0f236ed36a43c7ed39f1165ee20191" + +[[package]] +name = "version-compare" +version = "0.0.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d63556a25bae6ea31b52e640d7c41d1ab27faba4ccb600013837a3d0b3994ca1" + +[[package]] +name = "version-compare" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "579a42fc0b8e0c63b76519a339be31bed574929511fa53c1a3acae26eb258f29" + +[[package]] +name = "version_check" +version = "0.9.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "49874b5167b65d7193b8aba1567f5c7d93d001cafc34600cee003eda787e483f" + +[[package]] +name = "virtual_display" +version = "0.1.0" +dependencies = [ + "cc", + "hbb_common", + "lazy_static", + "serde 1.0.159", + "serde_derive", + "thiserror", +] + +[[package]] +name = "walkdir" +version = "2.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "36df944cda56c7d8d8b7496af378e6b16de9284591917d307c9b4d313c44e698" +dependencies = [ + "same-file", + "winapi-util", +] + +[[package]] +name = "wasi" +version = "0.11.0+wasi-snapshot-preview1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9c8d87e72b64a3b4db28d11ce29237c246188f4f51057d65a7eab63b7987e423" + +[[package]] +name = "wasm-bindgen" +version = "0.2.84" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "31f8dcbc21f30d9b8f2ea926ecb58f6b91192c17e9d33594b3df58b2007ca53b" +dependencies = [ + "cfg-if 1.0.0", + "wasm-bindgen-macro", +] + +[[package]] +name = "wasm-bindgen-backend" +version = "0.2.84" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "95ce90fd5bcc06af55a641a86428ee4229e44e07033963a2290a8e241607ccb9" +dependencies = [ + "bumpalo", + "log", + "once_cell", + "proc-macro2", + "quote", + "syn 1.0.109", + "wasm-bindgen-shared", +] + +[[package]] +name = "wasm-bindgen-macro" +version = "0.2.84" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4c21f77c0bedc37fd5dc21f897894a5ca01e7bb159884559461862ae90c0b4c5" +dependencies = [ + "quote", + "wasm-bindgen-macro-support", +] + +[[package]] +name = "wasm-bindgen-macro-support" +version = "0.2.84" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2aff81306fcac3c7515ad4e177f521b5c9a15f2b08f4e32d823066102f35a5f6" +dependencies = [ + "proc-macro2", + "quote", + "syn 1.0.109", + "wasm-bindgen-backend", + "wasm-bindgen-shared", +] + +[[package]] +name = "wasm-bindgen-shared" +version = "0.2.84" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0046fef7e28c3804e5e38bfa31ea2a0f73905319b677e57ebe37e49358989b5d" + +[[package]] +name = "web-sys" +version = "0.3.61" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e33b99f4b23ba3eec1a53ac264e35a755f00e966e0065077d6027c0f575b0b97" +dependencies = [ + "js-sys", + "wasm-bindgen", +] + +[[package]] +name = "webm" +version = "1.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ecb047148a12ef1fd8ab26302bca7e82036f005c3073b48e17cc1b44ec577136" +dependencies = [ + "webm-sys", +] + +[[package]] +name = "webm-sys" +version = "1.0.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0ded6ec82ccf51fe265b0b2b1579cac839574ed910c17baac58e807f8a9de7f3" +dependencies = [ + "cc", +] + +[[package]] +name = "webpki" +version = "0.22.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f095d78192e208183081cc07bc5515ef55216397af48b873e5edcd72637fa1bd" +dependencies = [ + "ring", + "untrusted", +] + +[[package]] +name = "weezl" +version = "0.1.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9193164d4de03a926d909d3bc7c30543cecb35400c02114792c2cae20d5e2dbb" + +[[package]] +name = "which" +version = "3.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d011071ae14a2f6671d0b74080ae0cd8ebf3a6f8c9589a2cd45f23126fe29724" +dependencies = [ + "failure", + "libc", +] + +[[package]] +name = "which" +version = "4.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2441c784c52b289a054b7201fc93253e288f094e2f4be9058343127c4226a269" +dependencies = [ + "either", + "libc", + "once_cell", +] + +[[package]] +name = "whoami" +version = "1.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2c70234412ca409cc04e864e89523cb0fc37f5e1344ebed5a3ebf4192b6b9f68" +dependencies = [ + "wasm-bindgen", + "web-sys", +] + +[[package]] +name = "widestring" +version = "0.4.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c168940144dd21fd8046987c16a46a33d5fc84eec29ef9dcddc2ac9e31526b7c" + +[[package]] +name = "widestring" +version = "1.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "653f141f39ec16bba3c5abe400a0c60da7468261cc2cbf36805022876bc721a8" + +[[package]] +name = "winapi" +version = "0.2.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "167dc9d6949a9b857f3451275e911c3f44255842c1f7a76f33c55103a909087a" + +[[package]] +name = "winapi" +version = "0.3.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5c839a674fcd7a98952e593242ea400abe93992746761e38641405d28b00f419" +dependencies = [ + "winapi-i686-pc-windows-gnu", + "winapi-x86_64-pc-windows-gnu", +] + +[[package]] +name = "winapi-build" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2d315eee3b34aca4797b2da6b13ed88266e6d612562a0c46390af8299fc699bc" + +[[package]] +name = "winapi-i686-pc-windows-gnu" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ac3b87c63620426dd9b991e5ce0329eff545bccbbb34f3be09ff6fb6ab51b7b6" + +[[package]] +name = "winapi-util" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "70ec6ce85bb158151cae5e5c87f95a8e97d2c0c4b001223f33a334e3ce5de178" +dependencies = [ + "winapi 0.3.9", +] + +[[package]] +name = "winapi-wsapoll" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "44c17110f57155602a80dca10be03852116403c9ff3cd25b079d666f2aa3df6e" +dependencies = [ + "winapi 0.3.9", +] + +[[package]] +name = "winapi-x86_64-pc-windows-gnu" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "712e227841d057c1ee1cd2fb22fa7e5a5461ae8e48fa2ca79ec42cfc1931183f" + +[[package]] +name = "windows" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e686886bc078bc1b0b600cac0147aadb815089b6e4da64016cbd754b6342700f" +dependencies = [ + "windows-targets 0.48.0", +] + +[[package]] +name = "windows-service" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0c643e10139d127d30d6d753398c8a6f0a43532e8370f6c9d29ebbff29b984ab" +dependencies = [ + "bitflags", + "err-derive", + "widestring 0.4.3", + "winapi 0.3.9", +] + +[[package]] +name = "windows-sys" +version = "0.28.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "82ca39602d5cbfa692c4b67e3bcbb2751477355141c1ed434c94da4186836ff6" +dependencies = [ + "windows_aarch64_msvc 0.28.0", + "windows_i686_gnu 0.28.0", + "windows_i686_msvc 0.28.0", + "windows_x86_64_gnu 0.28.0", + "windows_x86_64_msvc 0.28.0", +] + +[[package]] +name = "windows-sys" +version = "0.42.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5a3e1820f08b8513f676f7ab6c1f99ff312fb97b553d30ff4dd86f9f15728aa7" +dependencies = [ + "windows_aarch64_gnullvm 0.42.2", + "windows_aarch64_msvc 0.42.2", + "windows_i686_gnu 0.42.2", + "windows_i686_msvc 0.42.2", + "windows_x86_64_gnu 0.42.2", + "windows_x86_64_gnullvm 0.42.2", + "windows_x86_64_msvc 0.42.2", +] + +[[package]] +name = "windows-sys" +version = "0.45.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "75283be5efb2831d37ea142365f009c02ec203cd29a3ebecbc093d52315b66d0" +dependencies = [ + "windows-targets 0.42.2", +] + +[[package]] +name = "windows-sys" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "677d2418bec65e3338edb076e806bc1ec15693c5d0104683f2efe857f61056a9" +dependencies = [ + "windows-targets 0.48.0", +] + +[[package]] +name = "windows-targets" +version = "0.42.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8e5180c00cd44c9b1c88adb3693291f1cd93605ded80c250a75d472756b4d071" +dependencies = [ + "windows_aarch64_gnullvm 0.42.2", + "windows_aarch64_msvc 0.42.2", + "windows_i686_gnu 0.42.2", + "windows_i686_msvc 0.42.2", + "windows_x86_64_gnu 0.42.2", + "windows_x86_64_gnullvm 0.42.2", + "windows_x86_64_msvc 0.42.2", +] + +[[package]] +name = "windows-targets" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7b1eb6f0cd7c80c79759c929114ef071b87354ce476d9d94271031c0497adfd5" +dependencies = [ + "windows_aarch64_gnullvm 0.48.0", + "windows_aarch64_msvc 0.48.0", + "windows_i686_gnu 0.48.0", + "windows_i686_msvc 0.48.0", + "windows_x86_64_gnu 0.48.0", + "windows_x86_64_gnullvm 0.48.0", + "windows_x86_64_msvc 0.48.0", +] + +[[package]] +name = "windows-win" +version = "2.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8d4243ec23afe4e9b4e668b3c0a0e973f1b8265f6a46223cfcbc16fd267480c0" +dependencies = [ + "winapi 0.3.9", +] + +[[package]] +name = "windows_aarch64_gnullvm" +version = "0.42.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "597a5118570b68bc08d8d59125332c54f1ba9d9adeedeef5b99b02ba2b0698f8" + +[[package]] +name = "windows_aarch64_gnullvm" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "91ae572e1b79dba883e0d315474df7305d12f569b400fcf90581b06062f7e1bc" + +[[package]] +name = "windows_aarch64_msvc" +version = "0.28.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "52695a41e536859d5308cc613b4a022261a274390b25bd29dfff4bf08505f3c2" + +[[package]] +name = "windows_aarch64_msvc" +version = "0.42.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e08e8864a60f06ef0d0ff4ba04124db8b0fb3be5776a5cd47641e942e58c4d43" + +[[package]] +name = "windows_aarch64_msvc" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b2ef27e0d7bdfcfc7b868b317c1d32c641a6fe4629c171b8928c7b08d98d7cf3" + +[[package]] +name = "windows_i686_gnu" +version = "0.28.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f54725ac23affef038fecb177de6c9bf065787c2f432f79e3c373da92f3e1d8a" + +[[package]] +name = "windows_i686_gnu" +version = "0.42.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c61d927d8da41da96a81f029489353e68739737d3beca43145c8afec9a31a84f" + +[[package]] +name = "windows_i686_gnu" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "622a1962a7db830d6fd0a69683c80a18fda201879f0f447f065a3b7467daa241" + +[[package]] +name = "windows_i686_msvc" +version = "0.28.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "51d5158a43cc43623c0729d1ad6647e62fa384a3d135fd15108d37c683461f64" + +[[package]] +name = "windows_i686_msvc" +version = "0.42.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "44d840b6ec649f480a41c8d80f9c65108b92d89345dd94027bfe06ac444d1060" + +[[package]] +name = "windows_i686_msvc" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4542c6e364ce21bf45d69fdd2a8e455fa38d316158cfd43b3ac1c5b1b19f8e00" + +[[package]] +name = "windows_x86_64_gnu" +version = "0.28.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bc31f409f565611535130cfe7ee8e6655d3fa99c1c61013981e491921b5ce954" + +[[package]] +name = "windows_x86_64_gnu" +version = "0.42.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8de912b8b8feb55c064867cf047dda097f92d51efad5b491dfb98f6bbb70cb36" + +[[package]] +name = "windows_x86_64_gnu" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ca2b8a661f7628cbd23440e50b05d705db3686f894fc9580820623656af974b1" + +[[package]] +name = "windows_x86_64_gnullvm" +version = "0.42.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "26d41b46a36d453748aedef1486d5c7a85db22e56aff34643984ea85514e94a3" + +[[package]] +name = "windows_x86_64_gnullvm" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7896dbc1f41e08872e9d5e8f8baa8fdd2677f29468c4e156210174edc7f7b953" + +[[package]] +name = "windows_x86_64_msvc" +version = "0.28.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3f2b8c7cbd3bfdddd9ab98769f9746a7fad1bca236554cd032b78d768bc0e89f" + +[[package]] +name = "windows_x86_64_msvc" +version = "0.42.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9aec5da331524158c6d1a4ac0ab1541149c0b9505fde06423b02f5ef0106b9f0" + +[[package]] +name = "windows_x86_64_msvc" +version = "0.48.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1a515f5799fe4961cb532f983ce2b23082366b898e52ffbce459c86f67c8378a" + +[[package]] +name = "winnow" +version = "0.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ae8970b36c66498d8ff1d66685dc86b91b29db0c7739899012f63a63814b4b28" +dependencies = [ + "memchr", +] + +[[package]] +name = "winreg" +version = "0.6.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b2986deb581c4fe11b621998a5e53361efe6b48a151178d0cd9eeffa4dc6acc9" +dependencies = [ + "winapi 0.3.9", +] + +[[package]] +name = "winreg" +version = "0.10.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "80d0f4e272c85def139476380b12f9ac60926689dd2e01d4923222f40580869d" +dependencies = [ + "winapi 0.3.9", +] + +[[package]] +name = "winres" +version = "0.1.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b68db261ef59e9e52806f688020631e987592bd83619edccda9c47d42cde4f6c" +dependencies = [ + "toml 0.5.11", +] + +[[package]] +name = "ws2_32-sys" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d59cefebd0c892fa2dd6de581e937301d8552cb44489cdff035c6187cb63fa5e" +dependencies = [ + "winapi 0.2.8", + "winapi-build", +] + +[[package]] +name = "x11" +version = "2.21.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "502da5464ccd04011667b11c435cb992822c2c0dbde1770c988480d312a0db2e" +dependencies = [ + "libc", + "pkg-config", +] + +[[package]] +name = "x11-clipboard" +version = "0.5.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "473068b7b80ac86a18328824f1054e5e007898c47b5bbc281bd7abe32bc3653c" +dependencies = [ + "xcb", +] + +[[package]] +name = "x11rb" +version = "0.9.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6e99be55648b3ae2a52342f9a870c0e138709a3493261ce9b469afe6e4df6d8a" +dependencies = [ + "gethostname", + "nix 0.22.3", + "winapi 0.3.9", + "winapi-wsapoll", +] + +[[package]] +name = "xcb" +version = "0.10.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "771e2b996df720cd1c6dd9ff90f62d91698fd3610cc078388d0564bdd6622a9c" +dependencies = [ + "libc", + "log", + "quick-xml", +] + +[[package]] +name = "zstd" +version = "0.9.2+zstd.1.5.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2390ea1bf6c038c39674f22d95f0564725fc06034a47129179810b2fc58caa54" +dependencies = [ + "zstd-safe", +] + +[[package]] +name = "zstd-safe" +version = "4.1.3+zstd.1.5.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e99d81b99fb3c2c2c794e3fe56c305c63d5173a16a46b5850b07c935ffc7db79" +dependencies = [ + "libc", + "zstd-sys", +] + +[[package]] +name = "zstd-sys" +version = "1.6.2+zstd.1.5.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2daf2f248d9ea44454bfcb2516534e8b8ad2fc91bf818a1885495fc42bc8ac9f" +dependencies = [ + "cc", + "libc", +] diff --git a/pkgs/applications/networking/remote/rustdesk/cargo.patch b/pkgs/applications/networking/remote/rustdesk/cargo.patch index e2124f26171..eb8495f14ca 100644 --- a/pkgs/applications/networking/remote/rustdesk/cargo.patch +++ b/pkgs/applications/networking/remote/rustdesk/cargo.patch @@ -1,38 +1,3 @@ -diff --git a/Cargo.lock b/Cargo.lock -index fb17c7e..ef157e5 100644 ---- a/Cargo.lock -+++ b/Cargo.lock -@@ -2332,10 +2332,10 @@ dependencies = [ - [[package]] - name = "magnum-opus" - version = "0.4.0" --source = "git+https://github.com/open-trade/magnum-opus#3c3d0b86ae95c84930bebffe4bcb03b3bd83342b" -+source = "git+https://github.com/TheRadioGuy/magnum-opus#171e1d021004626f7444d1e39b98f50bc3cb2604" - dependencies = [ -- "bindgen", -- "target_build_utils", -+ "libc", -+ "opusic-sys", - ] - - [[package]] -@@ -2796,6 +2796,16 @@ version = "0.1.5" - source = "registry+https://github.com/rust-lang/crates.io-index" - checksum = "ff011a302c396a5197692431fc1948019154afc178baf7d8e37367442a4601cf" - -+[[package]] -+name = "opusic-sys" -+version = "0.3.6" -+source = "registry+https://github.com/rust-lang/crates.io-index" -+checksum = "5eace752ce07a037241dba8f02c654799f051e431b27028056bcb480e83b54f5" -+dependencies = [ -+ "cmake", -+ "libc", -+] -+ - [[package]] - name = "os_str_bytes" - version = "6.0.0" diff --git a/Cargo.toml b/Cargo.toml index 1b715bd..960e8da 100644 --- a/Cargo.toml diff --git a/pkgs/applications/networking/remote/rustdesk/default.nix b/pkgs/applications/networking/remote/rustdesk/default.nix index fbbef2dc619..bb3b8ffed61 100644 --- a/pkgs/applications/networking/remote/rustdesk/default.nix +++ b/pkgs/applications/networking/remote/rustdesk/default.nix @@ -1,5 +1,4 @@ { lib -, stdenv , fetchFromGitHub , makeDesktopItem , copyDesktopItems @@ -8,7 +7,6 @@ , cmake , yasm , nasm -, zip , pkg-config , clang , gtk3 @@ -43,9 +41,22 @@ rustPlatform.buildRustPackage rec { ./fix-for-rust-1.65.diff ]; - cargoSha256 = "sha256-1OMWEk+DerltF7kwdo4d04rbgIFLHBRq3vZaL7jtrdE="; + LIBCLANG_PATH = "${llvmPackages.libclang.lib}/lib"; - LIBCLANG_PATH="${llvmPackages.libclang.lib}/lib"; + cargoLock = { + lockFile = ./Cargo.lock; + outputHashes = { + "confy-0.4.0" = "sha256-e91cvEixhpPzIthAxzTa3fDY6eCsHUy/eZQAqs7QTDo="; + "parity-tokio-ipc-0.7.3-1" = "sha256-eULJePtBu0iBI3It/bPH0h82Obsb1PJALgwYwrnCFYI="; + "rdev-0.5.0-2" = "sha256-7CEZ2wIM4QAPfY1tGKqXfHplTaxHnccVqFRPjY21Svo="; + "tokio-socks-0.5.1-1" = "sha256-45QQ6FrhGU9uEhbKXTKd/mY6MDumO6p46NmlakdyDQk="; + "libappindicator-0.6.1" = "sha256-JGnnZrcwbh8WJ6+/4bYhfD3HvgF2C7XaaGb6TaMRWdw="; + "magnum-opus-0.4.0" = "sha256-U5uuN4YolOYDnFNbtPpwYefcBDTUUyioui0UCcW8dyo="; + "rust-pulsectl-0.2.12" = "sha256-8jXTspWvjONFcvw9/Z8C43g4BuGZ3rsG32tvLMQbtbM="; + "sciter-rs-0.5.57" = "sha256-ZZnZDhMjK0LjgmK0da1yvB0uoKueLhhhQtzmjoN+1R0="; + "systray-0.4.1" = "sha256-p1PMr/8oS6zHx4+Ng4zCqt0xZ57cq3wAu6/agyWq5Jw="; + }; + }; # Change magnus-opus version to upstream so that it does not use # vcpkg for libopus since it does not work. @@ -55,35 +66,37 @@ rustPlatform.buildRustPackage rec { # Manually simulate a vcpkg installation so that it can link the libaries # properly. - postUnpack = let - vcpkg_target = "x64-linux"; + postUnpack = + let + vcpkg_target = "x64-linux"; - updates_vcpkg_file = writeText "update_vcpkg_rustdesk" + updates_vcpkg_file = writeText "update_vcpkg_rustdesk" + '' + Package : libyuv + Architecture : ${vcpkg_target} + Version : 1.0 + Status : is installed + + Package : libvpx + Architecture : ${vcpkg_target} + Version : 1.0 + Status : is installed + ''; + in '' - Package : libyuv - Architecture : ${vcpkg_target} - Version : 1.0 - Status : is installed + export VCPKG_ROOT="$TMP/vcpkg"; - Package : libvpx - Architecture : ${vcpkg_target} - Version : 1.0 - Status : is installed + mkdir -p $VCPKG_ROOT/.vcpkg-root + mkdir -p $VCPKG_ROOT/installed/${vcpkg_target}/lib + mkdir -p $VCPKG_ROOT/installed/vcpkg/updates + ln -s ${updates_vcpkg_file} $VCPKG_ROOT/installed/vcpkg/status + mkdir -p $VCPKG_ROOT/installed/vcpkg/info + touch $VCPKG_ROOT/installed/vcpkg/info/libyuv_1.0_${vcpkg_target}.list + touch $VCPKG_ROOT/installed/vcpkg/info/libvpx_1.0_${vcpkg_target}.list + + ln -s ${libvpx.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/ + ln -s ${libyuv.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/ ''; - in '' - export VCPKG_ROOT="$TMP/vcpkg"; - - mkdir -p $VCPKG_ROOT/.vcpkg-root - mkdir -p $VCPKG_ROOT/installed/${vcpkg_target}/lib - mkdir -p $VCPKG_ROOT/installed/vcpkg/updates - ln -s ${updates_vcpkg_file} $VCPKG_ROOT/installed/vcpkg/status - mkdir -p $VCPKG_ROOT/installed/vcpkg/info - touch $VCPKG_ROOT/installed/vcpkg/info/libyuv_1.0_${vcpkg_target}.list - touch $VCPKG_ROOT/installed/vcpkg/info/libvpx_1.0_${vcpkg_target}.list - - ln -s ${libvpx.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/ - ln -s ${libyuv.out}/lib/* $VCPKG_ROOT/installed/${vcpkg_target}/lib/ - ''; nativeBuildInputs = [ pkg-config cmake makeWrapper copyDesktopItems yasm nasm clang wrapGAppsHook ]; buildInputs = [ alsa-lib pulseaudio libXfixes libxcb xdotool gtk3 libvpx libopus libXtst libyuv ]; @@ -91,15 +104,22 @@ rustPlatform.buildRustPackage rec { # Checks require an active X display. doCheck = false; - desktopItems = [ (makeDesktopItem { - name = "rustdesk"; - exec = meta.mainProgram; - icon = "rustdesk"; - desktopName = "RustDesk"; - comment = meta.description; - genericName = "Remote Desktop"; - categories = ["Network"]; - }) ]; + desktopItems = [ + (makeDesktopItem { + name = "rustdesk"; + exec = meta.mainProgram; + icon = "rustdesk"; + desktopName = "RustDesk"; + comment = meta.description; + genericName = "Remote Desktop"; + categories = [ "Network" ]; + }) + ]; + + postPatch = '' + rm Cargo.lock + ln -s ${./Cargo.lock} Cargo.lock + ''; # Add static ui resources and libsciter to same folder as binary so that it # can find them. @@ -122,7 +142,7 @@ rustPlatform.buildRustPackage rec { description = "Yet another remote desktop software"; homepage = "https://rustdesk.com"; license = licenses.gpl3Only; - maintainers = with maintainers; [ leixb ]; + maintainers = with maintainers; [ ocfox leixb ]; platforms = [ "x86_64-linux" ]; mainProgram = "rustdesk"; }; From 4f255d1d045b20267f9e391a3b733d043c4d637f Mon Sep 17 00:00:00 2001 From: Ashish SHUKLA Date: Sat, 8 Apr 2023 10:27:12 +0200 Subject: [PATCH 24/32] goawk: 1.21.0 -> 1.22.0 --- pkgs/tools/text/goawk/default.nix | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/pkgs/tools/text/goawk/default.nix b/pkgs/tools/text/goawk/default.nix index a9d71f1faae..a3ecc22b28d 100644 --- a/pkgs/tools/text/goawk/default.nix +++ b/pkgs/tools/text/goawk/default.nix @@ -2,13 +2,13 @@ buildGoModule rec { pname = "goawk"; - version = "1.21.0"; + version = "1.22.0"; src = fetchFromGitHub { owner = "benhoyt"; repo = "goawk"; rev = "v${version}"; - sha256 = "sha256-I6KmNPFD8kkYDyek8lR1ZS7biPA/LYGwJqMoA2fG7Wg="; + sha256 = "sha256-8UoYGAmYmC0hcxLfkNac3flKyPBT/xsykuy+TvVosBI="; }; vendorSha256 = null; From 55fbe3d799155a45663b8cc616f3e610d3a6fc32 Mon Sep 17 00:00:00 2001 From: Michael Weiss Date: Sat, 8 Apr 2023 15:23:08 +0200 Subject: [PATCH 25/32] chromiumBeta: 112.0.5615.49 -> 113.0.5672.24 --- .../networking/browsers/chromium/upstream-info.json | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/pkgs/applications/networking/browsers/chromium/upstream-info.json b/pkgs/applications/networking/browsers/chromium/upstream-info.json index afb638bce8e..12919687fc2 100644 --- a/pkgs/applications/networking/browsers/chromium/upstream-info.json +++ b/pkgs/applications/networking/browsers/chromium/upstream-info.json @@ -19,15 +19,15 @@ } }, "beta": { - "version": "112.0.5615.49", - "sha256": "0hgzbbmz40235binfn3vkkxzvwxilaxg04dclqrz980z7hvkgzfx", - "sha256bin64": "146gd9csj08d1ygwwh6gyqqbi7d34mhv3vv7wv4a8z9hn7jhdifg", + "version": "113.0.5672.24", + "sha256": "1z7yv5lqi1n4ycymkf0kz1v8ig06n430a37ik8hri3a6db8r9lv8", + "sha256bin64": "1y9jaw47dgphqr2l8yc36s7k9lf4qrbmfll1d2d1zdjd5ma3slab", "deps": { "gn": { - "version": "2023-02-17", + "version": "2023-03-18", "url": "https://gn.googlesource.com/gn", - "rev": "b25a2f8c2d33f02082f0f258350f5e22c0973108", - "sha256": "075p4jwk1apvwmqmvhwfw5f669ci7nxwjq9mz5aa2g5lz4fkdm4c" + "rev": "41fef642de70ecdcaaa26be96d56a0398f95abd4", + "sha256": "12w4g2dl58283allclpi1c4i6ih9v2xvdb9hpbmfda12v8lizmlq" } } }, From 3472b48fe67e1c51f508dfcad560f395ab39a3a2 Mon Sep 17 00:00:00 2001 From: Michael Weiss Date: Sat, 8 Apr 2023 15:23:09 +0200 Subject: [PATCH 26/32] chromiumDev: 113.0.5672.24 -> 114.0.5696.0 --- .../networking/browsers/chromium/upstream-info.json | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/pkgs/applications/networking/browsers/chromium/upstream-info.json b/pkgs/applications/networking/browsers/chromium/upstream-info.json index afb638bce8e..cc609a2cd87 100644 --- a/pkgs/applications/networking/browsers/chromium/upstream-info.json +++ b/pkgs/applications/networking/browsers/chromium/upstream-info.json @@ -32,9 +32,9 @@ } }, "dev": { - "version": "113.0.5672.24", - "sha256": "1z7yv5lqi1n4ycymkf0kz1v8ig06n430a37ik8hri3a6db8r9lv8", - "sha256bin64": "0byksvk781gmh5fmjmc77jh19gvkzadf78yr9b4c42las44g4pn4", + "version": "114.0.5696.0", + "sha256": "0hcf971azy3jjsl211qk53b6nk0f4pzf2dwfv3rjh4f6fa3nbwai", + "sha256bin64": "1ssj633vdg3ghd87az28q262ly1fk7rc3k70l1531xvj2030ijrl", "deps": { "gn": { "version": "2023-03-18", From b345e08f5d02ec31687d090c1e61820476925c60 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Sandro=20J=C3=A4ckel?= Date: Sat, 8 Apr 2023 17:38:38 +0200 Subject: [PATCH 27/32] python310Packages.bottle: disable timing sensitive test ____________________________ TestSendFile.test_ims _____________________________ self = def test_ims(self): """ SendFile: If-Modified-Since""" request.environ['HTTP_IF_MODIFIED_SINCE'] = time.strftime("%a, %d %b %Y %H:%M:%S GMT", time.gmtime()) res = static_file(basename, root=root) self.assertEqual(304, res.status_code) self.assertEqual(int(os.stat(__file__).st_mtime), parse_date(res.headers['Last-Modified'])) > self.assertAlmostEqual(int(time.time()), parse_date(res.headers['Date'])) E AssertionError: 1680561734 != 1680561733.0 within 7 places (1.0 difference) test_sendfile.py:78: AssertionError --- pkgs/development/python-modules/bottle/default.nix | 2 ++ 1 file changed, 2 insertions(+) diff --git a/pkgs/development/python-modules/bottle/default.nix b/pkgs/development/python-modules/bottle/default.nix index 3d12ba7e3ae..7a97b999023 100644 --- a/pkgs/development/python-modules/bottle/default.nix +++ b/pkgs/development/python-modules/bottle/default.nix @@ -27,6 +27,8 @@ buildPythonPackage rec { "test_delete_cookie" "test_error" "test_error_in_generator_callback" + # timing sensitive + "test_ims" ]; __darwinAllowLocalNetworking = true; From bc3272f51e3cf4433f6ffe62c8307e92c35fb997 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Bj=C3=B8rn=20Forsman?= Date: Sat, 1 Apr 2023 12:45:26 +0200 Subject: [PATCH 28/32] nixos/installation-device.nix: improve comment about ssh login root is not the only user that can login (user "nixos" can too), so generalize the wording. --- nixos/modules/profiles/installation-device.nix | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/nixos/modules/profiles/installation-device.nix b/nixos/modules/profiles/installation-device.nix index 980720691a4..514ec06645e 100644 --- a/nixos/modules/profiles/installation-device.nix +++ b/nixos/modules/profiles/installation-device.nix @@ -65,8 +65,8 @@ with lib; start the graphical user interface. ''; - # We run sshd by default. Login via root is only possible after adding a - # password via "passwd" or by adding a ssh key to /home/nixos/.ssh/authorized_keys. + # We run sshd by default. Login is only possible after adding a + # password via "passwd" or by adding a ssh key to ~/.ssh/authorized_keys. # The latter one is particular useful if keys are manually added to # installation device for head-less systems i.e. arm boards by manually # mounting the storage in a different system. From 54731a8cea6d8b9d4fcb9ca7cbfc75ae1e07a171 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Bj=C3=B8rn=20Forsman?= Date: Sat, 1 Apr 2023 12:33:27 +0200 Subject: [PATCH 29/32] nixos/installer: update getty help message I think this is clearer. --- nixos/modules/profiles/installation-device.nix | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/nixos/modules/profiles/installation-device.nix b/nixos/modules/profiles/installation-device.nix index 514ec06645e..32884f4b875 100644 --- a/nixos/modules/profiles/installation-device.nix +++ b/nixos/modules/profiles/installation-device.nix @@ -52,9 +52,9 @@ with lib; services.getty.helpLine = '' The "nixos" and "root" accounts have empty passwords. - An ssh daemon is running. You then must set a password - for either "root" or "nixos" with `passwd` or add an ssh key - to /home/nixos/.ssh/authorized_keys be able to login. + To log in over ssh you must set a password for either "nixos" or "root" + with `passwd` (prefix with `sudo` for "root"), or add your public key to + /home/nixos/.ssh/authorized_keys or /root/.ssh/authorized_keys. If you need a wireless connection, type `sudo systemctl start wpa_supplicant` and configure a From 584278eface16359dec7839e001b80a7db7347d9 Mon Sep 17 00:00:00 2001 From: Kasper Kondzielski Date: Sat, 8 Apr 2023 20:16:10 +0200 Subject: [PATCH 30/32] noice-nvim: Remove nvim-notify from noice-nvim dependencies *!This will change behavior for the end user of noice as they will have to add notify plugin manually!* Notify plugin is no longer a hard dependency for noice plugin and there is no need to ship noice with it. In addition to that some plugins will perform auto discovery of plugins that are on the runtime path and in case they find notify they might want to use it. There is no easy way to remove nvim-notify from the runtime path without doing an override. https://github.com/folke/noice.nvim#%EF%B8%8F-requirements --- pkgs/applications/editors/vim/plugins/overrides.nix | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pkgs/applications/editors/vim/plugins/overrides.nix b/pkgs/applications/editors/vim/plugins/overrides.nix index e15d497a80a..c7a4b2fc96e 100644 --- a/pkgs/applications/editors/vim/plugins/overrides.nix +++ b/pkgs/applications/editors/vim/plugins/overrides.nix @@ -706,7 +706,7 @@ self: super: { }); noice-nvim = super.noice-nvim.overrideAttrs(old: { - dependencies = with self; [ nui-nvim nvim-notify ]; + dependencies = with self; [ nui-nvim ]; }); null-ls-nvim = super.null-ls-nvim.overrideAttrs (old: { From 0981043da85f5efdf5057ce52b0f30f5811987a6 Mon Sep 17 00:00:00 2001 From: Adam Joseph Date: Sat, 8 Apr 2023 12:09:07 -0700 Subject: [PATCH 31/32] lyx: fix cross python3 is a nativeBuildInput, not a buildInput --- pkgs/applications/misc/lyx/default.nix | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/pkgs/applications/misc/lyx/default.nix b/pkgs/applications/misc/lyx/default.nix index 921ae570bca..10dbb2c6a74 100644 --- a/pkgs/applications/misc/lyx/default.nix +++ b/pkgs/applications/misc/lyx/default.nix @@ -18,9 +18,9 @@ mkDerivation rec { ''; # LaTeX is used from $PATH, as people often want to have it with extra pkgs - nativeBuildInputs = [ pkg-config makeWrapper ]; + nativeBuildInputs = [ pkg-config makeWrapper python3 ]; buildInputs = [ - qtbase qtsvg python3 file/*for libmagic*/ bc + qtbase qtsvg file/*for libmagic*/ bc hunspell # enchant ]; From 6685004cd5d872185df2ccb758c15809f1ca05ef Mon Sep 17 00:00:00 2001 From: Adam Joseph Date: Wed, 5 Apr 2023 18:58:56 -0700 Subject: [PATCH 32/32] libiio: fix cross compilation --- pkgs/development/libraries/libiio/default.nix | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/pkgs/development/libraries/libiio/default.nix b/pkgs/development/libraries/libiio/default.nix index 815f23a0527..d217fbd8238 100644 --- a/pkgs/development/libraries/libiio/default.nix +++ b/pkgs/development/libraries/libiio/default.nix @@ -37,19 +37,19 @@ stdenv.mkDerivation rec { flex bison pkg-config - ]; + python + ] ++ lib.optional python.isPy3k python.pkgs.setuptools; buildInputs = [ - python libxml2 libusb1 - ] ++ lib.optional python.isPy3k python.pkgs.setuptools - ++ lib.optional avahiSupport avahi + ] ++ lib.optional avahiSupport avahi ++ lib.optional stdenv.isLinux libaio ++ lib.optionals stdenv.isDarwin [ CFNetwork CoreServices ]; cmakeFlags = [ "-DUDEV_RULES_INSTALL_DIR=${placeholder "out"}/lib/udev/rules.d" + "-DPython_EXECUTABLE=${python.pythonForBuild.interpreter}" "-DPYTHON_BINDINGS=on" # osx framework is disabled, # the linux-like directory structure is used for proper output splitting