From e790a98fde7790dfac2a30bb9308c34c8c8bb22f Mon Sep 17 00:00:00 2001 From: Thien Nguyen <58006629+1tnguyen@users.noreply.github.com> Date: Thu, 12 Dec 2024 16:38:51 +1100 Subject: [PATCH 1/6] Fix uninitialized memory issue for the result buffer in kron (#2475) Signed-off-by: Thien Nguyen --- runtime/nvqir/custatevec/CuStateVecCircuitSimulator.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/runtime/nvqir/custatevec/CuStateVecCircuitSimulator.cpp b/runtime/nvqir/custatevec/CuStateVecCircuitSimulator.cpp index 3f4d09bf76..5f4b2f4801 100644 --- a/runtime/nvqir/custatevec/CuStateVecCircuitSimulator.cpp +++ b/runtime/nvqir/custatevec/CuStateVecCircuitSimulator.cpp @@ -216,7 +216,8 @@ class CuStateVecCircuitSimulator void *newDeviceStateVector; HANDLE_CUDA_ERROR(cudaMalloc((void **)&newDeviceStateVector, stateDimension * sizeof(CudaDataType))); - + HANDLE_CUDA_ERROR(cudaMemset(newDeviceStateVector, 0, + stateDimension * sizeof(CudaDataType))); // Place the state data on device. Could be that // we just need the zero state, or the user could have provided one void *otherState; @@ -283,6 +284,8 @@ class CuStateVecCircuitSimulator void *newDeviceStateVector; HANDLE_CUDA_ERROR(cudaMalloc((void **)&newDeviceStateVector, stateDimension * sizeof(CudaDataType))); + HANDLE_CUDA_ERROR(cudaMemset(newDeviceStateVector, 0, + stateDimension * sizeof(CudaDataType))); constexpr int32_t threads_per_block = 256; uint32_t n_blocks = (stateDimension + threads_per_block - 1) / threads_per_block; From aae02de7da304fd1d83b121aad860d3baa28f6ca Mon Sep 17 00:00:00 2001 From: Bettina Heim Date: Thu, 12 Dec 2024 11:23:39 +0100 Subject: [PATCH 2/6] Updating docs version index (#2470) Signed-off-by: Bettina Heim --- .github/workflows/docker_images.yml | 2 +- .../main_divisive_clustering.py | 2 +- .../applications/python/vqe_advanced.ipynb | 2 +- .../examples/python/building_kernels.ipynb | 2 +- .../examples/python/building_kernels.py | 4 ++- .../examples/python/executing_kernels.ipynb | 2 +- docs/sphinx/releases.rst | 31 ++++++++++++++++--- 7 files changed, 35 insertions(+), 10 deletions(-) diff --git a/.github/workflows/docker_images.yml b/.github/workflows/docker_images.yml index 5b4717c570..0e62b2a302 100644 --- a/.github/workflows/docker_images.yml +++ b/.github/workflows/docker_images.yml @@ -694,7 +694,7 @@ jobs: fi image_tag=`docker inspect $cudaq_image --format='{{json .Config.Labels}}' | jq -r '."org.opencontainers.image.version"'` - docs_version="CUDA_QUANTUM_VERSION=${image_tag%-base}" + docs_version="CUDA_QUANTUM_VERSION=$(echo $image_tag | sed -re 's/^(cu[0-9]+-)?(.*)-base$/\2/')" docker image rm $cudaq_image docker image prune --force diff --git a/docs/sphinx/applications/python/divisive_clustering_src/main_divisive_clustering.py b/docs/sphinx/applications/python/divisive_clustering_src/main_divisive_clustering.py index 6ada0c0cde..4065d93755 100644 --- a/docs/sphinx/applications/python/divisive_clustering_src/main_divisive_clustering.py +++ b/docs/sphinx/applications/python/divisive_clustering_src/main_divisive_clustering.py @@ -27,7 +27,7 @@ type=str, choices=["qpp-cpu", "nvidia", "nvidia-mgpu"], help= - "Quantum simulator backend. Default is qpp-cpu. See https://nvidia.github.io/cuda-quantum/0.6.0/using/simulators.html for more options.", + "Quantum simulator backend. Default is qpp-cpu. See https://nvidia.github.io/cuda-quantum for more options.", ) argparser.add_argument( "-d", diff --git a/docs/sphinx/applications/python/vqe_advanced.ipynb b/docs/sphinx/applications/python/vqe_advanced.ipynb index 188df3af45..3135b93829 100644 --- a/docs/sphinx/applications/python/vqe_advanced.ipynb +++ b/docs/sphinx/applications/python/vqe_advanced.ipynb @@ -331,7 +331,7 @@ "name": "stdout", "output_type": "stream", "text": [ - "[\u001b[38;2;255;000;000mwarning\u001b[0m] Target \u001b[38;2;000;000;255mnvidia-mqpu\u001b[0m: \u001b[38;2;000;000;255mThis target is deprecating. Please use the 'nvidia' target with option 'mqpu,fp32' or 'mqpu' (fp32 is the default precision option) by adding the command line option '--target-option mqpu,fp32' or passing it as cudaq.set_target('nvidia', option='mqpu,fp32') in Python. Please refer to CUDA-Q \u001b]8;;https://nvidia.github.io/cuda-quantum/latest/using/backends/platform.html#nvidia-mqpu-platform\u001b\\documentation\u001b]8;;\u001b\\ for more information.\u001b[0m\n" + "[\u001b[38;2;255;000;000mwarning\u001b[0m] Target \u001b[38;2;000;000;255mnvidia-mqpu\u001b[0m: \u001b[38;2;000;000;255mThis target is deprecating. Please use the 'nvidia' target with option 'mqpu,fp32' or 'mqpu' (fp32 is the default precision option) by adding the command line option '--target-option mqpu,fp32' or passing it as cudaq.set_target('nvidia', option='mqpu,fp32') in Python. Please refer to CUDA-Q \u001b]8;;https://nvidia.github.io/cuda-quantum/latest/using/backends/platform\u001b\\documentation\u001b]8;;\u001b\\ for more information.\u001b[0m\n" ] } ], diff --git a/docs/sphinx/examples/python/building_kernels.ipynb b/docs/sphinx/examples/python/building_kernels.ipynb index 366b9f626f..b847b4936e 100644 --- a/docs/sphinx/examples/python/building_kernels.ipynb +++ b/docs/sphinx/examples/python/building_kernels.ipynb @@ -145,7 +145,7 @@ "### Applying Gates\n", "\n", "\n", - "After a kernel is constructed, gates can be applied to start building out a quantum circuit. All the predefined gates in CUDA-Q can be found [here](https://nvidia.github.io/cuda-quantum/latest/api/default_ops.html#unitary-operations-on-qubits).\n", + "After a kernel is constructed, gates can be applied to start building out a quantum circuit. All the predefined gates in CUDA-Q can be found [here](https://nvidia.github.io/cuda-quantum/latest/api/default_ops).\n", "\n", "\n", "Gates can be applied to all qubits in a register:" diff --git a/docs/sphinx/examples/python/building_kernels.py b/docs/sphinx/examples/python/building_kernels.py index 2228d1045b..35583fbb9b 100644 --- a/docs/sphinx/examples/python/building_kernels.py +++ b/docs/sphinx/examples/python/building_kernels.py @@ -100,7 +100,9 @@ def kernel(state: cudaq.State): # ### Applying Gates # # -# After a kernel is constructed, gates can be applied to start building out a quantum circuit. All the predefined gates in CUDA-Q can be found [here](https://nvidia.github.io/cuda-quantum/latest/api/default_ops.html#unitary-operations-on-qubits). +# After a kernel is constructed, gates can be applied to start building out a quantum circuit. +# All the predefined gates in CUDA-Q can be found here: +# https://nvidia.github.io/cuda-quantum/api/default_ops. # # # Gates can be applied to all qubits in a register: diff --git a/docs/sphinx/examples/python/executing_kernels.ipynb b/docs/sphinx/examples/python/executing_kernels.ipynb index 66a7afee3b..0d3250c8c1 100644 --- a/docs/sphinx/examples/python/executing_kernels.ipynb +++ b/docs/sphinx/examples/python/executing_kernels.ipynb @@ -78,7 +78,7 @@ "source": [ "Note that there is a subtle difference between how `sample` is executed with the target device set to a simulator or with the target device set to a QPU. In simulation mode, the quantum state is built once and then sampled $s$ times where $s$ equals the `shots_count`. In hardware execution mode, the quantum state collapses upon measurement and hence needs to be rebuilt over and over again.\n", "\n", - "There are a number of helpful tools that can be found in the API [here](https://nvidia.github.io/cuda-quantum/latest/api/languages/python_api.html#cudaq.SampleResult) to process the `Sample_Result` object produced by `sample`." + "There are a number of helpful tools that can be found in the [API docs](https://nvidia.github.io/cuda-quantum/latest/api/languages/python_api) to process the `Sample_Result` object produced by `sample`." ] }, { diff --git a/docs/sphinx/releases.rst b/docs/sphinx/releases.rst index 6a8a353462..de3e12d78b 100644 --- a/docs/sphinx/releases.rst +++ b/docs/sphinx/releases.rst @@ -4,12 +4,35 @@ CUDA-Q Releases **latest** -The latest version of CUDA-Q is on the main branch of our `GitHub repository `__ and is also available as a Docker image. More information about installing the nightly builds can be found :doc:`here ` +The latest version of CUDA-Q is on the main branch of our `GitHub repository `__ +and is also available as a Docker image. More information about installing the nightly builds can be found +:doc:`here ` - `Docker image (nightly builds) `__ - `Documentation `__ - `Examples `__ +**0.9.0** + +We are very excited to share a new toolset added for modeling and manipulating the dynamics of physical systems. +The new API allows to define and execute a time evolution under arbitrary operators. For more information, take +a look at the `docs `__. +The 0.9.0 release furthermore includes a range of contribution to add new backends to CUDA-Q, including backends +from `Anyon Technologies `__, +`Ferimioniq `__, and +`QuEra Computing `__, +as well as updates to existing backends from `ORCA `__ +and `OQC `__. +We hope you enjoy the new features - also check out our new notebooks and examples to dive into CUDA-Q. + +- `Docker image `__ +- `Python wheel `__ +- `C++ installer `__ +- `Documentation `__ +- `Examples `__ + +The full change log can be found `here `__. + **0.8.0** The 0.8.0 release adds a range of changes to improve the ease of use and performance with CUDA-Q. @@ -17,7 +40,7 @@ The changes listed below highlight some of what we think will be the most useful to know about. While the listed changes do not capture all of the great contributions, we would like to extend many thanks for every contribution, in particular those from external contributors. -- `Docker image `__ +- `Docker image `__ - `Python wheel `__ - `C++ installer `__ - `Documentation `__ @@ -30,7 +53,7 @@ The full change log can be found `here `__. +`here `__. It furthermore adds a range of bug fixes and changes the Python wheel installation instructions. - `Docker image `__ @@ -46,7 +69,7 @@ The full change log can be found `here `, giving you access to our most powerful GPU-accelerated simulators even if you don't have an NVIDIA GPU. With 0.7.0, we have furthermore greatly increased expressiveness of the Python and C++ language frontends. -Check out our `documentation `__ +Check out our `documentation `__ to get started with the new Python syntax support we have added, and `follow our blog `__ to learn more about the new setup and its performance benefits. From 09f5e1f0d89f96c9827f3fb4f731adea01e6c21a Mon Sep 17 00:00:00 2001 From: Eric Schweitz Date: Thu, 12 Dec 2024 10:07:10 -0800 Subject: [PATCH 3/6] Convert static variable to a constexpr. (#2471) Make this static data structure constant for efficiency and to eliminate an initialization ctor. Signed-off-by: Eric Schweitz --- lib/Optimizer/CodeGen/ConvertToQIRProfile.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/lib/Optimizer/CodeGen/ConvertToQIRProfile.cpp b/lib/Optimizer/CodeGen/ConvertToQIRProfile.cpp index 2ae90d302b..58526dc692 100644 --- a/lib/Optimizer/CodeGen/ConvertToQIRProfile.cpp +++ b/lib/Optimizer/CodeGen/ConvertToQIRProfile.cpp @@ -516,7 +516,7 @@ namespace { /// trivial pass only does this preparation work. It performs no analysis and /// does not rewrite function body's, etc. -static const std::vector measurementFunctionNames{ +static constexpr std::array measurementFunctionNames{ cudaq::opt::QIRMeasureBody, cudaq::opt::QIRMeasure, cudaq::opt::QIRMeasureToRegister}; @@ -564,7 +564,7 @@ struct QIRProfilePreparationPass func.getFunctionType().getParams(), module); // Apply irreversible attribute to measurement functions - for (auto &funcName : measurementFunctionNames) { + for (auto *funcName : measurementFunctionNames) { Operation *op = SymbolTable::lookupSymbolIn(module, funcName); auto funcOp = llvm::dyn_cast_if_present(op); if (funcOp) { From 95134fdaeb8f94fc60833182e8af966daeb25c40 Mon Sep 17 00:00:00 2001 From: Pradnya Khalate <148914294+khalatepradnya@users.noreply.github.com> Date: Thu, 12 Dec 2024 11:35:55 -0800 Subject: [PATCH 4/6] Enable nightly integration tests for `infleqtion` target (#2469) Signed-off-by: Pradnya Khalate --- .github/workflows/integration_tests.yml | 46 +++++++++++++++++++++++++ 1 file changed, 46 insertions(+) diff --git a/.github/workflows/integration_tests.yml b/.github/workflows/integration_tests.yml index 55d329b93f..5c5bc28300 100644 --- a/.github/workflows/integration_tests.yml +++ b/.github/workflows/integration_tests.yml @@ -17,6 +17,7 @@ on: options: - nightly - anyon + - infleqtion - ionq - iqm - oqc @@ -650,6 +651,51 @@ jobs: fi shell: bash + - name: Submit to Infleqtion test server + if: (success() || failure()) && (inputs.target == 'infleqtion' || github.event_name == 'schedule' || inputs.target == 'nightly') + run: | + echo "### Submit to Infleqtion server" >> $GITHUB_STEP_SUMMARY + export SUPERSTAQ_API_KEY='${{ secrets.SUPERSTAQ_API_KEY }}' + set +e # Allow script to keep going through errors + test_err_sum=0 + cpp_tests="docs/sphinx/targets/cpp/infleqtion.cpp" + for filename in $cpp_tests; do + [ -e "$filename" ] || echo "::error::Couldn't find file ($filename)" + nvq++ --target infleqtion $filename + test_status=$? + if [ $test_status -eq 0 ]; then + ./a.out + test_status=$? + if [ $test_status -eq 0 ]; then + echo ":white_check_mark: Successfully ran test: $filename" >> $GITHUB_STEP_SUMMARY + else + echo ":x: Test failed (failed to execute): $filename" >> $GITHUB_STEP_SUMMARY + test_err_sum=$((test_err_sum+1)) + fi + else + echo ":x: Test failed (failed to compile): $filename" >> $GITHUB_STEP_SUMMARY + test_err_sum=$((test_err_sum+1)) + fi + done + python_tests="docs/sphinx/targets/python/infleqtion.py" + for filename in $python_tests; do + [ -e "$filename" ] || echo "::error::Couldn't find file ($filename)" + python3 $filename 1> /dev/null + test_status=$? + if [ $test_status -eq 0 ]; then + echo ":white_check_mark: Successfully ran test: $filename" >> $GITHUB_STEP_SUMMARY + else + echo ":x: Test failed (failed to execute): $filename" >> $GITHUB_STEP_SUMMARY + test_err_sum=$((test_err_sum+1)) + fi + done + set -e # Re-enable exit code error checking + if [ ! $test_err_sum -eq 0 ]; then + echo "::error::${test_err_sum} tests failed. See step summary for a list of failures" + exit 1 + fi + shell: bash + - name: Submit to ${{ inputs.target }} # The full set of tests used by this step is currently only supported on # Quantinuum. The other supported tests are tested by the step above. From 105c05a15c5533cf2aadb0ede263f110f44426a0 Mon Sep 17 00:00:00 2001 From: Pradnya Khalate <148914294+khalatepradnya@users.noreply.github.com> Date: Thu, 12 Dec 2024 13:03:46 -0800 Subject: [PATCH 5/6] Enable nightly integration tests for `anyon` target (#2308) Signed-off-by: Pradnya Khalate --- .github/workflows/integration_tests.yml | 12 ++---------- 1 file changed, 2 insertions(+), 10 deletions(-) diff --git a/.github/workflows/integration_tests.yml b/.github/workflows/integration_tests.yml index 5c5bc28300..5619f5baea 100644 --- a/.github/workflows/integration_tests.yml +++ b/.github/workflows/integration_tests.yml @@ -336,11 +336,7 @@ jobs: fi - name: Setup anyon account - # This step is currently bypassed during nightly runs due to - # maintenance. Restore the if check to the original value when - # maintenance is complete. - #if: github.event_name == 'schedule' || inputs.target == 'nightly' || inputs.target == 'anyon' - if: inputs.target == 'anyon' + if: github.event_name == 'schedule' || inputs.target == 'nightly' || inputs.target == 'anyon' run: | curl -X POST --user "${{ secrets.ANYON_USERNAME }}:${{ secrets.ANYON_PASSWORD }}" -H "Content-Type: application/json" https://api.anyon.cloud:5000/login > credentials.json id_token=`cat credentials.json | jq -r '."id_token"'` @@ -349,11 +345,7 @@ jobs: echo "refresh: $refresh_token" >> ~/.anyon_config - name: QIR syntax check (Anyon) - # This step is currently bypassed during nightly runs due to - # maintenance. Restore the if check to the original value when - # maintenance is complete. - #if: github.event_name == 'schedule' || inputs.target == 'nightly' || inputs.target == 'anyon' - if: inputs.target == 'anyon' + if: github.event_name == 'schedule' || inputs.target == 'nightly' || inputs.target == 'anyon' run: | echo "### QIR syntax check (Anyon)" >> $GITHUB_STEP_SUMMARY export CUDAQ_LOG_LEVEL="info" From 2cf6b04d068aa40071a36f169795bebe28c61f3e Mon Sep 17 00:00:00 2001 From: Eric Schweitz Date: Thu, 12 Dec 2024 17:14:17 -0800 Subject: [PATCH 6/6] [core] Move canonical patterns out of tablegen. (#2473) * [core] Move canonical patterns out of tablegen. Fix #476. Eliminate the file Canonical.td and move/rewrite the patterns from that file. Signed-off-by: Eric Schweitz * Fix comments. Signed-off-by: Eric Schweitz * Move OpRewritePatterns to the new file. Signed-off-by: Eric Schweitz --------- Signed-off-by: Eric Schweitz --- .../Optimizer/Dialect/Quake/CMakeLists.txt | 4 - .../Optimizer/Dialect/Quake/Canonical.td | 67 --- lib/Optimizer/Dialect/Quake/CMakeLists.txt | 1 - .../Dialect/Quake/CanonicalPatterns.inc | 489 ++++++++++++++++++ lib/Optimizer/Dialect/Quake/QuakeOps.cpp | 410 +-------------- 5 files changed, 490 insertions(+), 481 deletions(-) delete mode 100644 include/cudaq/Optimizer/Dialect/Quake/Canonical.td create mode 100644 lib/Optimizer/Dialect/Quake/CanonicalPatterns.inc diff --git a/include/cudaq/Optimizer/Dialect/Quake/CMakeLists.txt b/include/cudaq/Optimizer/Dialect/Quake/CMakeLists.txt index d038abd040..6dca96dc83 100644 --- a/include/cudaq/Optimizer/Dialect/Quake/CMakeLists.txt +++ b/include/cudaq/Optimizer/Dialect/Quake/CMakeLists.txt @@ -9,7 +9,3 @@ add_cudaq_dialect(Quake quake) add_cudaq_interface(QuakeInterfaces) add_cudaq_dialect_doc(QuakeDialect quake) - -set(LLVM_TARGET_DEFINITIONS Canonical.td) -mlir_tablegen(Canonical.inc -gen-rewriters) -add_public_tablegen_target(CanonicalIncGen) diff --git a/include/cudaq/Optimizer/Dialect/Quake/Canonical.td b/include/cudaq/Optimizer/Dialect/Quake/Canonical.td deleted file mode 100644 index d7aec89e6f..0000000000 --- a/include/cudaq/Optimizer/Dialect/Quake/Canonical.td +++ /dev/null @@ -1,67 +0,0 @@ -/********************************************************** -*- tablegen -*- *** - * Copyright (c) 2022 - 2024 NVIDIA Corporation & Affiliates. * - * All rights reserved. * - * * - * This source code and the accompanying materials are made available under * - * the terms of the Apache License 2.0 which accompanies this distribution. * - ******************************************************************************/ - -#ifndef NVQPP_OPTIMIZER_DIALECT_QUAKE_CANONICAL -#define NVQPP_OPTIMIZER_DIALECT_QUAKE_CANONICAL - -include "mlir/IR/OpBase.td" -include "mlir/IR/PatternBase.td" -include "mlir/Dialect/Arith/IR/ArithOps.td" -include "cudaq/Optimizer/Dialect/Quake/QuakeOps.td" - -def KnownSizePred : Constraint< - CPred<"$0.getType().isa() && " - "$0.getType().cast().hasSpecifiedSize()">>; - -def UnknownSizePred : Constraint< - CPred<"$0.getType().isa() && " - "!$0.getType().cast().hasSpecifiedSize()">>; - -def createConstantOp : NativeCodeCall< - "$_builder.create($_loc, $0.getType()," - " $_builder.getIntegerAttr($0.getType()," - " $1.getType().cast().getSize()))">; - -// %4 = quake.veq_size %3 : (!quake.veq<10>) -> 164 -// ──────────────────────────────────────────────── -// %4 = constant 10 : i64 -def ForwardConstantVeqSizePattern : Pat< - (quake_VeqSizeOp:$res $veq), (createConstantOp $res, $veq), - [(KnownSizePred $veq)]>; - -def SizeIsPresentPred : Constraint(" - " $0[0].getDefiningOp())">>; - -def createAllocaOp : NativeCodeCall< - "quake::createConstantAlloca($_builder, $_loc, $0, $1)">; - -// %2 = constant 10 : i32 -// %3 = quake.alloca !quake.veq[%2 : i32] -// ─────────────────────────────────────────── -// %3 = quake.alloca !quake.veq<10> -def FuseConstantToAllocaPattern : Pat< - (quake_AllocaOp:$alloca $optSize), (createAllocaOp $alloca, $optSize), - [(SizeIsPresentPred $optSize)]>; - -def createExtractRefOp : NativeCodeCall< - "$_builder.create($_loc, $0," - " cast($1[0].getDefiningOp()).getValue()." - " cast().getInt())">; - -// %2 = constant 10 : i32 -// %3 = quake.extract_ref %1[%2] : (!quake.veq, i32) -> !quake.ref -// ─────────────────────────────────────────── -// %3 = quake.extract_ref %1[10] : (!quake.veq) -> !quake.ref -def FuseConstantToExtractRefPattern : Pat< - (quake_ExtractRefOp $veq, $index, $rawIndex), - (createExtractRefOp $veq, $index), - [(SizeIsPresentPred $index)]>; - -#endif diff --git a/lib/Optimizer/Dialect/Quake/CMakeLists.txt b/lib/Optimizer/Dialect/Quake/CMakeLists.txt index 87733716c4..bc55e40d52 100644 --- a/lib/Optimizer/Dialect/Quake/CMakeLists.txt +++ b/lib/Optimizer/Dialect/Quake/CMakeLists.txt @@ -16,7 +16,6 @@ add_cudaq_dialect_library(QuakeDialect QuakeDialectIncGen QuakeOpsIncGen QuakeTypesIncGen - CanonicalIncGen LINK_LIBS CCDialect diff --git a/lib/Optimizer/Dialect/Quake/CanonicalPatterns.inc b/lib/Optimizer/Dialect/Quake/CanonicalPatterns.inc new file mode 100644 index 0000000000..de1eada9be --- /dev/null +++ b/lib/Optimizer/Dialect/Quake/CanonicalPatterns.inc @@ -0,0 +1,489 @@ +/****************************************************************-*- C++ -*-**** + * Copyright (c) 2022 - 2024 NVIDIA Corporation & Affiliates. * + * All rights reserved. * + * * + * This source code and the accompanying materials are made available under * + * the terms of the Apache License 2.0 which accompanies this distribution. * + ******************************************************************************/ + +// These canonicalization patterns are used by the canonicalize pass and not +// shared for other uses. Generally speaking, these patterns should be trivial +// peephole optimizations that reduce the size and complexity of the input IR. + +// This file must be included after a `using namespace mlir;` as it uses bare +// identifiers from that namespace. + +namespace { + +// %4 = quake.veq_size %3 : (!quake.veq<10>) -> 164 +// ──────────────────────────────────────────────── +// %4 = constant 10 : i64 +struct ForwardConstantVeqSizePattern + : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::VeqSizeOp veqSize, + PatternRewriter &rewriter) const override { + auto veqTy = dyn_cast(veqSize.getVeq().getType()); + if (!veqTy) + return failure(); + if (!veqTy.hasSpecifiedSize()) + return failure(); + auto resTy = veqSize.getType(); + rewriter.replaceOpWithNewOp(veqSize, veqTy.getSize(), + resTy); + return success(); + } +}; + +// %2 = constant 10 : i32 +// %3 = quake.alloca !quake.veq[%2 : i32] +// ───────────────────────────────────────── +// %3 = quake.alloca !quake.veq<10> +struct FuseConstantToAllocaPattern : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::AllocaOp alloc, + PatternRewriter &rewriter) const override { + auto size = alloc.getSize(); + if (!size) + return failure(); + auto intCon = cudaq::opt::factory::getIntIfConstant(size); + if (!intCon) + return failure(); + auto veqTy = dyn_cast(alloc.getType()); + if (!veqTy) + return failure(); + if (veqTy.hasSpecifiedSize()) + return failure(); + auto loc = alloc.getLoc(); + auto resTy = alloc.getType(); + auto newAlloc = rewriter.create( + loc, static_cast(*intCon)); + rewriter.replaceOpWithNewOp(alloc, resTy, newAlloc); + return success(); + } +}; + +// %2 = constant 10 : i32 +// %3 = quake.extract_ref %1[%2] : (!quake.veq, i32) -> !quake.ref +// ────────────────────────────────────────────────────────────────── +// %3 = quake.extract_ref %1[10] : (!quake.veq) -> !quake.ref +struct FuseConstantToExtractRefPattern + : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::ExtractRefOp extract, + PatternRewriter &rewriter) const override { + auto index = extract.getIndex(); + if (!index) + return failure(); + auto intCon = cudaq::opt::factory::getIntIfConstant(index); + if (!intCon) + return failure(); + rewriter.replaceOpWithNewOp( + extract, extract.getVeq(), static_cast(*intCon)); + return success(); + } +}; + +// %4 = quake.concat %2, %3 : (!quake.ref, !quake.ref) -> !quake.veq<2> +// %7 = quake.extract_ref %4[0] : (!quake.veq<2>) -> !quake.ref +// ─────────────────────────────────────────── +// replace all use with %2 +struct ForwardConcatExtractPattern + : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::ExtractRefOp extract, + PatternRewriter &rewriter) const override { + auto veq = extract.getVeq(); + auto concatOp = veq.getDefiningOp(); + if (concatOp && extract.hasConstantIndex()) { + // Don't run this canonicalization if any of the operands + // to concat are of type veq. + auto concatQubits = concatOp.getQbits(); + for (auto qOp : concatQubits) + if (isa(qOp.getType())) + return failure(); + + // concat only has ref type operands. + auto index = extract.getConstantIndex(); + if (index < concatQubits.size()) { + auto qOpValue = concatQubits[index]; + if (isa(qOpValue.getType())) { + rewriter.replaceOp(extract, {qOpValue}); + return success(); + } + } + } + return failure(); + } +}; + +// %2 = quake.concat %1 : (!quake.ref) -> !quake.veq<1> +// %3 = quake.extract_ref %2[0] : (!quake.veq<1>) -> !quake.ref +// quake.* %3 ... +// ─────────────────────────────────────────── +// quake.* %1 ... +struct ForwardConcatExtractSingleton + : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::ExtractRefOp extract, + PatternRewriter &rewriter) const override { + if (auto concat = extract.getVeq().getDefiningOp()) + if (concat.getType().getSize() == 1 && extract.hasConstantIndex() && + extract.getConstantIndex() == 0) { + assert(concat.getQbits().size() == 1 && concat.getQbits()[0]); + extract.getResult().replaceUsesWithIf( + concat.getQbits()[0], [&](OpOperand &use) { + if (Operation *user = use.getOwner()) + return isQuakeOperation(user); + return false; + }); + return success(); + } + return failure(); + } +}; + +// %7 = quake.concat %4 : (!quake.veq<2>) -> !quake.veq<2> +// ─────────────────────────────────────────── +// removed +struct ConcatNoOpPattern : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::ConcatOp concat, + PatternRewriter &rewriter) const override { + // Remove concat veq -> veq + // or + // concat ref -> ref + auto qubitsToConcat = concat.getQbits(); + if (qubitsToConcat.size() > 1) + return failure(); + + // We only want to handle veq -> veq here. + if (isa(qubitsToConcat.front().getType())) { + return failure(); + } + + // Do not handle anything where we don't know the sizes. + auto retTy = concat.getResult().getType(); + if (auto veqTy = dyn_cast(retTy)) + if (!veqTy.hasSpecifiedSize()) + // This could be a folded quake.relax_size op. + return failure(); + + rewriter.replaceOp(concat, qubitsToConcat); + return success(); + } +}; + +// %8 = quake.concat %4, %5, %6 : (!quake.ref, !quake.veq<4>, +// !quake.veq<2>) -> !quake.veq +// ─────────────────────────────────────────────────────────── +// %.8 = quake.concat %4, %5, %6 : (!quake.ref, !quake.veq<4>, +// !quake.veq<2>) -> !quake.veq<7> +// %8 = quake.relax_size %.8 : (!quake.veq<7>) -> !quake.veq +struct ConcatSizePattern : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::ConcatOp concat, + PatternRewriter &rewriter) const override { + if (concat.getType().hasSpecifiedSize()) + return failure(); + + // Walk the arguments and sum them, if possible. + std::size_t sum = 0; + for (auto opnd : concat.getQbits()) { + if (auto veqTy = dyn_cast(opnd.getType())) { + if (!veqTy.hasSpecifiedSize()) + return failure(); + sum += veqTy.getSize(); + continue; + } + assert(isa(opnd.getType())); + sum++; + } + + // Leans into the relax_size canonicalization pattern. + auto *ctx = rewriter.getContext(); + auto loc = concat.getLoc(); + auto newTy = quake::VeqType::get(ctx, sum); + Value newOp = + rewriter.create(loc, newTy, concat.getQbits()); + auto noSizeTy = quake::VeqType::getUnsized(ctx); + rewriter.replaceOpWithNewOp(concat, noSizeTy, newOp); + return success(); + } +}; + +// %7 = quake.make_struq %5, %6 : (!quake.veq, !quake.veq) -> +// !quake.struq, !quake.veq> +// %8 = quake.get_member %7[1] : (!quake.struq, +// !quake.veq>) -> !quake.veq +// ─────────────────────────────────────────────────────────── +// replace uses of %8 with %6 +struct BypassMakeStruq : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::GetMemberOp getMem, + PatternRewriter &rewriter) const override { + auto makeStruq = getMem.getStruq().getDefiningOp(); + if (!makeStruq) + return failure(); + auto toStrTy = cast(getMem.getStruq().getType()); + std::uint32_t idx = getMem.getIndex(); + Value from = makeStruq.getOperand(idx); + auto toTy = toStrTy.getMembers()[idx]; + if (from.getType() != toTy) + rewriter.replaceOpWithNewOp(getMem, toTy, from); + else + rewriter.replaceOp(getMem, from); + return success(); + } +}; + +// %22 = quake.init_state %1, %2 : (!quake.veq, T) -> !quake.veq +// ──────────────────────────────────────────────────────────────────── +// %.22 = quake.init_state %1, %2 : (!quake.veq, T) -> !quake.veq +// %22 = quake.relax_size %.22 : (!quake.veq) -> !quake.veq +struct ForwardAllocaTypePattern + : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::InitializeStateOp initState, + PatternRewriter &rewriter) const override { + if (auto isTy = dyn_cast(initState.getType())) + if (!isTy.hasSpecifiedSize()) { + auto targ = initState.getTargets(); + if (auto targTy = dyn_cast(targ.getType())) + if (targTy.hasSpecifiedSize()) { + auto newInit = rewriter.create( + initState.getLoc(), targTy, targ, initState.getState()); + rewriter.replaceOpWithNewOp(initState, isTy, + newInit); + return success(); + } + } + + // Remove any intervening cast to !cc.ptr> ops. + if (auto stateCast = + initState.getState().getDefiningOp()) + if (auto ptrTy = dyn_cast(stateCast.getType())) { + auto eleTy = ptrTy.getElementType(); + if (auto arrTy = dyn_cast(eleTy)) + if (arrTy.isUnknownSize()) { + rewriter.replaceOpWithNewOp( + initState, initState.getTargets().getType(), + initState.getTargets(), stateCast.getValue()); + return success(); + } + } + return failure(); + } +}; + +// %3 = quake.subveq %0, 4, 10 : (!quake.veq<12>, i64, i64) -> !quake.veq +// ────────────────────────────────────────────────────────────────────────── +// %.3 = quake.subveq %0, 4, 10 : (!quake.veq<12>, i64, i64) -> !quake.veq<7> +// %3 = quake.relax_size %.3 : (!quake.veq<7>) -> !quake.veq +struct FixUnspecifiedSubveqPattern : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::SubVeqOp subveq, + PatternRewriter &rewriter) const override { + auto veqTy = dyn_cast(subveq.getType()); + if (veqTy && veqTy.hasSpecifiedSize()) + return failure(); + if (!(subveq.hasConstantLowerBound() && subveq.hasConstantUpperBound())) + return failure(); + auto *ctx = rewriter.getContext(); + std::size_t size = + subveq.getConstantUpperBound() - subveq.getConstantLowerBound() + 1u; + auto szVecTy = quake::VeqType::get(ctx, size); + auto loc = subveq.getLoc(); + auto subv = rewriter.create( + loc, szVecTy, subveq.getVeq(), subveq.getLower(), subveq.getUpper(), + subveq.getRawLower(), subveq.getRawUpper()); + rewriter.replaceOpWithNewOp(subveq, veqTy, subv); + return success(); + } +}; + +// %1 = constant 4 : i64 +// %2 = constant 10 : i64 +// %3 = quake.subveq %0, %1, %2 : (!quake.veq<12>, i64, i64) -> !quake.veq +// ────────────────────────────────────────────────────────────────────────── +// %3 = quake.subveq %0, 4, 10 : (!quake.veq<12>, i64, i64) -> !quake.veq<7> +struct FuseConstantToSubveqPattern : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::SubVeqOp subveq, + PatternRewriter &rewriter) const override { + if (subveq.hasConstantLowerBound() && subveq.hasConstantUpperBound()) + return failure(); + bool regen = false; + std::int64_t lo = subveq.getConstantLowerBound(); + Value loVal = subveq.getLower(); + if (!subveq.hasConstantLowerBound()) + if (auto olo = cudaq::opt::factory::getIntIfConstant(subveq.getLower())) { + regen = true; + loVal = nullptr; + lo = *olo; + } + + std::int64_t hi = subveq.getConstantUpperBound(); + Value hiVal = subveq.getUpper(); + if (!subveq.hasConstantUpperBound()) + if (auto ohi = cudaq::opt::factory::getIntIfConstant(subveq.getUpper())) { + regen = true; + hiVal = nullptr; + hi = *ohi; + } + + if (!regen) + return failure(); + rewriter.replaceOpWithNewOp( + subveq, subveq.getType(), subveq.getVeq(), loVal, hiVal, lo, hi); + return success(); + } +}; + +// Replace subveq operations that extract the entire original register with the +// original register. +struct RemoveSubVeqNoOpPattern : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::SubVeqOp subVeqOp, + PatternRewriter &rewriter) const override { + auto origVeq = subVeqOp.getVeq(); + // The original veq size must be known + auto veqType = dyn_cast(origVeq.getType()); + if (!veqType.hasSpecifiedSize()) + return failure(); + if (!(subVeqOp.hasConstantLowerBound() && subVeqOp.hasConstantUpperBound())) + return failure(); + + // If the subveq is the whole register, than the start value must be 0. + if (subVeqOp.getConstantLowerBound() != 0) + return failure(); + + // If the sizes are equal, then replace + if (veqType.getSize() != subVeqOp.getConstantUpperBound() + 1) + return failure(); + + // this subveq is the whole original register, hence a no-op + rewriter.replaceOp(subVeqOp, origVeq); + return success(); + } +}; + +// %11 = quake.init_state %_, %_ : (!quake.veq<2>, T1) -> !quake.veq +// %12 = quake.veq_size %11 : (!quake.veq) -> i64 +// ──────────────────────────────────────────────────────────────────── +// %11 = quake.init_state %_, %_ : (!quake.veq<2>, T1) -> !quake.veq +// %12 = constant 2 : i64 +struct FoldInitStateSizePattern : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::VeqSizeOp veqSize, + PatternRewriter &rewriter) const override { + Value veq = veqSize.getVeq(); + if (auto initState = veq.getDefiningOp()) + if (auto veqTy = + dyn_cast(initState.getTargets().getType())) + if (veqTy.hasSpecifiedSize()) { + std::size_t numQubits = veqTy.getSize(); + rewriter.replaceOpWithNewOp(veqSize, numQubits, + veqSize.getType()); + return success(); + } + return failure(); + } +}; + +// If there is no operation that modifies the wire after it gets unwrapped and +// before it is wrapped, then the wrap operation is a nop and can be +// eliminated. +struct KillDeadWrapPattern : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::WrapOp wrap, + PatternRewriter &rewriter) const override { + if (auto unwrap = wrap.getWireValue().getDefiningOp()) + rewriter.eraseOp(wrap); + return success(); + } +}; + +template +struct MergeRotationPattern : public OpRewritePattern { + using Base = OpRewritePattern; + using Base::Base; + + LogicalResult matchAndRewrite(OP rotate, + PatternRewriter &rewriter) const override { + auto wireTy = quake::WireType::get(rewriter.getContext()); + if (rotate.getTarget(0).getType() != wireTy || + !rotate.getControls().empty()) + return failure(); + assert(!rotate.getNegatedQubitControls()); + auto input = rotate.getTarget(0).template getDefiningOp(); + if (!input || !input.getControls().empty()) + return failure(); + assert(!input.getNegatedQubitControls()); + + // At this point, we have + // %input = quake.rotate %angle1, %wire + // %rotate = quake.rotate %angle2, %input + // Replace those ops with + // %new = quake.rotate (%angle1 + %angle2), %wire + auto loc = rotate.getLoc(); + auto angle1 = input.getParameter(0); + auto angle2 = rotate.getParameter(0); + if (angle1.getType() != angle2.getType()) + return failure(); + auto adjAttr = rotate.getIsAdjAttr(); + auto newAngle = [&]() -> Value { + if (input.isAdj() == rotate.isAdj()) + return rewriter.create(loc, angle1, angle2); + // One is adjoint, so it should be subtracted from the other. + if (input.isAdj()) + return rewriter.create(loc, angle2, angle1); + adjAttr = input.getIsAdjAttr(); + return rewriter.create(loc, angle1, angle2); + }(); + rewriter.replaceOpWithNewOp(rotate, rotate.getResultTypes(), adjAttr, + ValueRange{newAngle}, ValueRange{}, + ValueRange{input.getTarget(0)}, + rotate.getNegatedQubitControlsAttr()); + return success(); + } +}; + +// Forward the argument to a relax_size to the users for all users that are +// quake operations. All quake ops that take a sized veq argument are +// polymorphic on all veq types. If the op is not a quake op, then maintain +// strong typing. +struct ForwardRelaxedSizePattern : public OpRewritePattern { + using OpRewritePattern::OpRewritePattern; + + LogicalResult matchAndRewrite(quake::RelaxSizeOp relax, + PatternRewriter &rewriter) const override { + auto inpVec = relax.getInputVec(); + bool replaced = false; + rewriter.replaceOpWithIf(relax, inpVec, [&](OpOperand &use) { + bool res = false; + if (Operation *user = use.getOwner()) + res = isQuakeOperation(user) && !isa(user); + replaced = replaced || res; + return res; + }); + // return success if and only if at least one use was replaced. + return success(replaced); + }; +}; + +} // namespace diff --git a/lib/Optimizer/Dialect/Quake/QuakeOps.cpp b/lib/Optimizer/Dialect/Quake/QuakeOps.cpp index 1f2c3bd06d..d2da75d99b 100644 --- a/lib/Optimizer/Dialect/Quake/QuakeOps.cpp +++ b/lib/Optimizer/Dialect/Quake/QuakeOps.cpp @@ -23,9 +23,7 @@ using namespace mlir; -namespace { -#include "cudaq/Optimizer/Dialect/Quake/Canonical.inc" -} // namespace +#include "CanonicalPatterns.inc" static LogicalResult verifyWireResultsAreLinear(Operation *op) { for (Value v : op->getOpResults()) @@ -311,73 +309,6 @@ LogicalResult quake::BorrowWireOp::verify() { // Concat //===----------------------------------------------------------------------===// -namespace { -// %7 = quake.concat %4 : (!quake.veq<2>) -> !quake.veq<2> -// ─────────────────────────────────────────── -// removed -struct ConcatNoOpPattern : public OpRewritePattern { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::ConcatOp concat, - PatternRewriter &rewriter) const override { - // Remove concat veq -> veq - // or - // concat ref -> ref - auto qubitsToConcat = concat.getQbits(); - if (qubitsToConcat.size() > 1) - return failure(); - - // We only want to handle veq -> veq here. - if (isa(qubitsToConcat.front().getType())) { - return failure(); - } - - // Do not handle anything where we don't know the sizes. - auto retTy = concat.getResult().getType(); - if (auto veqTy = dyn_cast(retTy)) - if (!veqTy.hasSpecifiedSize()) - // This could be a folded quake.relax_size op. - return failure(); - - rewriter.replaceOp(concat, qubitsToConcat); - return success(); - } -}; - -struct ConcatSizePattern : public OpRewritePattern { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::ConcatOp concat, - PatternRewriter &rewriter) const override { - if (concat.getType().hasSpecifiedSize()) - return failure(); - - // Walk the arguments and sum them, if possible. - std::size_t sum = 0; - for (auto opnd : concat.getQbits()) { - if (auto veqTy = dyn_cast(opnd.getType())) { - if (!veqTy.hasSpecifiedSize()) - return failure(); - sum += veqTy.getSize(); - continue; - } - assert(isa(opnd.getType())); - sum++; - } - - // Leans into the relax_size canonicalization pattern. - auto *ctx = rewriter.getContext(); - auto loc = concat.getLoc(); - auto newTy = quake::VeqType::get(ctx, sum); - Value newOp = - rewriter.create(loc, newTy, concat.getQbits()); - auto noSizeTy = quake::VeqType::getUnsized(ctx); - rewriter.replaceOpWithNewOp(concat, noSizeTy, newOp); - return success(); - }; -}; -} // namespace - void quake::ConcatOp::getCanonicalizationPatterns(RewritePatternSet &patterns, MLIRContext *context) { patterns.add(context); @@ -418,69 +349,6 @@ void printRawIndex(OpAsmPrinter &printer, OP refOp, Value index, printer << rawIndex.getValue(); } -namespace { -// %4 = quake.concat %2, %3 : (!quake.ref, !quake.ref) -> !quake.veq<2> -// %7 = quake.extract_ref %4[0] : (!quake.veq<2>) -> !quake.ref -// ─────────────────────────────────────────── -// replace all use with %2 -struct ForwardConcatExtractPattern - : public OpRewritePattern { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::ExtractRefOp extract, - PatternRewriter &rewriter) const override { - auto veq = extract.getVeq(); - auto concatOp = veq.getDefiningOp(); - if (concatOp && extract.hasConstantIndex()) { - // Don't run this canonicalization if any of the operands - // to concat are of type veq. - auto concatQubits = concatOp.getQbits(); - for (auto qOp : concatQubits) - if (isa(qOp.getType())) - return failure(); - - // concat only has ref type operands. - auto index = extract.getConstantIndex(); - if (index < concatQubits.size()) { - auto qOpValue = concatQubits[index]; - if (isa(qOpValue.getType())) { - rewriter.replaceOp(extract, {qOpValue}); - return success(); - } - } - } - return failure(); - } -}; - -// %2 = quake.concat %1 : (!quake.ref) -> !quake.veq<1> -// %3 = quake.extract_ref %2[0] : (!quake.veq<1>) -> !quake.ref -// quake.* %3 ... -// ─────────────────────────────────────────── -// quake.* %1 ... -struct ForwardConcatExtractSingleton - : public OpRewritePattern { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::ExtractRefOp extract, - PatternRewriter &rewriter) const override { - if (auto concat = extract.getVeq().getDefiningOp()) - if (concat.getType().getSize() == 1 && extract.hasConstantIndex() && - extract.getConstantIndex() == 0) { - assert(concat.getQbits().size() == 1 && concat.getQbits()[0]); - extract.getResult().replaceUsesWithIf( - concat.getQbits()[0], [&](OpOperand &use) { - if (Operation *user = use.getOwner()) - return isQuakeOperation(user); - return false; - }); - return success(); - } - return failure(); - } -}; -} // namespace - void quake::ExtractRefOp::getCanonicalizationPatterns( RewritePatternSet &patterns, MLIRContext *context) { patterns.add { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::GetMemberOp getMem, - PatternRewriter &rewriter) const override { - if (auto makeStruq = - getMem.getStruq().getDefiningOp()) { - auto toStrTy = cast(getMem.getStruq().getType()); - std::uint32_t idx = getMem.getIndex(); - Value from = makeStruq.getOperand(idx); - auto toTy = toStrTy.getMembers()[idx]; - if (from.getType() != toTy) { - rewriter.replaceOpWithNewOp(getMem, toTy, from); - } else { - rewriter.replaceOp(getMem, from); - } - return success(); - } - return failure(); - } -}; -} // namespace - void quake::GetMemberOp::getCanonicalizationPatterns( RewritePatternSet &patterns, MLIRContext *context) { patterns.add(context); @@ -575,48 +419,6 @@ LogicalResult quake::InitializeStateOp::verify() { return success(); } -namespace { -// %22 = quake.init_state %1, %2 : (!quake.veq, T) -> !quake.veq -// ──────────────────────────────────────────────────────────────────── -// %22' = quake.init_state %1, %2 : (!quake.veq, T) -> !quake.veq -// %22 = quake.relax_size %22' : (!quake.veq) -> !quake.veq -struct ForwardAllocaTypePattern - : public OpRewritePattern { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::InitializeStateOp initState, - PatternRewriter &rewriter) const override { - if (auto isTy = dyn_cast(initState.getType())) - if (!isTy.hasSpecifiedSize()) { - auto targ = initState.getTargets(); - if (auto targTy = dyn_cast(targ.getType())) - if (targTy.hasSpecifiedSize()) { - auto newInit = rewriter.create( - initState.getLoc(), targTy, targ, initState.getState()); - rewriter.replaceOpWithNewOp(initState, isTy, - newInit); - return success(); - } - } - - // Remove any intervening cast to !cc.ptr> ops. - if (auto stateCast = - initState.getState().getDefiningOp()) - if (auto ptrTy = dyn_cast(stateCast.getType())) { - auto eleTy = ptrTy.getElementType(); - if (auto arrTy = dyn_cast(eleTy)) - if (arrTy.isUnknownSize()) { - rewriter.replaceOpWithNewOp( - initState, initState.getTargets().getType(), - initState.getTargets(), stateCast.getValue()); - return success(); - } - } - return failure(); - } -}; -} // namespace - void quake::InitializeStateOp::getCanonicalizationPatterns( RewritePatternSet &patterns, MLIRContext *context) { patterns.add(context); @@ -652,30 +454,6 @@ LogicalResult quake::RelaxSizeOp::verify() { return success(); } -namespace { -// Forward the argument to a relax_size to the users for all users that are -// quake operations. All quake ops that take a sized veq argument are -// polymorphic on all veq types. If the op is not a quake op, then maintain -// strong typing. -struct ForwardRelaxedSizePattern : public RewritePattern { - ForwardRelaxedSizePattern(MLIRContext *context) - : RewritePattern("quake.relax_size", 1, context, {}) {} - - LogicalResult matchAndRewrite(Operation *op, - PatternRewriter &rewriter) const override { - auto relax = cast(op); - auto inpVec = relax.getInputVec(); - Value result = relax.getResult(); - result.replaceUsesWithIf(inpVec, [&](OpOperand &use) { - if (Operation *user = use.getOwner()) - return isQuakeOperation(user) && !isa(user); - return false; - }); - return success(); - }; -}; -} // namespace - void quake::RelaxSizeOp::getCanonicalizationPatterns( RewritePatternSet &patterns, MLIRContext *context) { patterns.add(context); @@ -709,103 +487,6 @@ LogicalResult quake::SubVeqOp::verify() { return success(); } -namespace { -// %3 = quake.subveq %0, 4, 10 : (!quake.veq<12>, i64, i64) -> !quake.veq -// ───────────────────────────────────────────────────────────────────────────── -// %new3 = quake.subveq %0, 4, 10 : (!quake.veq<12>, i64, i64) -> !quake.veq<7> -// %3 = quake.relax_size %new3 : (!quake.veq<7>) -> !quake.veq -struct FixUnspecifiedSubveqPattern : public OpRewritePattern { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::SubVeqOp subveq, - PatternRewriter &rewriter) const override { - auto veqTy = dyn_cast(subveq.getType()); - if (veqTy && veqTy.hasSpecifiedSize()) - return failure(); - if (!(subveq.hasConstantLowerBound() && subveq.hasConstantUpperBound())) - return failure(); - auto *ctx = rewriter.getContext(); - std::size_t size = - subveq.getConstantUpperBound() - subveq.getConstantLowerBound() + 1u; - auto szVecTy = quake::VeqType::get(ctx, size); - auto loc = subveq.getLoc(); - auto subv = rewriter.create( - loc, szVecTy, subveq.getVeq(), subveq.getLower(), subveq.getUpper(), - subveq.getRawLower(), subveq.getRawUpper()); - rewriter.replaceOpWithNewOp(subveq, veqTy, subv); - return success(); - } -}; - -// %1 = constant 4 : i64 -// %2 = constant 10 : i64 -// %3 = quake.subveq %0, %1, %2 : (!quake.veq<12>, i64, i64) -> !quake.veq -// ───────────────────────────────────────────────────────────────────────────── -// %3 = quake.subveq %0, 4, 10 : (!quake.veq<12>, i64, i64) -> !quake.veq<7> -struct FuseConstantToSubveqPattern : public OpRewritePattern { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::SubVeqOp subveq, - PatternRewriter &rewriter) const override { - if (subveq.hasConstantLowerBound() && subveq.hasConstantUpperBound()) - return failure(); - bool regen = false; - std::int64_t lo = subveq.getConstantLowerBound(); - Value loVal = subveq.getLower(); - if (!subveq.hasConstantLowerBound()) - if (auto olo = cudaq::opt::factory::getIntIfConstant(subveq.getLower())) { - regen = true; - loVal = nullptr; - lo = *olo; - } - - std::int64_t hi = subveq.getConstantUpperBound(); - Value hiVal = subveq.getUpper(); - if (!subveq.hasConstantUpperBound()) - if (auto ohi = cudaq::opt::factory::getIntIfConstant(subveq.getUpper())) { - regen = true; - hiVal = nullptr; - hi = *ohi; - } - - if (!regen) - return failure(); - rewriter.replaceOpWithNewOp( - subveq, subveq.getType(), subveq.getVeq(), loVal, hiVal, lo, hi); - return success(); - } -}; - -// Replace subveq operations that extract the entire original register with the -// original register. -struct RemoveSubVeqNoOpPattern : public OpRewritePattern { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::SubVeqOp subVeqOp, - PatternRewriter &rewriter) const override { - auto origVeq = subVeqOp.getVeq(); - // The original veq size must be known - auto veqType = dyn_cast(origVeq.getType()); - if (!veqType.hasSpecifiedSize()) - return failure(); - if (!(subVeqOp.hasConstantLowerBound() && subVeqOp.hasConstantUpperBound())) - return failure(); - - // If the subveq is the whole register, than the start value must be 0. - if (subVeqOp.getConstantLowerBound() != 0) - return failure(); - - // If the sizes are equal, then replace - if (veqType.getSize() != subVeqOp.getConstantUpperBound() + 1) - return failure(); - - // this subveq is the whole original register, hence a no-op - rewriter.replaceOp(subVeqOp, origVeq); - return success(); - } -}; -} // namespace - void quake::SubVeqOp::getCanonicalizationPatterns(RewritePatternSet &patterns, MLIRContext *context) { patterns.add { - using OpRewritePattern::OpRewritePattern; - - // %11 = quake.init_state %_, %_ : (!quake.veq<2>, T1) -> !quake.veq - // %12 = quake.veq_size %11 : (!quake.veq) -> i64 - // ──────────────────────────────────────────────────────────────────── - // %11 = quake.init_state %_, %_ : (!quake.veq<2>, T1) -> !quake.veq - // %12 = constant 2 : i64 - LogicalResult matchAndRewrite(quake::VeqSizeOp veqSize, - PatternRewriter &rewriter) const override { - Value veq = veqSize.getVeq(); - if (auto initState = veq.getDefiningOp()) - if (auto veqTy = - dyn_cast(initState.getTargets().getType())) - if (veqTy.hasSpecifiedSize()) { - std::size_t numQubits = veqTy.getSize(); - rewriter.replaceOpWithNewOp(veqSize, numQubits, - veqSize.getType()); - return success(); - } - return failure(); - } -}; -} // namespace - void quake::VeqSizeOp::getCanonicalizationPatterns(RewritePatternSet &patterns, MLIRContext *context) { patterns.add( @@ -852,22 +507,6 @@ void quake::VeqSizeOp::getCanonicalizationPatterns(RewritePatternSet &patterns, // WrapOp //===----------------------------------------------------------------------===// -namespace { -// If there is no operation that modifies the wire after it gets unwrapped and -// before it is wrapped, then the wrap operation is a nop and can be -// eliminated. -struct KillDeadWrapPattern : public OpRewritePattern { - using OpRewritePattern::OpRewritePattern; - - LogicalResult matchAndRewrite(quake::WrapOp wrap, - PatternRewriter &rewriter) const override { - if (auto unwrap = wrap.getWireValue().getDefiningOp()) - rewriter.eraseOp(wrap); - return success(); - } -}; -} // namespace - void quake::WrapOp::getCanonicalizationPatterns(RewritePatternSet &patterns, MLIRContext *context) { patterns.add(context); @@ -1040,53 +679,6 @@ void quake::RxOp::getOperatorMatrix(Matrix &matrix) { -1i * std::sin(theta / 2.), std::cos(theta / 2.)}); } -namespace { -template -struct MergeRotationPattern : public OpRewritePattern { - using Base = OpRewritePattern; - using Base::Base; - - LogicalResult matchAndRewrite(OP rotate, - PatternRewriter &rewriter) const override { - auto wireTy = quake::WireType::get(rewriter.getContext()); - if (rotate.getTarget(0).getType() != wireTy || - !rotate.getControls().empty()) - return failure(); - assert(!rotate.getNegatedQubitControls()); - auto input = rotate.getTarget(0).template getDefiningOp(); - if (!input || !input.getControls().empty()) - return failure(); - assert(!input.getNegatedQubitControls()); - - // At this point, we have - // %input = quake.rotate %angle1, %wire - // %rotate = quake.rotate %angle2, %input - // Replace those ops with - // %new = quake.rotate (%angle1 + %angle2), %wire - auto loc = rotate.getLoc(); - auto angle1 = input.getParameter(0); - auto angle2 = rotate.getParameter(0); - if (angle1.getType() != angle2.getType()) - return failure(); - auto adjAttr = rotate.getIsAdjAttr(); - auto newAngle = [&]() -> Value { - if (input.isAdj() == rotate.isAdj()) - return rewriter.create(loc, angle1, angle2); - // One is adjoint, so it should be subtracted from the other. - if (input.isAdj()) - return rewriter.create(loc, angle2, angle1); - adjAttr = input.getIsAdjAttr(); - return rewriter.create(loc, angle1, angle2); - }(); - rewriter.replaceOpWithNewOp(rotate, rotate.getResultTypes(), adjAttr, - ValueRange{newAngle}, ValueRange{}, - ValueRange{input.getTarget(0)}, - rotate.getNegatedQubitControlsAttr()); - return success(); - } -}; -} // namespace - void quake::RxOp::getCanonicalizationPatterns(RewritePatternSet &patterns, MLIRContext *context) { patterns.add>(context);