From 4f37e20f285e4df00808639bc4f192a04bff8f04 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 8 Oct 2024 19:04:26 +0000
Subject: [PATCH 01/66] initial commit

---
 .../lightning_gpu/StateVectorCudaManaged.hpp  | 413 ++++++++++++------
 .../lightning_gpu/bindings/LGPUBindings.hpp   |  51 ++-
 .../lightning_gpu/_state_vector.py            |  27 +-
 .../lightning_gpu/lightning_gpu.py            |   1 +
 4 files changed, 354 insertions(+), 138 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index 9e68592df..67427f969 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -352,17 +352,17 @@ class StateVectorCudaManaged
         } else if (opName == "GlobalPhase") {
             globalPhaseStateVector(adjoint, params[0]);
         } else if (native_gates_.find(opName) != native_gates_.end()) {
-            applyParametricPauliGate({opName}, ctrls, tgts, params.front(),
-                                     adjoint);
+            applyParametricPauliGate_({opName}, ctrls, tgts, params.front(),
+                                      adjoint);
         } else if (opName == "Rot" || opName == "CRot") {
             if (adjoint) {
                 auto rot_matrix =
                     cuGates::getRot<CFP_t>(params[2], params[1], params[0]);
-                applyDeviceMatrixGate(rot_matrix.data(), ctrls, tgts, true);
+                applyDeviceMatrixGate_(rot_matrix.data(), ctrls, tgts, true);
             } else {
                 auto rot_matrix =
                     cuGates::getRot<CFP_t>(params[0], params[1], params[2]);
-                applyDeviceMatrixGate(rot_matrix.data(), ctrls, tgts, false);
+                applyDeviceMatrixGate_(rot_matrix.data(), ctrls, tgts, false);
             }
         } else if (opName == "Matrix") {
             DataBuffer<CFP_t, int> d_matrix{
@@ -375,8 +375,8 @@ class StateVectorCudaManaged
                                                        ctrls.rend()};
             const std::vector<std::size_t> tgts_local{tgts.rbegin(),
                                                       tgts.rend()};
-            applyDeviceMatrixGate(d_matrix.getData(), ctrls_local, tgts_local,
-                                  adjoint);
+            applyDeviceMatrixGate_(d_matrix.getData(), ctrls_local, tgts_local,
+                                   adjoint);
         } else if (par_gates_.find(opName) != par_gates_.end()) {
             par_gates_.at(opName)(wires, adjoint, params);
         } else { // No offloadable function call; defer to matrix passing
@@ -395,7 +395,7 @@ class StateVectorCudaManaged
             } else if (!gate_cache_.gateExists(opName, par[0])) {
                 gate_cache_.add_gate(opName, par[0], gate_matrix);
             }
-            applyDeviceMatrixGate(
+            applyDeviceMatrixGate_(
                 gate_cache_.get_gate_device_ptr(opName, par[0]), ctrls_local,
                 tgts_local, adjoint);
         }
@@ -405,27 +405,141 @@ class StateVectorCudaManaged
      * @brief Apply a single gate to the state vector.
      *
      * @param opName Name of gate to apply.
-     * @param controlled_wires Control wires.
-     * @param controlled_values Control values (false or true).
-     * @param wires Wires to apply gate to.
-     * @param inverse Indicates whether to use adjoint of gate.
+     * @param ctrls Control wires.
+     * @param ctrls_values Control values (false or true).
+     * @param tgts Wires to apply gate to.
+     * @param adjoint Indicates whether to use adjoint of gate.
      * @param params Optional parameter list for parametric gates.
-     * @param params Optional std gate matrix if opName doesn't exist.
+     * @param gate_matrix Optional gate matrix vector on the host if opName
+     * doesn't exist.
      */
-    template <template <typename...> class complex_t>
-    void
-    applyOperation(const std::string &opName,
-                   const std::vector<std::size_t> &controlled_wires,
-                   const std::vector<bool> &controlled_values,
-                   const std::vector<std::size_t> &wires, bool inverse = false,
-                   const std::vector<Precision> &params = {0.0},
-                   const std::vector<complex_t<Precision>> &gate_matrix = {}) {
-        PL_ABORT_IF_NOT(controlled_wires.empty(),
-                        "Controlled kernels not implemented.");
-        PL_ABORT_IF_NOT(controlled_wires.size() == controlled_values.size(),
-                        "`controlled_wires` must have the same size as "
-                        "`controlled_values`.");
-        applyOperation(opName, wires, inverse, params, gate_matrix);
+    void applyOperation(const std::string &opName,
+                        const std::vector<std::size_t> &ctrls,
+                        const std::vector<bool> &ctrls_values,
+                        const std::vector<std::size_t> &tgts,
+                        bool adjoint = false,
+                        const std::vector<Precision> &params = {0.0},
+                        const std::vector<ComplexT> &gate_matrix = {}) {
+        PL_ABORT_IF(ctrls.size() != ctrls_values.size(),
+                    "`ctrls` and `ctrls_values` must have the same size.");
+        std::vector<int> ctrlsInt(ctrls.size());
+        std::vector<int> tgtsInt(tgts.size());
+        std::vector<int> ctrls_valuesInt(ctrls.size());
+
+        std::transform(
+            ctrls.begin(), ctrls.end(), ctrlsInt.begin(), [&](std::size_t x) {
+                return static_cast<int>(BaseType::getNumQubits() - 1 - x);
+            });
+        std::transform(
+            tgts.begin(), tgts.end(), tgtsInt.begin(), [&](std::size_t x) {
+                return static_cast<int>(BaseType::getNumQubits() - 1 - x);
+            });
+
+        std::transform(ctrls_values.begin(), ctrls_values.end(),
+                       ctrls_valuesInt.begin(),
+                       [&](bool x) { return static_cast<int>(x); });
+        if (opName == "MultiRZ") {
+            const std::vector<std::string> names(tgts.size(), {"RZ"});
+            applyParametricPauliGeneralGate_(names, ctrlsInt, ctrls_valuesInt,
+                                             tgtsInt, params.front(), adjoint);
+        } else if (native_gates_.find(opName) != native_gates_.end()) {
+            applyParametricPauliGeneralGate_({opName}, ctrlsInt,
+                                             ctrls_valuesInt, tgtsInt,
+                                             params.front(), adjoint);
+        } else if (opName == "Rot" || opName == "CRot") {
+            if (adjoint) {
+                auto rot_matrix =
+                    cuGates::getRot<CFP_t>(params[2], params[1], params[0]);
+                applyDeviceGeneralGate_(rot_matrix.data(), ctrlsInt, tgtsInt,
+                                        ctrls_valuesInt, true);
+            } else {
+                auto rot_matrix =
+                    cuGates::getRot<CFP_t>(params[0], params[1], params[2]);
+                applyDeviceGeneralGate_(rot_matrix.data(), ctrlsInt, tgtsInt,
+                                        ctrls_valuesInt, false);
+            }
+        } else if (opName == "Matrix") {
+            DataBuffer<CFP_t, int> d_matrix{
+                gate_matrix.size(), BaseType::getDataBuffer().getDevTag(),
+                true};
+            d_matrix.CopyHostDataToGpu(gate_matrix.data(), d_matrix.getLength(),
+                                       false);
+
+            applyDeviceGeneralGate_(d_matrix.getData(), ctrlsInt, tgtsInt,
+                                    ctrls_valuesInt, adjoint);
+        } else if (par_gates_.find(opName) != par_gates_.end()) {
+            auto &gateMap =
+                cuGates::DynamicGateDataAccess<PrecisionT>::getInstance();
+            auto &&matrix_cu = gateMap.getGateData(opName, params);
+
+            gate_cache_.add_gate(opName, params[0], matrix_cu);
+
+            applyDeviceGeneralGate_(
+                gate_cache_.get_gate_device_ptr(opName, params[0]), ctrlsInt,
+                tgtsInt, ctrls_valuesInt, adjoint);
+        } else { // No offloadable function call; defer to matrix passing
+            auto &&par =
+                (params.empty()) ? std::vector<Precision>{0.0} : params;
+            if (!gate_cache_.gateExists(opName, par[0]) &&
+                gate_matrix.empty()) {
+                std::string message = "Currently unsupported gate: " + opName;
+                throw LightningException(message);
+            } else if (!gate_cache_.gateExists(opName, par[0])) {
+                std::vector<CFP_t> matrix_cu(gate_matrix.size());
+                std::transform(
+                    gate_matrix.begin(), gate_matrix.end(), matrix_cu.begin(),
+                    [](const std::complex<Precision> &x) {
+                        return cuUtil::complexToCu<std::complex<Precision>>(x);
+                    });
+
+                gate_cache_.add_gate(opName, par[0], matrix_cu);
+            }
+            applyDeviceGeneralGate_(
+                gate_cache_.get_gate_device_ptr(opName, par[0]), ctrlsInt,
+                tgtsInt, ctrls_valuesInt, adjoint);
+        }
+    }
+
+    /**
+     * @brief Apply a single gate to the state vector.
+     *
+     * @param gate_matrix Gate matrix data (in row-major format).
+     * @param ctrls Control wires.
+     * @param ctrls_values Control values (false or true).
+     * @param tgts Target wires to apply gate to.
+     * @param inverse Indicates whether to use adjoint of gate.
+     */
+    void applyControlledMatrix(const ComplexT *gate_matrix,
+                               const std::size_t matrix_size,
+                               const std::vector<std::size_t> &ctrls,
+                               const std::vector<bool> &ctrls_values,
+                               const std::vector<std::size_t> &tgts,
+                               bool inverse = false) {
+        PL_ABORT_IF(ctrls.size() != ctrls_values.size(),
+                    "`ctrls` and `ctrls_values` must have the same size.");
+        DataBuffer<CFP_t, int> d_matrix{
+            matrix_size, BaseType::getDataBuffer().getDevTag(), true};
+        d_matrix.CopyHostDataToGpu(gate_matrix, matrix_size, false);
+        // ensure wire indexing correctly preserved for tensor-observables
+        std::vector<int> ctrlsInt(ctrls.size());
+        std::vector<int> tgtsInt(tgts.size());
+        std::vector<int> ctrls_valuesInt(ctrls.size());
+
+        std::transform(
+            ctrls.begin(), ctrls.end(), ctrlsInt.begin(), [&](std::size_t x) {
+                return static_cast<int>(BaseType::getNumQubits() - 1 - x);
+            });
+        std::transform(
+            tgts.begin(), tgts.end(), tgtsInt.begin(), [&](std::size_t x) {
+                return static_cast<int>(BaseType::getNumQubits() - 1 - x);
+            });
+
+        std::transform(ctrls_values.begin(), ctrls_values.end(),
+                       ctrls_valuesInt.begin(),
+                       [&](bool x) { return static_cast<int>(x); });
+
+        applyDeviceGeneralGate_(d_matrix.getData(), ctrlsInt, tgtsInt,
+                                ctrls_valuesInt, inverse);
     }
 
     /**
@@ -500,65 +614,65 @@ class StateVectorCudaManaged
                             bool adjoint) {
         static const std::string name{"PauliX"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applyPauliY(const std::vector<std::size_t> &wires,
                             bool adjoint) {
         static const std::string name{"PauliY"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applyPauliZ(const std::vector<std::size_t> &wires,
                             bool adjoint) {
         static const std::string name{"PauliZ"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applyHadamard(const std::vector<std::size_t> &wires,
                               bool adjoint) {
         static const std::string name{"Hadamard"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applyS(const std::vector<std::size_t> &wires, bool adjoint) {
         static const std::string name{"S"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applyT(const std::vector<std::size_t> &wires, bool adjoint) {
         static const std::string name{"T"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applyRX(const std::vector<std::size_t> &wires, bool adjoint,
                         Precision param) {
         static const std::vector<std::string> name{{"RX"}};
-        applyParametricPauliGate(name, {wires.begin(), wires.end() - 1},
-                                 {wires.back()}, param, adjoint);
+        applyParametricPauliGate_(name, {wires.begin(), wires.end() - 1},
+                                  {wires.back()}, param, adjoint);
     }
     inline void applyRY(const std::vector<std::size_t> &wires, bool adjoint,
                         Precision param) {
         static const std::vector<std::string> name{{"RY"}};
-        applyParametricPauliGate(name, {wires.begin(), wires.end() - 1},
-                                 {wires.back()}, param, adjoint);
+        applyParametricPauliGate_(name, {wires.begin(), wires.end() - 1},
+                                  {wires.back()}, param, adjoint);
     }
     inline void applyRZ(const std::vector<std::size_t> &wires, bool adjoint,
                         Precision param) {
         static const std::vector<std::string> name{{"RZ"}};
-        applyParametricPauliGate(name, {wires.begin(), wires.end() - 1},
-                                 {wires.back()}, param, adjoint);
+        applyParametricPauliGate_(name, {wires.begin(), wires.end() - 1},
+                                  {wires.back()}, param, adjoint);
     }
     inline void applyRot(const std::vector<std::size_t> &wires, bool adjoint,
                          Precision param0, Precision param1, Precision param2) {
@@ -578,53 +692,53 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(gate_key,
                                  cuGates::getPhaseShift<CFP_t>(param));
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
 
     /* two-qubit gates */
     inline void applyCNOT(const std::vector<std::size_t> &wires, bool adjoint) {
         static const std::string name{"CNOT"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applyCY(const std::vector<std::size_t> &wires, bool adjoint) {
         static const std::string name{"CY"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applyCZ(const std::vector<std::size_t> &wires, bool adjoint) {
         static const std::string name{"CZ"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applySWAP(const std::vector<std::size_t> &wires, bool adjoint) {
         static const std::string name{"SWAP"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param), {},
+                               wires, adjoint);
     }
     inline void applyIsingXX(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
         static const std::vector<std::string> names(wires.size(), {"RX"});
-        applyParametricPauliGate(names, {}, wires, param, adjoint);
+        applyParametricPauliGate_(names, {}, wires, param, adjoint);
     }
     inline void applyIsingYY(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
         static const std::vector<std::string> names(wires.size(), {"RY"});
-        applyParametricPauliGate(names, {}, wires, param, adjoint);
+        applyParametricPauliGate_(names, {}, wires, param, adjoint);
     }
     inline void applyIsingZZ(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
         static const std::vector<std::string> names(wires.size(), {"RZ"});
-        applyParametricPauliGate(names, {}, wires, param, adjoint);
+        applyParametricPauliGate_(names, {}, wires, param, adjoint);
     }
     inline void applyIsingXY(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
@@ -633,8 +747,8 @@ class StateVectorCudaManaged
         if (!gate_cache_.gateExists(gate_key)) {
             gate_cache_.add_gate(gate_key, cuGates::getIsingXY<CFP_t>(param));
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
     }
     inline void applyCRot(const std::vector<std::size_t> &wires, bool adjoint,
                           const std::vector<Precision> &params) {
@@ -672,8 +786,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(gate_key,
                                  cuGates::getSingleExcitation<CFP_t>(param));
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
     }
     inline void
     applySingleExcitationMinus(const std::vector<std::size_t> &wires,
@@ -684,8 +798,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(
                 gate_key, cuGates::getSingleExcitationMinus<CFP_t>(param));
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
     }
     inline void applySingleExcitationPlus(const std::vector<std::size_t> &wires,
                                           bool adjoint, Precision param) {
@@ -695,8 +809,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(
                 gate_key, cuGates::getSingleExcitationPlus<CFP_t>(param));
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
     }
 
     /* three-qubit gates */
@@ -704,42 +818,42 @@ class StateVectorCudaManaged
                              bool adjoint) {
         static const std::string name{"Toffoli"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.begin(), wires.end() - 1}, {wires.back()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.begin(), wires.end() - 1}, {wires.back()},
+                               adjoint);
     }
     inline void applyCSWAP(const std::vector<std::size_t> &wires,
                            bool adjoint) {
         static const std::string name{"SWAP"};
         static const Precision param = 0.0;
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                              {wires.front()}, {wires.begin() + 1, wires.end()},
-                              adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                               {wires.front()},
+                               {wires.begin() + 1, wires.end()}, adjoint);
     }
 
     /* four-qubit gates */
     inline void applyDoubleExcitation(const std::vector<std::size_t> &wires,
                                       bool adjoint, Precision param) {
         auto &&mat = cuGates::getDoubleExcitation<CFP_t>(param);
-        applyDeviceMatrixGate(mat.data(), {}, wires, adjoint);
+        applyDeviceMatrixGate_(mat.data(), {}, wires, adjoint);
     }
     inline void
     applyDoubleExcitationMinus(const std::vector<std::size_t> &wires,
                                bool adjoint, Precision param) {
         auto &&mat = cuGates::getDoubleExcitationMinus<CFP_t>(param);
-        applyDeviceMatrixGate(mat.data(), {}, wires, adjoint);
+        applyDeviceMatrixGate_(mat.data(), {}, wires, adjoint);
     }
     inline void applyDoubleExcitationPlus(const std::vector<std::size_t> &wires,
                                           bool adjoint, Precision param) {
         auto &&mat = cuGates::getDoubleExcitationPlus<CFP_t>(param);
-        applyDeviceMatrixGate(mat.data(), {}, wires, adjoint);
+        applyDeviceMatrixGate_(mat.data(), {}, wires, adjoint);
     }
 
     /* Multi-qubit gates */
     inline void applyMultiRZ(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
         const std::vector<std::string> names(wires.size(), {"RZ"});
-        applyParametricPauliGate(names, {}, wires, param, adjoint);
+        applyParametricPauliGate_(names, {}, wires, param, adjoint);
     }
 
     /* Gate generators */
@@ -805,8 +919,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(gate_key,
                                  cuGates::getGeneratorIsingXX<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return -static_cast<PrecisionT>(0.5);
     }
     inline PrecisionT
@@ -818,8 +932,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(gate_key,
                                  cuGates::getGeneratorIsingYY<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return -static_cast<PrecisionT>(0.5);
     }
     inline PrecisionT
@@ -831,8 +945,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(gate_key,
                                  cuGates::getGeneratorIsingZZ<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return -static_cast<PrecisionT>(0.5);
     }
 
@@ -845,8 +959,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(gate_key,
                                  cuGates::getGeneratorIsingXY<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return static_cast<PrecisionT>(0.5);
     }
 
@@ -933,8 +1047,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(
                 gate_key, cuGates::getGeneratorSingleExcitation<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return -static_cast<PrecisionT>(0.5);
     }
     inline PrecisionT
@@ -947,8 +1061,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(
                 gate_key, cuGates::getGeneratorSingleExcitationMinus<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return -static_cast<PrecisionT>(0.5);
     }
     inline PrecisionT
@@ -961,8 +1075,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(
                 gate_key, cuGates::getGeneratorSingleExcitationPlus<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return -static_cast<PrecisionT>(0.5);
     }
 
@@ -976,8 +1090,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(
                 gate_key, cuGates::getGeneratorDoubleExcitation<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return -static_cast<PrecisionT>(0.5);
     }
     inline PrecisionT
@@ -990,8 +1104,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(
                 gate_key, cuGates::getGeneratorDoubleExcitationMinus<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return -static_cast<PrecisionT>(0.5);
     }
     inline PrecisionT
@@ -1004,8 +1118,8 @@ class StateVectorCudaManaged
             gate_cache_.add_gate(
                 gate_key, cuGates::getGeneratorDoubleExcitationPlus<CFP_t>());
         }
-        applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(gate_key), {},
-                              wires, adjoint);
+        applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(gate_key), {},
+                               wires, adjoint);
         return -static_cast<PrecisionT>(0.5);
     }
 
@@ -1014,8 +1128,8 @@ class StateVectorCudaManaged
         static const std::string name{"PauliZ"};
         static const Precision param = 0.0;
         for (const auto &w : wires) {
-            applyDeviceMatrixGate(gate_cache_.get_gate_device_ptr(name, param),
-                                  {}, {w}, adjoint);
+            applyDeviceMatrixGate_(gate_cache_.get_gate_device_ptr(name, param),
+                                   {}, {w}, adjoint);
         }
         return -static_cast<PrecisionT>(0.5);
     }
@@ -1414,12 +1528,10 @@ class StateVectorCudaManaged
      * @param tgts target wires.
      * @param use_adjoint Take adjoint of operation.
      */
-    void applyParametricPauliGate(const std::vector<std::string> &pauli_words,
-                                  std::vector<std::size_t> ctrls,
-                                  std::vector<std::size_t> tgts,
-                                  Precision param, bool use_adjoint = false) {
-        int nIndexBits = BaseType::getNumQubits();
-
+    void applyParametricPauliGate_(const std::vector<std::string> &pauli_words,
+                                   std::vector<std::size_t> ctrls,
+                                   std::vector<std::size_t> tgts,
+                                   Precision param, bool use_adjoint = false) {
         std::vector<int> ctrlsInt(ctrls.size());
         std::vector<int> tgtsInt(tgts.size());
 
@@ -1433,6 +1545,29 @@ class StateVectorCudaManaged
                 return static_cast<int>(BaseType::getNumQubits() - 1 - x);
             });
 
+        const std::vector<int> ctrls_valuesInt(ctrls.size(), 1);
+
+        applyParametricPauliGeneralGate_(pauli_words, ctrlsInt, ctrls_valuesInt,
+                                         tgtsInt, param, use_adjoint);
+    }
+
+    /**
+     * @brief Apply a parametric Pauli gate using custateVec calls.
+     *
+     * @param pauli_words List of Pauli words representing operation.
+     * @param ctrls Control wires
+     * @param ctrls_values Control values
+     * @param tgts target wires.
+     * @param param Rotation angle.
+     * @param use_adjoint Take adjoint of operation.
+     */
+    void applyParametricPauliGeneralGate_(
+        const std::vector<std::string> &pauli_words,
+        const std::vector<int> &ctrlsInt,
+        const std::vector<int> &ctrls_valuesInt, const std::vector<int> tgtsInt,
+        Precision param, bool use_adjoint = false) {
+        int nIndexBits = BaseType::getNumQubits();
+
         cudaDataType_t data_type;
 
         if constexpr (std::is_same_v<CFP_t, cuDoubleComplex> ||
@@ -1457,10 +1592,10 @@ class StateVectorCudaManaged
             /* double */ local_angle,
             /* const custatevecPauli_t* */ pauli_enums.data(),
             /* const int32_t* */ tgtsInt.data(),
-            /* const uint32_t */ tgts.size(),
+            /* const uint32_t */ tgtsInt.size(),
             /* const int32_t* */ ctrlsInt.data(),
-            /* const int32_t* */ nullptr,
-            /* const uint32_t */ ctrls.size()));
+            /* const int32_t* */ ctrls_valuesInt.data(),
+            /* const uint32_t */ ctrlsInt.size()));
         PL_CUDA_IS_SUCCESS(cudaStreamSynchronize(
             BaseType::getDataBuffer().getDevTag().getStreamID()));
     }
@@ -1477,14 +1612,10 @@ class StateVectorCudaManaged
      * @param tgts Target qubits.
      * @param use_adjoint Use adjoint of given gate.
      */
-    void applyDeviceMatrixGate(const CFP_t *matrix,
-                               const std::vector<std::size_t> &ctrls,
-                               const std::vector<std::size_t> &tgts,
-                               bool use_adjoint = false) {
-        void *extraWorkspace = nullptr;
-        std::size_t extraWorkspaceSizeInBytes = 0;
-        int nIndexBits = BaseType::getNumQubits();
-
+    void applyDeviceMatrixGate_(const CFP_t *matrix,
+                                const std::vector<std::size_t> &ctrls,
+                                const std::vector<std::size_t> &tgts,
+                                bool use_adjoint = false) {
         std::vector<int> ctrlsInt(ctrls.size());
         std::vector<int> tgtsInt(tgts.size());
 
@@ -1497,6 +1628,34 @@ class StateVectorCudaManaged
                 return static_cast<int>(BaseType::getNumQubits() - 1 - x);
             });
 
+        const std::vector<int> ctrls_values(ctrls.size(), 1);
+
+        applyDeviceGeneralGate_(matrix, ctrlsInt, tgtsInt, ctrls_values,
+                                use_adjoint);
+    }
+
+    /**
+     * @brief Apply a given host or device-stored array representing the gate
+     * `matrix` to the state vector at qubit indices given by `tgts` and
+     * control-lines given by `ctrls`. The adjoint can be taken by setting
+     * `use_adjoint` to true.
+     *
+     * @param matrix Device data array in row-major order representing
+     * a given gate.
+     * @param ctrls Control line qubits.
+     * @param tgts Target qubits.
+     * @param ctrls_values Control values.
+     * @param use_adjoint Use adjoint of given gate. Defaults to false.
+     */
+    void applyDeviceGeneralGate_(const CFP_t *matrix,
+                                 const std::vector<int> &ctrls,
+                                 const std::vector<int> &tgts,
+                                 const std::vector<int> &ctrls_values,
+                                 bool use_adjoint = false) {
+        void *extraWorkspace = nullptr;
+        std::size_t extraWorkspaceSizeInBytes = 0;
+        int nIndexBits = BaseType::getNumQubits();
+
         cudaDataType_t data_type;
         custatevecComputeType_t compute_type;
 
@@ -1544,10 +1703,10 @@ class StateVectorCudaManaged
             /* cudaDataType_t */ data_type,
             /* custatevecMatrixLayout_t */ CUSTATEVEC_MATRIX_LAYOUT_ROW,
             /* const int32_t */ use_adjoint,
-            /* const int32_t* */ tgtsInt.data(),
+            /* const int32_t* */ tgts.data(),
             /* const uint32_t */ tgts.size(),
-            /* const int32_t* */ ctrlsInt.data(),
-            /* const int32_t* */ nullptr,
+            /* const int32_t* */ ctrls.data(),
+            /* const int32_t* */ ctrls_values.data(),
             /* const uint32_t */ ctrls.size(),
             /* custatevecComputeType_t */ compute_type,
             /* void* */ extraWorkspace,
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp
index c361bd6ed..a8e393883 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp
@@ -51,9 +51,55 @@ using StateVectorBackends =
                               StateVectorCudaManaged<double>, void>;
 
 /**
- * @brief Get a gate kernel map for a statevector.
+ * @brief Register controlled matrix kernel.
  */
+template <class StateVectorT>
+void applyControlledMatrix(
+    StateVectorT &st,
+    const py::array_t<std::complex<typename StateVectorT::PrecisionT>,
+                      py::array::c_style | py::array::forcecast> &matrix,
+    const std::vector<std::size_t> &controlled_wires,
+    const std::vector<bool> &controlled_values,
+    const std::vector<std::size_t> &wires, bool inverse = false) {
+    using ComplexT = typename StateVectorT::ComplexT;
+    st.applyControlledMatrix(
+        static_cast<const ComplexT *>(matrix.request().ptr),
+        static_cast<const std::size_t>(matrix.request().size), controlled_wires,
+        controlled_values, wires, inverse);
+}
+
+template <class StateVectorT, class PyClass>
+void registerControlledGate(PyClass &pyclass) {
+    using PrecisionT =
+        typename StateVectorT::PrecisionT; // Statevector's precision
+    using ParamT = PrecisionT;             // Parameter's data precision
 
+    using Pennylane::Gates::ControlledGateOperation;
+    using Pennylane::Util::for_each_enum;
+    namespace Constant = Pennylane::Gates::Constant;
+
+    for_each_enum<ControlledGateOperation>(
+        [&pyclass](ControlledGateOperation gate_op) {
+            using Pennylane::Util::lookup;
+            const auto gate_name =
+                std::string(lookup(Constant::controlled_gate_names, gate_op));
+            const std::string doc = "Apply the " + gate_name + " gate.";
+            auto func = [gate_name = gate_name](
+                            StateVectorT &sv,
+                            const std::vector<std::size_t> &controlled_wires,
+                            const std::vector<bool> &controlled_values,
+                            const std::vector<std::size_t> &wires, bool inverse,
+                            const std::vector<ParamT> &params) {
+                sv.applyOperation(gate_name, controlled_wires,
+                                  controlled_values, wires, inverse, params);
+            };
+            pyclass.def(gate_name.c_str(), func, doc.c_str());
+        });
+}
+
+/**
+ * @brief Get a gate kernel map for a statevector.
+ */
 template <class StateVectorT, class PyClass>
 void registerBackendClassSpecificBindings(PyClass &pyclass) {
     using PrecisionT =
@@ -65,6 +111,7 @@ void registerBackendClassSpecificBindings(PyClass &pyclass) {
                                  py::array::c_style | py::array::forcecast>;
 
     registerGatesForStateVector<StateVectorT>(pyclass);
+    registerControlledGate<StateVectorT>(pyclass);
 
     pyclass
         .def(py::init<std::size_t>())              // qubits, device
@@ -96,6 +143,8 @@ void registerBackendClassSpecificBindings(PyClass &pyclass) {
             },
             "Set State Vector on GPU with values for the state vector and "
             "wires on the host memory.")
+        .def("applyControlledMatrix", &applyControlledMatrix<StateVectorT>,
+             "Apply controlled operation")
         .def(
             "DeviceToDevice",
             [](StateVectorT &sv, const StateVectorT &other, bool async) {
diff --git a/pennylane_lightning/lightning_gpu/_state_vector.py b/pennylane_lightning/lightning_gpu/_state_vector.py
index a00044356..e8d180196 100644
--- a/pennylane_lightning/lightning_gpu/_state_vector.py
+++ b/pennylane_lightning/lightning_gpu/_state_vector.py
@@ -237,15 +237,24 @@ def _apply_lightning_controlled(self, operation):
         """
         state = self.state_vector
 
+        basename = operation.base.name
+        method = getattr(state, f"{basename}", None)
         control_wires = list(operation.control_wires)
         control_values = operation.control_values
-        name = operation.name
-        # Apply GlobalPhase
-        inv = False
-        param = operation.parameters[0]
-        wires = self.wires.indices(operation.wires)
-        matrix = global_phase_diagonal(param, self.wires, control_wires, control_values)
-        state.apply(name, wires, inv, [[param]], matrix)
+        target_wires = list(operation.target_wires)
+        if method is not None:  # apply n-controlled specialized gate
+            inv = False
+            param = operation.parameters
+            method(control_wires, control_values, target_wires, inv, param)
+        else:  # apply gate as an n-controlled matrix
+            method = getattr(state, "applyControlledMatrix")
+            method(
+                qml.matrix(operation.base),
+                control_wires,
+                control_values,
+                target_wires,
+                False,
+            )
 
     def _apply_lightning_midmeasure(self):
         """Execute a MidMeasureMP operation and return the sample in mid_measurements.
@@ -292,9 +301,7 @@ def _apply_lightning(
             if method is not None:  # apply specialized gate
                 param = operation.parameters
                 method(wires, invert_param, param)
-            elif isinstance(operation, qml.ops.Controlled) and isinstance(
-                operation.base, qml.GlobalPhase
-            ):  # apply n-controlled gate
+            elif isinstance(operation, qml.ops.Controlled):  # apply n-controlled gate
                 # LGPU do not support the controlled gates except for GlobalPhase
                 self._apply_lightning_controlled(operation)
             else:  # apply gate as a matrix
diff --git a/pennylane_lightning/lightning_gpu/lightning_gpu.py b/pennylane_lightning/lightning_gpu/lightning_gpu.py
index 2b295c499..2a8f22517 100644
--- a/pennylane_lightning/lightning_gpu/lightning_gpu.py
+++ b/pennylane_lightning/lightning_gpu/lightning_gpu.py
@@ -87,6 +87,7 @@
         "MultiRZ",
         "GlobalPhase",
         "C(GlobalPhase)",
+        "C(MultiRZ)",
         "Hadamard",
         "S",
         "Adjoint(S)",

From dad7de87ea394bce4b984843efd1bd3eb9e1b2d1 Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Tue, 8 Oct 2024 19:05:04 +0000
Subject: [PATCH 02/66] Auto update version from '0.39.0-dev40' to
 '0.39.0-dev41'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index c6e0b5050..ab5a1f0f5 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev40"
+__version__ = "0.39.0-dev41"

From dd0307a7cb35003c12b2e6cd9dedad115acb121f Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Wed, 16 Oct 2024 17:48:33 +0000
Subject: [PATCH 03/66] Auto update version from '0.39.0-dev45' to
 '0.39.0-dev46'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index 93371827c..37f7ba1bb 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev45"
+__version__ = "0.39.0-dev46"

From d81874ba7c3ed9094f7a959831e38cc8e406faab Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Wed, 16 Oct 2024 21:32:18 +0000
Subject: [PATCH 04/66] fix applyControlledMatrix

---
 .../src/simulators/lightning_gpu/StateVectorCudaManaged.hpp   | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index 866dc3188..8827a87b8 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -526,11 +526,11 @@ class StateVectorCudaManaged
         std::vector<int> ctrls_valuesInt(ctrls.size());
 
         std::transform(
-            ctrls.begin(), ctrls.end(), ctrlsInt.begin(), [&](std::size_t x) {
+            ctrls.rbegin(), ctrls.rend(), ctrlsInt.begin(), [&](std::size_t x) {
                 return static_cast<int>(BaseType::getNumQubits() - 1 - x);
             });
         std::transform(
-            tgts.begin(), tgts.end(), tgtsInt.begin(), [&](std::size_t x) {
+            tgts.rbegin(), tgts.rend(), tgtsInt.begin(), [&](std::size_t x) {
                 return static_cast<int>(BaseType::getNumQubits() - 1 - x);
             });
 

From ed2df4579d4832c0856fd503879b1609e58943d5 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Wed, 16 Oct 2024 23:12:49 +0000
Subject: [PATCH 05/66] update GlobalPhase and CGlobalPhase gate support

---
 .../lightning_gpu/StateVectorCudaManaged.hpp  |  72 ++--------
 .../src/simulators/lightning_gpu/initSV.cu    | 129 ------------------
 2 files changed, 12 insertions(+), 189 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index 8827a87b8..2d8a0f39d 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -67,25 +67,6 @@ extern void setBasisState_CUDA(cuDoubleComplex *sv, cuDoubleComplex &value,
                                const std::size_t index, bool async,
                                cudaStream_t stream_id);
 
-extern void globalPhaseStateVector_CUDA(cuComplex *sv, std::size_t num_sv,
-                                        cuComplex phase,
-                                        std::size_t thread_per_block,
-                                        cudaStream_t stream_id);
-extern void globalPhaseStateVector_CUDA(cuDoubleComplex *sv, std::size_t num_sv,
-                                        cuDoubleComplex phase,
-                                        std::size_t thread_per_block,
-                                        cudaStream_t stream_id);
-
-extern void cGlobalPhaseStateVector_CUDA(cuComplex *sv, std::size_t num_sv,
-                                         bool adjoint, cuComplex *phase,
-                                         std::size_t thread_per_block,
-                                         cudaStream_t stream_id);
-extern void cGlobalPhaseStateVector_CUDA(cuDoubleComplex *sv,
-                                         std::size_t num_sv, bool adjoint,
-                                         cuDoubleComplex *phase,
-                                         std::size_t thread_per_block,
-                                         cudaStream_t stream_id);
-
 /**
  * @brief Managed memory CUDA state-vector class using custateVec backed
  * gate-calls.
@@ -257,44 +238,6 @@ class StateVectorCudaManaged
                                     use_async);
     }
 
-    /**
-     * @brief Multiplies the state-vector by a global phase.
-     *
-     * @param adjoint Indicates whether to use adjoint of gate.
-     * @param param Complex phase generator.
-     */
-    template <std::size_t thread_per_block = 256>
-    void globalPhaseStateVector(const bool adjoint, const Precision param) {
-        auto stream_id = BaseType::getDataBuffer().getDevTag().getStreamID();
-        std::complex<Precision> phase =
-            std::exp(std::complex<Precision>{0, (adjoint) ? param : -param});
-        auto cuPhase = complexToCu(phase);
-        globalPhaseStateVector_CUDA(BaseType::getData(), BaseType::getLength(),
-                                    cuPhase, thread_per_block, stream_id);
-    }
-
-    /**
-     * @brief Multiplies the state-vector by a controlled global phase.
-     *
-     * @param phase Controlled complex phase vector.
-     */
-    template <std::size_t thread_per_block = 256>
-    void cGlobalPhaseStateVector(const bool adjoint,
-                                 const std::vector<CFP_t> &phase,
-                                 const bool async = false) {
-        PL_ABORT_IF_NOT(BaseType::getLength() == phase.size(),
-                        "The state-vector data must have the same size as the "
-                        "controlled-phase data.")
-        auto device_id = BaseType::getDataBuffer().getDevTag().getDeviceID();
-        auto stream_id = BaseType::getDataBuffer().getDevTag().getStreamID();
-        DataBuffer<CFP_t, int> d_phase{phase.size(), device_id, stream_id,
-                                       true};
-        d_phase.CopyHostDataToGpu(phase.data(), d_phase.getLength(), async);
-        cGlobalPhaseStateVector_CUDA(BaseType::getData(), BaseType::getLength(),
-                                     adjoint, d_phase.getData(),
-                                     thread_per_block, stream_id);
-    }
-
     /**
      * @brief Apply a single gate to the state-vector. Offloads to custatevec
      * specific API calls if available. If unable, attempts to use prior cached
@@ -347,10 +290,14 @@ class StateVectorCudaManaged
                                             wires.end()};
         if (opName == "Identity") {
             return;
-        } else if (opName == "C(GlobalPhase)") {
-            cGlobalPhaseStateVector(adjoint, gate_matrix);
         } else if (opName == "GlobalPhase") {
-            globalPhaseStateVector(adjoint, params[0]);
+            //TODO: Check the performance GlobalPhase implementation via scaling the statevector
+            const std::vector<std::string> names(BaseType::getNumQubits(), "I");
+            std::vector<std::size_t> tgts_all(BaseType::getNumQubits());
+            std::iota(tgts_all.begin(), tgts_all.end(), 0);
+            std::reverse(tgts_all.begin(), tgts_all.end());
+            applyParametricPauliGate_(names, {}, tgts_all, 2 * params[0],
+                                      adjoint);
         } else if (native_gates_.find(opName) != native_gates_.end()) {
             applyParametricPauliGate_({opName}, ctrls, tgts, params.front(),
                                       adjoint);
@@ -442,6 +389,11 @@ class StateVectorCudaManaged
             const std::vector<std::string> names(tgts.size(), {"RZ"});
             applyParametricPauliGeneralGate_(names, ctrlsInt, ctrls_valuesInt,
                                              tgtsInt, params.front(), adjoint);
+        } else if (opName == "GlobalPhase") {
+            const std::vector<std::string> names(tgts.size(), "I");
+            applyParametricPauliGeneralGate_(names, ctrlsInt, ctrls_valuesInt,
+                                             tgtsInt, 2 * params[0], adjoint);
+
         } else if (native_gates_.find(opName) != native_gates_.end()) {
             applyParametricPauliGeneralGate_({opName}, ctrlsInt,
                                              ctrls_valuesInt, tgtsInt,
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/initSV.cu b/pennylane_lightning/core/src/simulators/lightning_gpu/initSV.cu
index 8a62e89e8..381a34915 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/initSV.cu
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/initSV.cu
@@ -106,96 +106,6 @@ void setStateVector_CUDA_call(GPUDataT *sv, index_type &num_indices,
     PL_CUDA_IS_SUCCESS(cudaGetLastError());
 }
 
-/**
- * @brief The CUDA kernel that multiplies the state vector data on GPU device
- * by a global phase.
- *
- * @param sv Complex data pointer of state vector on device.
- * @param num_sv Number of state vector elements.
- * @param phase Complex data pointer of input values (on device).
- */
-template <class GPUDataT, class index_type>
-__global__ void globalPhaseStateVectorkernel(GPUDataT *sv, index_type num_sv,
-                                             GPUDataT phase) {
-    const unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
-    if (i < num_sv) {
-        sv[i] = Cmul(sv[i], phase);
-    }
-}
-
-/**
- * @brief The CUDA kernel call wrapper.
- *
- * @param sv Complex data pointer of state vector on device.
- * @param num_sv Number of state vector elements.
- * @param phase Constant complex phase.
- * @param thread_per_block Number of threads set per block.
- * @param stream_id Stream id of CUDA calls
- */
-template <class GPUDataT, class index_type>
-void globalPhaseStateVector_CUDA_call(GPUDataT *sv, index_type num_sv,
-                                      GPUDataT phase,
-                                      std::size_t thread_per_block,
-                                      cudaStream_t stream_id) {
-    auto dv = std::div(static_cast<long>(num_sv), thread_per_block);
-    std::size_t num_blocks = dv.quot + (dv.rem == 0 ? 0 : 1);
-    const std::size_t block_per_grid = (num_blocks == 0 ? 1 : num_blocks);
-    dim3 blockSize(thread_per_block, 1, 1);
-    dim3 gridSize(block_per_grid, 1);
-
-    globalPhaseStateVectorkernel<GPUDataT, index_type>
-        <<<gridSize, blockSize, 0, stream_id>>>(sv, num_sv, phase);
-    PL_CUDA_IS_SUCCESS(cudaGetLastError());
-}
-
-/**
- * @brief The CUDA kernel that multiplies the state vector data on GPU device
- * by a controlled global phase.
- *
- * @param sv Complex data pointer of state vector on device.
- * @param num_sv Number of state vector elements.
- * @param phase Complex data pointer of controlled global phase values (on
- * device).
- */
-template <class GPUDataT, class index_type, bool adjoint = false>
-__global__ void cGlobalPhaseStateVectorkernel(GPUDataT *sv, index_type num_sv,
-                                              GPUDataT *phase) {
-    const unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
-    if (i < num_sv) {
-        if constexpr (adjoint) {
-            sv[i] = Cmul(sv[i], Conj(phase[i]));
-        } else {
-            sv[i] = Cmul(sv[i], phase[i]);
-        }
-    }
-}
-
-/**
- * @brief The CUDA kernel call wrapper.
- *
- * @param sv Complex data pointer of state vector on device.
- * @param num_sv Number of state vector elements.
- * @param phase Complex data pointer of controlled global phase values (on
- * device).
- * @param thread_per_block Number of threads set per block.
- * @param stream_id Stream id of CUDA calls
- */
-template <class GPUDataT, class index_type, bool adjoint = false>
-void cGlobalPhaseStateVector_CUDA_call(GPUDataT *sv, index_type num_sv,
-                                       GPUDataT *phase,
-                                       std::size_t thread_per_block,
-                                       cudaStream_t stream_id) {
-    auto dv = std::div(static_cast<long>(num_sv), thread_per_block);
-    std::size_t num_blocks = dv.quot + (dv.rem == 0 ? 0 : 1);
-    const std::size_t block_per_grid = (num_blocks == 0 ? 1 : num_blocks);
-    dim3 blockSize(thread_per_block, 1, 1);
-    dim3 gridSize(block_per_grid, 1);
-
-    cGlobalPhaseStateVectorkernel<GPUDataT, index_type, adjoint>
-        <<<gridSize, blockSize, 0, stream_id>>>(sv, num_sv, phase);
-    PL_CUDA_IS_SUCCESS(cudaGetLastError());
-}
-
 /**
  * @brief CUDA runtime API call wrapper.
  *
@@ -243,43 +153,4 @@ void setBasisState_CUDA(cuDoubleComplex *sv, cuDoubleComplex &value,
     setBasisState_CUDA_call(sv, value, index, async, stream_id);
 }
 
-void globalPhaseStateVector_CUDA(cuComplex *sv, std::size_t num_sv,
-                                 cuComplex phase, std::size_t thread_per_block,
-                                 cudaStream_t stream_id) {
-    globalPhaseStateVector_CUDA_call(sv, num_sv, phase, thread_per_block,
-                                     stream_id);
-}
-void globalPhaseStateVector_CUDA(cuDoubleComplex *sv, std::size_t num_sv,
-                                 cuDoubleComplex phase,
-                                 std::size_t thread_per_block,
-                                 cudaStream_t stream_id) {
-    globalPhaseStateVector_CUDA_call(sv, num_sv, phase, thread_per_block,
-                                     stream_id);
-}
-
-void cGlobalPhaseStateVector_CUDA(cuComplex *sv, std::size_t num_sv,
-                                  bool adjoint, cuComplex *phase,
-                                  std::size_t thread_per_block,
-                                  cudaStream_t stream_id) {
-    if (adjoint) {
-        cGlobalPhaseStateVector_CUDA_call<cuComplex, std::size_t, true>(
-            sv, num_sv, phase, thread_per_block, stream_id);
-    } else {
-        cGlobalPhaseStateVector_CUDA_call<cuComplex, std::size_t, false>(
-            sv, num_sv, phase, thread_per_block, stream_id);
-    }
-}
-void cGlobalPhaseStateVector_CUDA(cuDoubleComplex *sv, std::size_t num_sv,
-                                  bool adjoint, cuDoubleComplex *phase,
-                                  std::size_t thread_per_block,
-                                  cudaStream_t stream_id) {
-    if (adjoint) {
-        cGlobalPhaseStateVector_CUDA_call<cuDoubleComplex, std::size_t, true>(
-            sv, num_sv, phase, thread_per_block, stream_id);
-    } else {
-        cGlobalPhaseStateVector_CUDA_call<cuDoubleComplex, std::size_t, false>(
-            sv, num_sv, phase, thread_per_block, stream_id);
-    }
-}
-
 } // namespace Pennylane::LightningGPU
\ No newline at end of file

From e06cd311154842578889794de6fd17f74970c090 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Wed, 16 Oct 2024 23:24:52 +0000
Subject: [PATCH 06/66] fix applyControlledMatrix

---
 .../src/simulators/lightning_gpu/StateVectorCudaManaged.hpp | 5 +++--
 tests/lightning_qubit/test_adjoint_jacobian_class.py        | 2 +-
 tests/lightning_qubit/test_measurements_class.py            | 2 +-
 tests/test_adjoint_jacobian.py                              | 4 ++--
 tests/test_gates.py                                         | 6 +++---
 5 files changed, 10 insertions(+), 9 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index 2d8a0f39d..3c3fddd08 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -291,7 +291,8 @@ class StateVectorCudaManaged
         if (opName == "Identity") {
             return;
         } else if (opName == "GlobalPhase") {
-            //TODO: Check the performance GlobalPhase implementation via scaling the statevector
+            // TODO: Check the performance GlobalPhase implementation via
+            // scaling the statevector
             const std::vector<std::string> names(BaseType::getNumQubits(), "I");
             std::vector<std::size_t> tgts_all(BaseType::getNumQubits());
             std::iota(tgts_all.begin(), tgts_all.end(), 0);
@@ -486,7 +487,7 @@ class StateVectorCudaManaged
                 return static_cast<int>(BaseType::getNumQubits() - 1 - x);
             });
 
-        std::transform(ctrls_values.begin(), ctrls_values.end(),
+        std::transform(ctrls_values.rbegin(), ctrls_values.rend(),
                        ctrls_valuesInt.begin(),
                        [&](bool x) { return static_cast<int>(x); });
 
diff --git a/tests/lightning_qubit/test_adjoint_jacobian_class.py b/tests/lightning_qubit/test_adjoint_jacobian_class.py
index 7fbaa19b1..3bff0e64a 100644
--- a/tests/lightning_qubit/test_adjoint_jacobian_class.py
+++ b/tests/lightning_qubit/test_adjoint_jacobian_class.py
@@ -166,7 +166,7 @@ def test_not_expectation_return_type(self, lightning_sv):
             self.calculate_jacobian(lightning_sv(num_wires=1), tape)
 
     @pytest.mark.skipif(
-        device_name != "lightning.qubit",
+        device_name not in ("lightning.qubit", "lightning.gpu"),
         reason="N-controlled operations only implemented in lightning.qubit.",
     )
     @pytest.mark.parametrize("n_qubits", [1, 2, 3, 4])
diff --git a/tests/lightning_qubit/test_measurements_class.py b/tests/lightning_qubit/test_measurements_class.py
index 9a606e66a..e418936c5 100644
--- a/tests/lightning_qubit/test_measurements_class.py
+++ b/tests/lightning_qubit/test_measurements_class.py
@@ -799,7 +799,7 @@ def test_controlled_qubit_gates(self, operation, n_qubits, control_value, tol, l
                     assert np.allclose(result, expected, tol * 10)
 
     @pytest.mark.skipif(
-        device_name not in ("lightning.qubit", "lightning.tensor"),
+        device_name not in ("lightning.qubit", "lightning.tensor", "lightning.gpu"),
         reason="N-controlled operations only implemented in lightning.qubit.",
     )
     def test_controlled_qubit_unitary_from_op(self, tol, lightning_sv):
diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py
index 83713b1c6..8dff699b4 100644
--- a/tests/test_adjoint_jacobian.py
+++ b/tests/test_adjoint_jacobian.py
@@ -268,7 +268,7 @@ def test_Rot_gradient(self, stateprep, theta, dev):
         assert np.allclose(calculated_val, numeric_val, atol=tol, rtol=0)
 
     @pytest.mark.skipif(
-        device_name != "lightning.qubit",
+        device_name not in ("lightning.qubit", "lightning.gpu"),
         reason="N-controlled operations only implemented in lightning.qubit.",
     )
     @pytest.mark.parametrize("n_qubits", [1, 2, 3, 4])
@@ -817,7 +817,7 @@ def circuit(p):
             assert np.allclose(jac_ad, jac_bp, atol=tol, rtol=0)
 
     @pytest.mark.skipif(
-        device_name != "lightning.qubit",
+        device_name not in ("lightning.qubit", "lightning.gpu"),
         reason="N-controlled operations only implemented in lightning.qubit.",
     )
     @pytest.mark.parametrize(
diff --git a/tests/test_gates.py b/tests/test_gates.py
index da414c378..d586d35ac 100644
--- a/tests/test_gates.py
+++ b/tests/test_gates.py
@@ -372,7 +372,7 @@ def test_state_prep(n_targets, tol):
 
 
 @pytest.mark.skipif(
-    device_name != "lightning.qubit",
+    device_name not in ("lightning.qubit", "lightning.gpu"),
     reason="N-controlled operations only implemented in lightning.qubit.",
 )
 @pytest.mark.parametrize("control_value", [False, True])
@@ -497,7 +497,7 @@ def circuit():
 
 
 @pytest.mark.skipif(
-    device_name != "lightning.qubit",
+    device_name not in ("lightning.qubit", "lightning.gpu"),
     reason="N-controlled operations only implemented in lightning.qubit.",
 )
 def test_controlled_qubit_unitary_from_op(tol):
@@ -558,7 +558,7 @@ def test_paulirot(n_wires, n_targets, tol):
 
 
 @pytest.mark.skipif(
-    device_name not in ("lightning.qubit", "lightning.tensor"),
+    device_name not in ("lightning.qubit", "lightning.tensor", "lightning.gpu"),
     reason="N-controlled operations only implemented in lightning.qubit.",
 )
 @pytest.mark.parametrize("control_wires", range(4))

From ad6a17919d9b138ba13f1a5a69e9fa3bb09cd75b Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 14:02:10 +0000
Subject: [PATCH 07/66] update unit tests

---
 .../Test_StateVectorCudaManaged_Param.cpp     |  7 +++++--
 .../tests/Test_StateVectorCudaManaged.cpp     | 21 +++++++------------
 2 files changed, 12 insertions(+), 16 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp
index e2485910d..adc25eff5 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp
@@ -1564,7 +1564,6 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyControlledGlobalPhase",
     std::mt19937_64 re{1337};
     const std::size_t num_qubits = 3;
     const bool inverse = GENERATE(false, true);
-    const std::size_t index = GENERATE(0, 1, 2);
     /* The `phase` array contains the diagonal entries of the controlled-phase
        operator. It can be created in Python using the following command
 
@@ -1582,7 +1581,11 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyControlledGlobalPhase",
     auto sv_data = createRandomStateVectorData<TestType>(re, num_qubits);
     StateVectorCudaManaged<TestType> sv(
         reinterpret_cast<ComplexT *>(sv_data.data()), sv_data.size());
-    sv.applyOperation("C(GlobalPhase)", {index}, inverse, {}, phase);
+    std::vector<std::size_t> ctrls = {0, 1};
+    std::vector<bool> ctrl_vals = {0, 1};
+    std::vector<std::size_t> tgts = {2};
+    const TestType param = -M_PI_2;
+    sv.applyOperation("GlobalPhase", ctrls, ctrl_vals, tgts, inverse, {param});
     auto result_sv = sv.getDataVector();
     for (std::size_t j = 0; j < exp2(num_qubits); j++) {
         ComplexT tmp = (inverse) ? conj(phase[j]) : phase[j];
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
index 841074474..70100aa5e 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
@@ -188,20 +188,13 @@ TEMPLATE_PRODUCT_TEST_CASE("StateVectorCudaManaged::applyOperations",
             state_vector.applyOperations({"PauliX", "PauliY"}, {{0}, {1}},
                                          {false}),
             LightningException, "must all be equal"); // invalid inverse
-        PL_REQUIRE_THROWS_MATCHES(
-            state_vector.applyOperation("PauliX", std::vector<std::size_t>{0},
-                                        std::vector<bool>{false},
-                                        std::vector<std::size_t>{1}, false,
-                                        {0.0}, std::vector<ComplexT>{}),
-            LightningException,
-            "Controlled kernels not implemented."); // invalid controlled_wires
-        PL_REQUIRE_THROWS_MATCHES(
-            state_vector.applyOperation("PauliX", {}, std::vector<bool>{false},
-                                        std::vector<std::size_t>{1}, false,
-                                        {0.0}, std::vector<ComplexT>{}),
-            LightningException,
-            "`controlled_wires` must have the same size "
-            "as"); // invalid controlled_wires
+        PL_REQUIRE_THROWS_MATCHES(state_vector.applyOperation(
+                                      "PauliX", {}, std::vector<bool>{false},
+                                      std::vector<std::size_t>{1}, false, {0.0},
+                                      std::vector<ComplexT>{}),
+                                  LightningException,
+                                  "`ctrls` and `ctrls_values` must have the "
+                                  "same size."); // invalid controlled_wires
     }
 
     SECTION("Test invalid arguments with parameters") {

From c4b6d858f24176f512111a31fb6e8fc052d1d52c Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 14:02:26 +0000
Subject: [PATCH 08/66] make format

---
 .../tests/Test_StateVectorCudaManaged.cpp          | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
index 70100aa5e..c14afb4e3 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
@@ -188,13 +188,13 @@ TEMPLATE_PRODUCT_TEST_CASE("StateVectorCudaManaged::applyOperations",
             state_vector.applyOperations({"PauliX", "PauliY"}, {{0}, {1}},
                                          {false}),
             LightningException, "must all be equal"); // invalid inverse
-        PL_REQUIRE_THROWS_MATCHES(state_vector.applyOperation(
-                                      "PauliX", {}, std::vector<bool>{false},
-                                      std::vector<std::size_t>{1}, false, {0.0},
-                                      std::vector<ComplexT>{}),
-                                  LightningException,
-                                  "`ctrls` and `ctrls_values` must have the "
-                                  "same size."); // invalid controlled_wires
+        PL_REQUIRE_THROWS_MATCHES(
+            state_vector.applyOperation("PauliX", {}, std::vector<bool>{false},
+                                        std::vector<std::size_t>{1}, false,
+                                        {0.0}, std::vector<ComplexT>{}),
+            LightningException,
+            "`ctrls` and `ctrls_values` must have the "
+            "same size."); // invalid controlled_wires
     }
 
     SECTION("Test invalid arguments with parameters") {

From 3ce16e6841eb3450ce99f62a0af85e8bdb67b405 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 14:54:03 +0000
Subject: [PATCH 09/66] update GlobalPhase support with `scaleC_CUDA`

---
 .../lightning_gpu/StateVectorCudaManaged.hpp     | 16 ++++++++--------
 1 file changed, 8 insertions(+), 8 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index 3c3fddd08..9c7ee3103 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -291,14 +291,14 @@ class StateVectorCudaManaged
         if (opName == "Identity") {
             return;
         } else if (opName == "GlobalPhase") {
-            // TODO: Check the performance GlobalPhase implementation via
-            // scaling the statevector
-            const std::vector<std::string> names(BaseType::getNumQubits(), "I");
-            std::vector<std::size_t> tgts_all(BaseType::getNumQubits());
-            std::iota(tgts_all.begin(), tgts_all.end(), 0);
-            std::reverse(tgts_all.begin(), tgts_all.end());
-            applyParametricPauliGate_(names, {}, tgts_all, 2 * params[0],
-                                      adjoint);
+            PrecisionT param = adjoint ? -params[0] : params[0];
+            CFP_t scale_factor{std::cos(param), -std::sin(param)};
+            scaleC_CUDA<CFP_t, CFP_t, int>(
+                scale_factor, BaseType::getDataBuffer().getData(),
+                BaseType::getDataBuffer().getLength(),
+                BaseType::getDataBuffer().getDevTag().getDeviceID(),
+                BaseType::getDataBuffer().getDevTag().getStreamID(),
+                getCublasCaller());
         } else if (native_gates_.find(opName) != native_gates_.end()) {
             applyParametricPauliGate_({opName}, ctrls, tgts, params.front(),
                                       adjoint);

From 008c4f30eee9c2950b294cf2da9596c6de82cb1e Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 15:30:12 +0000
Subject: [PATCH 10/66] remove global_phase_diagonal

---
 pennylane_lightning/lightning_gpu/_state_vector.py | 2 --
 1 file changed, 2 deletions(-)

diff --git a/pennylane_lightning/lightning_gpu/_state_vector.py b/pennylane_lightning/lightning_gpu/_state_vector.py
index dc4732bb1..93acf859b 100644
--- a/pennylane_lightning/lightning_gpu/_state_vector.py
+++ b/pennylane_lightning/lightning_gpu/_state_vector.py
@@ -43,7 +43,6 @@
 from pennylane.wires import Wires
 
 # pylint: disable=ungrouped-imports
-from pennylane_lightning.core._serialize import global_phase_diagonal
 from pennylane_lightning.core._state_vector_base import LightningBaseStateVector
 
 from ._measurements import LightningGPUMeasurements
@@ -331,7 +330,6 @@ def _apply_lightning(
                 param = operation.parameters
                 method(wires, invert_param, param)
             elif isinstance(operation, qml.ops.Controlled):  # apply n-controlled gate
-                # LGPU do not support the controlled gates except for GlobalPhase
                 self._apply_lightning_controlled(operation)
             else:  # apply gate as a matrix
                 try:

From 015b50afaf15805dad2c4451ae81d3df1cebcb72 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 17:03:39 +0000
Subject: [PATCH 11/66] remove unused lines

---
 .../lightning_gpu/StateVectorCudaManaged.hpp          | 11 +----------
 1 file changed, 1 insertion(+), 10 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index 9c7ee3103..330f7c67b 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -399,7 +399,7 @@ class StateVectorCudaManaged
             applyParametricPauliGeneralGate_({opName}, ctrlsInt,
                                              ctrls_valuesInt, tgtsInt,
                                              params.front(), adjoint);
-        } else if (opName == "Rot" || opName == "CRot") {
+        } else if (opName == "Rot") {
             if (adjoint) {
                 auto rot_matrix =
                     cuGates::getRot<CFP_t>(params[2], params[1], params[0]);
@@ -411,15 +411,6 @@ class StateVectorCudaManaged
                 applyDeviceGeneralGate_(rot_matrix.data(), ctrlsInt, tgtsInt,
                                         ctrls_valuesInt, false);
             }
-        } else if (opName == "Matrix") {
-            DataBuffer<CFP_t, int> d_matrix{
-                gate_matrix.size(), BaseType::getDataBuffer().getDevTag(),
-                true};
-            d_matrix.CopyHostDataToGpu(gate_matrix.data(), d_matrix.getLength(),
-                                       false);
-
-            applyDeviceGeneralGate_(d_matrix.getData(), ctrlsInt, tgtsInt,
-                                    ctrls_valuesInt, adjoint);
         } else if (par_gates_.find(opName) != par_gates_.end()) {
             auto &gateMap =
                 cuGates::DynamicGateDataAccess<PrecisionT>::getInstance();

From 02eb2904438d09c988a270381d2a6315b113442b Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 18:29:55 +0000
Subject: [PATCH 12/66] initial commit

---
 .../lightning_gpu/StateVectorCudaMPI.hpp      |   9 +
 .../lightning_gpu/StateVectorCudaManaged.hpp  | 184 ++++++++----------
 .../lightning_gpu/bindings/LGPUBindings.hpp   |  11 ++
 .../Test_StateVectorCudaManaged_Param.cpp     |   9 +-
 .../mpi/Test_StateVectorCudaMPI_Param.cpp     |  13 ++
 .../src/simulators/lightning_gpu/initSV.cu    | 132 +------------
 .../tests/Test_StateVectorCudaManaged.cpp     |   9 +-
 .../lightning_gpu/_state_vector.py            |  10 +-
 8 files changed, 134 insertions(+), 243 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
index 964c5e69c..f9c0f0116 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
@@ -386,6 +386,15 @@ class StateVectorCudaMPI final
 
         if (opName == "Identity") {
             return;
+        } else if (opName == "GlobalPhase") {
+            PrecisionT param = adjoint ? -params[0] : params[0];
+            CFP_t scale_factor{std::cos(param), -std::sin(param)};
+            scaleC_CUDA<CFP_t, CFP_t, int>(
+                scale_factor, BaseType::getDataBuffer().getData(),
+                BaseType::getDataBuffer().getLength(),
+                BaseType::getDataBuffer().getDevTag().getDeviceID(),
+                BaseType::getDataBuffer().getDevTag().getStreamID(),
+                getCublasCaller());
         } else if (native_gates_.find(opName) != native_gates_.end()) {
             applyParametricPauliGate({opName}, ctrls, tgts, params.front(),
                                      adjoint);
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index f5aeb4abb..2e31bbe2f 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -67,25 +67,6 @@ extern void setBasisState_CUDA(cuDoubleComplex *sv, cuDoubleComplex &value,
                                const std::size_t index, bool async,
                                cudaStream_t stream_id);
 
-extern void globalPhaseStateVector_CUDA(cuComplex *sv, std::size_t num_sv,
-                                        cuComplex phase,
-                                        std::size_t thread_per_block,
-                                        cudaStream_t stream_id);
-extern void globalPhaseStateVector_CUDA(cuDoubleComplex *sv, std::size_t num_sv,
-                                        cuDoubleComplex phase,
-                                        std::size_t thread_per_block,
-                                        cudaStream_t stream_id);
-
-extern void cGlobalPhaseStateVector_CUDA(cuComplex *sv, std::size_t num_sv,
-                                         bool adjoint, cuComplex *phase,
-                                         std::size_t thread_per_block,
-                                         cudaStream_t stream_id);
-extern void cGlobalPhaseStateVector_CUDA(cuDoubleComplex *sv,
-                                         std::size_t num_sv, bool adjoint,
-                                         cuDoubleComplex *phase,
-                                         std::size_t thread_per_block,
-                                         cudaStream_t stream_id);
-
 /**
  * @brief Managed memory CUDA state-vector class using custateVec backed
  * gate-calls.
@@ -257,44 +238,6 @@ class StateVectorCudaManaged
                                     use_async);
     }
 
-    /**
-     * @brief Multiplies the state-vector by a global phase.
-     *
-     * @param adjoint Indicates whether to use adjoint of gate.
-     * @param param Complex phase generator.
-     */
-    template <std::size_t thread_per_block = 256>
-    void globalPhaseStateVector(const bool adjoint, const Precision param) {
-        auto stream_id = BaseType::getDataBuffer().getDevTag().getStreamID();
-        std::complex<Precision> phase =
-            std::exp(std::complex<Precision>{0, (adjoint) ? param : -param});
-        auto cuPhase = complexToCu(phase);
-        globalPhaseStateVector_CUDA(BaseType::getData(), BaseType::getLength(),
-                                    cuPhase, thread_per_block, stream_id);
-    }
-
-    /**
-     * @brief Multiplies the state-vector by a controlled global phase.
-     *
-     * @param phase Controlled complex phase vector.
-     */
-    template <std::size_t thread_per_block = 256>
-    void cGlobalPhaseStateVector(const bool adjoint,
-                                 const std::vector<CFP_t> &phase,
-                                 const bool async = false) {
-        PL_ABORT_IF_NOT(BaseType::getLength() == phase.size(),
-                        "The state-vector data must have the same size as the "
-                        "controlled-phase data.")
-        auto device_id = BaseType::getDataBuffer().getDevTag().getDeviceID();
-        auto stream_id = BaseType::getDataBuffer().getDevTag().getStreamID();
-        DataBuffer<CFP_t, int> d_phase{phase.size(), device_id, stream_id,
-                                       true};
-        d_phase.CopyHostDataToGpu(phase.data(), d_phase.getLength(), async);
-        cGlobalPhaseStateVector_CUDA(BaseType::getData(), BaseType::getLength(),
-                                     adjoint, d_phase.getData(),
-                                     thread_per_block, stream_id);
-    }
-
     /**
      * @brief Apply a single gate to the state-vector. Offloads to custatevec
      * specific API calls if available. If unable, attempts to use prior cached
@@ -347,13 +290,18 @@ class StateVectorCudaManaged
                                             wires.end()};
         if (opName == "Identity") {
             return;
-        } else if (opName == "C(GlobalPhase)") {
-            cGlobalPhaseStateVector(adjoint, gate_matrix);
         } else if (opName == "GlobalPhase") {
-            globalPhaseStateVector(adjoint, params[0]);
+            PrecisionT param = adjoint ? -params[0] : params[0];
+            CFP_t scale_factor{std::cos(param), -std::sin(param)};
+            scaleC_CUDA<CFP_t, CFP_t, int>(
+                scale_factor, BaseType::getDataBuffer().getData(),
+                BaseType::getDataBuffer().getLength(),
+                BaseType::getDataBuffer().getDevTag().getDeviceID(),
+                BaseType::getDataBuffer().getDevTag().getStreamID(),
+                getCublasCaller());
         } else if (native_gates_.find(opName) != native_gates_.end()) {
-            applyParametricPauliGate({opName}, ctrls, tgts, params.front(),
-                                     adjoint);
+            applyParametricPauliGate_({opName}, ctrls, tgts, params.front(),
+                                      adjoint);
         } else if (opName == "Rot" || opName == "CRot") {
             if (adjoint) {
                 auto rot_matrix =
@@ -407,25 +355,44 @@ class StateVectorCudaManaged
      * @param opName Name of gate to apply.
      * @param controlled_wires Control wires.
      * @param controlled_values Control values (false or true).
-     * @param wires Wires to apply gate to.
-     * @param inverse Indicates whether to use adjoint of gate.
+     * @param tgt_wires Wires to apply gate to.
+     * @param adjoint Indicates whether to use adjoint of gate.
      * @param params Optional parameter list for parametric gates.
      * @param params Optional std gate matrix if opName doesn't exist.
      */
-    template <template <typename...> class complex_t>
-    void
-    applyOperation(const std::string &opName,
-                   const std::vector<std::size_t> &controlled_wires,
-                   const std::vector<bool> &controlled_values,
-                   const std::vector<std::size_t> &wires, bool inverse = false,
-                   const std::vector<Precision> &params = {0.0},
-                   const std::vector<complex_t<Precision>> &gate_matrix = {}) {
-        PL_ABORT_IF_NOT(controlled_wires.empty(),
-                        "Controlled kernels not implemented.");
-        PL_ABORT_IF_NOT(controlled_wires.size() == controlled_values.size(),
-                        "`controlled_wires` must have the same size as "
-                        "`controlled_values`.");
-        applyOperation(opName, wires, inverse, params, gate_matrix);
+    void applyOperation(
+        const std::string &opName, const std::vector<std::size_t> &ctrl_wires,
+        const std::vector<bool> &ctrl_values,
+        const std::vector<std::size_t> &tgt_wires, bool adjoint = false,
+        const std::vector<Precision> &params = {0.0},
+        [[maybe_unused]] const std::vector<ComplexT> &gate_matrix = {}) {
+        PL_ABORT_IF_NOT(opName == "GlobalPhase",
+                        "Only GlobalPhase gate is supported.");
+        PL_ABORT_IF(ctrl_wires.size() != ctrl_values.size(),
+                    "`ctrls` and `ctrls_values` must have the same size.");
+        std::vector<int> ctrlsInt(ctrl_wires.size());
+        std::vector<int> tgtsInt(tgt_wires.size());
+        std::vector<int> ctrls_valuesInt(ctrl_wires.size());
+
+        std::transform(ctrl_wires.begin(), ctrl_wires.end(), ctrlsInt.begin(),
+                       [&](std::size_t x) {
+                           return static_cast<int>(BaseType::getNumQubits() -
+                                                   1 - x);
+                       });
+        std::transform(tgt_wires.begin(), tgt_wires.end(), tgtsInt.begin(),
+                       [&](std::size_t x) {
+                           return static_cast<int>(BaseType::getNumQubits() -
+                                                   1 - x);
+                       });
+
+        std::transform(ctrl_values.begin(), ctrl_values.end(),
+                       ctrls_valuesInt.begin(),
+                       [&](bool x) { return static_cast<int>(x); });
+        if (opName == "GlobalPhase") {
+            const std::vector<std::string> names(tgt_wires.size(), "I");
+            applyParametricPauliGeneralGate_(names, ctrlsInt, ctrls_valuesInt,
+                                             tgtsInt, 2 * params[0], adjoint);
+        }
     }
 
     /**
@@ -595,20 +562,20 @@ class StateVectorCudaManaged
     inline void applyRX(const std::vector<std::size_t> &wires, bool adjoint,
                         Precision param) {
         static const std::vector<std::string> name{{"RX"}};
-        applyParametricPauliGate(name, {wires.begin(), wires.end() - 1},
-                                 {wires.back()}, param, adjoint);
+        applyParametricPauliGate_(name, {wires.begin(), wires.end() - 1},
+                                  {wires.back()}, param, adjoint);
     }
     inline void applyRY(const std::vector<std::size_t> &wires, bool adjoint,
                         Precision param) {
         static const std::vector<std::string> name{{"RY"}};
-        applyParametricPauliGate(name, {wires.begin(), wires.end() - 1},
-                                 {wires.back()}, param, adjoint);
+        applyParametricPauliGate_(name, {wires.begin(), wires.end() - 1},
+                                  {wires.back()}, param, adjoint);
     }
     inline void applyRZ(const std::vector<std::size_t> &wires, bool adjoint,
                         Precision param) {
         static const std::vector<std::string> name{{"RZ"}};
-        applyParametricPauliGate(name, {wires.begin(), wires.end() - 1},
-                                 {wires.back()}, param, adjoint);
+        applyParametricPauliGate_(name, {wires.begin(), wires.end() - 1},
+                                  {wires.back()}, param, adjoint);
     }
     inline void applyRot(const std::vector<std::size_t> &wires, bool adjoint,
                          Precision param0, Precision param1, Precision param2) {
@@ -664,17 +631,17 @@ class StateVectorCudaManaged
     inline void applyIsingXX(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
         static const std::vector<std::string> names(wires.size(), {"RX"});
-        applyParametricPauliGate(names, {}, wires, param, adjoint);
+        applyParametricPauliGate_(names, {}, wires, param, adjoint);
     }
     inline void applyIsingYY(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
         static const std::vector<std::string> names(wires.size(), {"RY"});
-        applyParametricPauliGate(names, {}, wires, param, adjoint);
+        applyParametricPauliGate_(names, {}, wires, param, adjoint);
     }
     inline void applyIsingZZ(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
         static const std::vector<std::string> names(wires.size(), {"RZ"});
-        applyParametricPauliGate(names, {}, wires, param, adjoint);
+        applyParametricPauliGate_(names, {}, wires, param, adjoint);
     }
     inline void applyIsingXY(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
@@ -789,7 +756,7 @@ class StateVectorCudaManaged
     inline void applyMultiRZ(const std::vector<std::size_t> &wires,
                              bool adjoint, Precision param) {
         const std::vector<std::string> names(wires.size(), {"RZ"});
-        applyParametricPauliGate(names, {}, wires, param, adjoint);
+        applyParametricPauliGate_(names, {}, wires, param, adjoint);
     }
 
     /* Gate generators */
@@ -1464,12 +1431,10 @@ class StateVectorCudaManaged
      * @param tgts target wires.
      * @param use_adjoint Take adjoint of operation.
      */
-    void applyParametricPauliGate(const std::vector<std::string> &pauli_words,
-                                  std::vector<std::size_t> ctrls,
-                                  std::vector<std::size_t> tgts,
-                                  Precision param, bool use_adjoint = false) {
-        int nIndexBits = BaseType::getNumQubits();
-
+    void applyParametricPauliGate_(const std::vector<std::string> &pauli_words,
+                                   std::vector<std::size_t> ctrls,
+                                   std::vector<std::size_t> tgts,
+                                   Precision param, bool use_adjoint = false) {
         std::vector<int> ctrlsInt(ctrls.size());
         std::vector<int> tgtsInt(tgts.size());
 
@@ -1483,6 +1448,29 @@ class StateVectorCudaManaged
                 return static_cast<int>(BaseType::getNumQubits() - 1 - x);
             });
 
+        const std::vector<int> ctrls_valuesInt(ctrls.size(), 1);
+
+        applyParametricPauliGeneralGate_(pauli_words, ctrlsInt, ctrls_valuesInt,
+                                         tgtsInt, param, use_adjoint);
+    }
+
+    /**
+     * @brief Apply a parametric Pauli gate using custateVec calls.
+     *
+     * @param pauli_words List of Pauli words representing operation.
+     * @param ctrls Control wires
+     * @param ctrls_values Control values
+     * @param tgts target wires.
+     * @param param Rotation angle.
+     * @param use_adjoint Take adjoint of operation.
+     */
+    void applyParametricPauliGeneralGate_(
+        const std::vector<std::string> &pauli_words,
+        const std::vector<int> &ctrlsInt,
+        const std::vector<int> &ctrls_valuesInt, const std::vector<int> tgtsInt,
+        Precision param, bool use_adjoint = false) {
+        int nIndexBits = BaseType::getNumQubits();
+
         cudaDataType_t data_type;
 
         if constexpr (std::is_same_v<CFP_t, cuDoubleComplex> ||
@@ -1491,14 +1479,12 @@ class StateVectorCudaManaged
         } else {
             data_type = CUDA_C_32F;
         }
-
         std::vector<custatevecPauli_t> pauli_enums;
         pauli_enums.reserve(pauli_words.size());
         for (const auto &pauli_str : pauli_words) {
             pauli_enums.push_back(native_gates_.at(pauli_str));
         }
         const auto local_angle = (use_adjoint) ? param / 2 : -param / 2;
-
         PL_CUSTATEVEC_IS_SUCCESS(custatevecApplyPauliRotation(
             /* custatevecHandle_t */ handle_.get(),
             /* void* */ BaseType::getData(),
@@ -1507,10 +1493,10 @@ class StateVectorCudaManaged
             /* double */ local_angle,
             /* const custatevecPauli_t* */ pauli_enums.data(),
             /* const int32_t* */ tgtsInt.data(),
-            /* const uint32_t */ tgts.size(),
+            /* const uint32_t */ tgtsInt.size(),
             /* const int32_t* */ ctrlsInt.data(),
-            /* const int32_t* */ nullptr,
-            /* const uint32_t */ ctrls.size()));
+            /* const int32_t* */ ctrls_valuesInt.data(),
+            /* const uint32_t */ ctrlsInt.size()));
         PL_CUDA_IS_SUCCESS(cudaStreamSynchronize(
             BaseType::getDataBuffer().getDevTag().getStreamID()));
     }
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp
index 145097b30..a755f61d8 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp
@@ -152,6 +152,17 @@ void registerBackendClassSpecificBindings(PyClass &pyclass) {
             "Initialize the statevector data to the |0...0> state")
         .def("collapse", &StateVectorT::collapse,
              "Collapse the statevector onto the 0 or 1 branch of a given wire.")
+        .def(
+            "apply",
+            [](StateVectorT &sv, const std::string &gate_name,
+               const std::vector<std::size_t> &controlled_wires,
+               const std::vector<bool> &controlled_values,
+               const std::vector<std::size_t> &wires, bool inverse,
+               const std::vector<ParamT> &params) {
+                sv.applyOperation(gate_name, controlled_wires,
+                                  controlled_values, wires, inverse, params);
+            },
+            "Apply operation via the gate matrix")
         .def(
             "apply",
             [](StateVectorT &sv, const std::string &str,
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp
index e2485910d..6a0de4efd 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp
@@ -1564,7 +1564,6 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyControlledGlobalPhase",
     std::mt19937_64 re{1337};
     const std::size_t num_qubits = 3;
     const bool inverse = GENERATE(false, true);
-    const std::size_t index = GENERATE(0, 1, 2);
     /* The `phase` array contains the diagonal entries of the controlled-phase
        operator. It can be created in Python using the following command
 
@@ -1582,7 +1581,11 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyControlledGlobalPhase",
     auto sv_data = createRandomStateVectorData<TestType>(re, num_qubits);
     StateVectorCudaManaged<TestType> sv(
         reinterpret_cast<ComplexT *>(sv_data.data()), sv_data.size());
-    sv.applyOperation("C(GlobalPhase)", {index}, inverse, {}, phase);
+    std::vector<std::size_t> ctrls = {0, 1};
+    std::vector<bool> ctrl_vals = {0, 1};
+    std::vector<std::size_t> tgts = {2};
+    const TestType param = -M_PI_2;
+    sv.applyOperation("GlobalPhase", ctrls, ctrl_vals, tgts, inverse, {param});
     auto result_sv = sv.getDataVector();
     for (std::size_t j = 0; j < exp2(num_qubits); j++) {
         ComplexT tmp = (inverse) ? conj(phase[j]) : phase[j];
@@ -1590,4 +1593,4 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyControlledGlobalPhase",
         CHECK((real(result_sv[j])) == Approx(real(tmp)));
         CHECK((imag(result_sv[j])) == Approx(imag(tmp)));
     }
-}
\ No newline at end of file
+}
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/mpi/Test_StateVectorCudaMPI_Param.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/mpi/Test_StateVectorCudaMPI_Param.cpp
index 17cf43e84..329b9f754 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/mpi/Test_StateVectorCudaMPI_Param.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/mpi/Test_StateVectorCudaMPI_Param.cpp
@@ -356,6 +356,19 @@ TEMPLATE_TEST_CASE("StateVectorCudaMPI::DoubleExcitationPlus",
         msb_4qbit, angle_1param);
 }
 
+TEMPLATE_TEST_CASE("StateVectorCudaMPI::GlobalPhase",
+                   "[StateVectorCudaMPI_Param]", float, double) {
+    PLGPU_MPI_TEST_GATE_OPS_PARAM(TestType, num_qubits,
+                                  applyDoubleExcitationPlus, "GlobalPhase",
+                                  lsb_4qbit, angle_1param);
+    PLGPU_MPI_TEST_GATE_OPS_PARAM(TestType, num_qubits,
+                                  applyDoubleExcitationPlus, "GlobalPhase",
+                                  mlsb_4qbit, angle_1param);
+    PLGPU_MPI_TEST_GATE_OPS_PARAM(TestType, num_qubits,
+                                  applyDoubleExcitationPlus, "GlobalPhase",
+                                  msb_4qbit, angle_1param);
+}
+
 TEMPLATE_TEST_CASE("LightningGPUMPI:applyOperation", "[LightningGPUMPI_Param]",
                    float, double) {
     using StateVectorT = StateVectorCudaMPI<TestType>;
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/initSV.cu b/pennylane_lightning/core/src/simulators/lightning_gpu/initSV.cu
index 8a62e89e8..10f56e6a8 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/initSV.cu
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/initSV.cu
@@ -106,96 +106,6 @@ void setStateVector_CUDA_call(GPUDataT *sv, index_type &num_indices,
     PL_CUDA_IS_SUCCESS(cudaGetLastError());
 }
 
-/**
- * @brief The CUDA kernel that multiplies the state vector data on GPU device
- * by a global phase.
- *
- * @param sv Complex data pointer of state vector on device.
- * @param num_sv Number of state vector elements.
- * @param phase Complex data pointer of input values (on device).
- */
-template <class GPUDataT, class index_type>
-__global__ void globalPhaseStateVectorkernel(GPUDataT *sv, index_type num_sv,
-                                             GPUDataT phase) {
-    const unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
-    if (i < num_sv) {
-        sv[i] = Cmul(sv[i], phase);
-    }
-}
-
-/**
- * @brief The CUDA kernel call wrapper.
- *
- * @param sv Complex data pointer of state vector on device.
- * @param num_sv Number of state vector elements.
- * @param phase Constant complex phase.
- * @param thread_per_block Number of threads set per block.
- * @param stream_id Stream id of CUDA calls
- */
-template <class GPUDataT, class index_type>
-void globalPhaseStateVector_CUDA_call(GPUDataT *sv, index_type num_sv,
-                                      GPUDataT phase,
-                                      std::size_t thread_per_block,
-                                      cudaStream_t stream_id) {
-    auto dv = std::div(static_cast<long>(num_sv), thread_per_block);
-    std::size_t num_blocks = dv.quot + (dv.rem == 0 ? 0 : 1);
-    const std::size_t block_per_grid = (num_blocks == 0 ? 1 : num_blocks);
-    dim3 blockSize(thread_per_block, 1, 1);
-    dim3 gridSize(block_per_grid, 1);
-
-    globalPhaseStateVectorkernel<GPUDataT, index_type>
-        <<<gridSize, blockSize, 0, stream_id>>>(sv, num_sv, phase);
-    PL_CUDA_IS_SUCCESS(cudaGetLastError());
-}
-
-/**
- * @brief The CUDA kernel that multiplies the state vector data on GPU device
- * by a controlled global phase.
- *
- * @param sv Complex data pointer of state vector on device.
- * @param num_sv Number of state vector elements.
- * @param phase Complex data pointer of controlled global phase values (on
- * device).
- */
-template <class GPUDataT, class index_type, bool adjoint = false>
-__global__ void cGlobalPhaseStateVectorkernel(GPUDataT *sv, index_type num_sv,
-                                              GPUDataT *phase) {
-    const unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
-    if (i < num_sv) {
-        if constexpr (adjoint) {
-            sv[i] = Cmul(sv[i], Conj(phase[i]));
-        } else {
-            sv[i] = Cmul(sv[i], phase[i]);
-        }
-    }
-}
-
-/**
- * @brief The CUDA kernel call wrapper.
- *
- * @param sv Complex data pointer of state vector on device.
- * @param num_sv Number of state vector elements.
- * @param phase Complex data pointer of controlled global phase values (on
- * device).
- * @param thread_per_block Number of threads set per block.
- * @param stream_id Stream id of CUDA calls
- */
-template <class GPUDataT, class index_type, bool adjoint = false>
-void cGlobalPhaseStateVector_CUDA_call(GPUDataT *sv, index_type num_sv,
-                                       GPUDataT *phase,
-                                       std::size_t thread_per_block,
-                                       cudaStream_t stream_id) {
-    auto dv = std::div(static_cast<long>(num_sv), thread_per_block);
-    std::size_t num_blocks = dv.quot + (dv.rem == 0 ? 0 : 1);
-    const std::size_t block_per_grid = (num_blocks == 0 ? 1 : num_blocks);
-    dim3 blockSize(thread_per_block, 1, 1);
-    dim3 gridSize(block_per_grid, 1);
-
-    cGlobalPhaseStateVectorkernel<GPUDataT, index_type, adjoint>
-        <<<gridSize, blockSize, 0, stream_id>>>(sv, num_sv, phase);
-    PL_CUDA_IS_SUCCESS(cudaGetLastError());
-}
-
 /**
  * @brief CUDA runtime API call wrapper.
  *
@@ -242,44 +152,4 @@ void setBasisState_CUDA(cuDoubleComplex *sv, cuDoubleComplex &value,
                         cudaStream_t stream_id) {
     setBasisState_CUDA_call(sv, value, index, async, stream_id);
 }
-
-void globalPhaseStateVector_CUDA(cuComplex *sv, std::size_t num_sv,
-                                 cuComplex phase, std::size_t thread_per_block,
-                                 cudaStream_t stream_id) {
-    globalPhaseStateVector_CUDA_call(sv, num_sv, phase, thread_per_block,
-                                     stream_id);
-}
-void globalPhaseStateVector_CUDA(cuDoubleComplex *sv, std::size_t num_sv,
-                                 cuDoubleComplex phase,
-                                 std::size_t thread_per_block,
-                                 cudaStream_t stream_id) {
-    globalPhaseStateVector_CUDA_call(sv, num_sv, phase, thread_per_block,
-                                     stream_id);
-}
-
-void cGlobalPhaseStateVector_CUDA(cuComplex *sv, std::size_t num_sv,
-                                  bool adjoint, cuComplex *phase,
-                                  std::size_t thread_per_block,
-                                  cudaStream_t stream_id) {
-    if (adjoint) {
-        cGlobalPhaseStateVector_CUDA_call<cuComplex, std::size_t, true>(
-            sv, num_sv, phase, thread_per_block, stream_id);
-    } else {
-        cGlobalPhaseStateVector_CUDA_call<cuComplex, std::size_t, false>(
-            sv, num_sv, phase, thread_per_block, stream_id);
-    }
-}
-void cGlobalPhaseStateVector_CUDA(cuDoubleComplex *sv, std::size_t num_sv,
-                                  bool adjoint, cuDoubleComplex *phase,
-                                  std::size_t thread_per_block,
-                                  cudaStream_t stream_id) {
-    if (adjoint) {
-        cGlobalPhaseStateVector_CUDA_call<cuDoubleComplex, std::size_t, true>(
-            sv, num_sv, phase, thread_per_block, stream_id);
-    } else {
-        cGlobalPhaseStateVector_CUDA_call<cuDoubleComplex, std::size_t, false>(
-            sv, num_sv, phase, thread_per_block, stream_id);
-    }
-}
-
-} // namespace Pennylane::LightningGPU
\ No newline at end of file
+} // namespace Pennylane::LightningGPU
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
index 841074474..6d2f17358 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
@@ -194,14 +194,15 @@ TEMPLATE_PRODUCT_TEST_CASE("StateVectorCudaManaged::applyOperations",
                                         std::vector<std::size_t>{1}, false,
                                         {0.0}, std::vector<ComplexT>{}),
             LightningException,
-            "Controlled kernels not implemented."); // invalid controlled_wires
+            "Only GlobalPhase gate is supported."); // invalid controlled_wires
         PL_REQUIRE_THROWS_MATCHES(
-            state_vector.applyOperation("PauliX", {}, std::vector<bool>{false},
+            state_vector.applyOperation("GlobalPhase", {},
+                                        std::vector<bool>{false},
                                         std::vector<std::size_t>{1}, false,
                                         {0.0}, std::vector<ComplexT>{}),
             LightningException,
-            "`controlled_wires` must have the same size "
-            "as"); // invalid controlled_wires
+            "`ctrls` and `ctrls_values` must have the "
+            "same size."); // invalid controlled_wires
     }
 
     SECTION("Test invalid arguments with parameters") {
diff --git a/pennylane_lightning/lightning_gpu/_state_vector.py b/pennylane_lightning/lightning_gpu/_state_vector.py
index 77e453778..a3787d40f 100644
--- a/pennylane_lightning/lightning_gpu/_state_vector.py
+++ b/pennylane_lightning/lightning_gpu/_state_vector.py
@@ -43,7 +43,6 @@
 from pennylane.wires import Wires
 
 # pylint: disable=ungrouped-imports
-from pennylane_lightning.core._serialize import global_phase_diagonal
 from pennylane_lightning.core._state_vector_base import LightningBaseStateVector
 
 from ._measurements import LightningGPUMeasurements
@@ -241,15 +240,14 @@ def _apply_lightning_controlled(self, operation):
         """
         state = self.state_vector
 
+        basename = operation.base.name
         control_wires = list(operation.control_wires)
         control_values = operation.control_values
-        name = operation.name
+        target_wires = list(operation.target_wires)
         # Apply GlobalPhase
         inv = False
-        param = operation.parameters[0]
-        wires = self.wires.indices(operation.wires)
-        matrix = global_phase_diagonal(param, self.wires, control_wires, control_values)
-        state.apply(name, wires, inv, [[param]], matrix)
+        param = operation.parameters
+        state.apply(basename, control_wires, control_values, target_wires, inv, param)
 
     def _apply_lightning_midmeasure(
         self, operation: MidMeasureMP, mid_measurements: dict, postselect_mode: str

From dd2dbab6ac22dbcbc6881f30f0ac737fa065428f Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Thu, 17 Oct 2024 18:30:51 +0000
Subject: [PATCH 13/66] Auto update version from '0.39.0-dev45' to
 '0.39.0-dev46'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index 93371827c..37f7ba1bb 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev45"
+__version__ = "0.39.0-dev46"

From a77fef35e392f27c55f616790e58aa98c4a453aa Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 19:06:13 +0000
Subject: [PATCH 14/66] add py unit tests for globalphase gate

---
 mpitests/test_apply.py | 11 +++++++++++
 1 file changed, 11 insertions(+)

diff --git a/mpitests/test_apply.py b/mpitests/test_apply.py
index 5987626f1..4fe12703b 100644
--- a/mpitests/test_apply.py
+++ b/mpitests/test_apply.py
@@ -235,6 +235,17 @@ def test_apply_operation_4gatequbit_1param_gate_qnode_param(
     ):
         apply_operation_gates_qnode_param(tol, dev_mpi, operation, par, Wires)
 
+    @pytest.mark.parametrize(
+        "operation",
+        [qml.GlobalPhase],
+    )
+    @pytest.mark.parametrize("par", [[0.13], [0.2], [0.3]])
+    def test_apply_global_phase(self, tol, operation, par, dev_mpi):
+        """Test applying the GlobalPhase operation."""
+        Wires = range(numQubits)
+
+        apply_operation_gates_qnode_param(tol, dev_mpi, operation, par, Wires)
+
     # BasisState test
     @pytest.mark.parametrize("operation", [qml.BasisState])
     @pytest.mark.parametrize("index", range(numQubits))

From dffb78f7d04331d48ee96fcac6832c9247083a57 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 19:25:50 +0000
Subject: [PATCH 15/66] raise error for mpi-cGlobalPhase

---
 mpitests/test_device.py                            | 11 +++++++++++
 pennylane_lightning/lightning_gpu/_state_vector.py |  2 ++
 2 files changed, 13 insertions(+)

diff --git a/mpitests/test_device.py b/mpitests/test_device.py
index dd783dbee..9169a0003 100644
--- a/mpitests/test_device.py
+++ b/mpitests/test_device.py
@@ -17,6 +17,8 @@
 # pylint: disable=protected-access,unused-variable,missing-function-docstring,c-extension-no-member
 
 import pennylane as qml
+from pennylane import DeviceError
+from pennylane.tape import QuantumScript
 import pytest
 from conftest import LightningDevice as ld
 from conftest import device_name
@@ -52,3 +54,12 @@ def test_unsupported_mpi_buf_size():
         match="Number of processes should be smaller than the number of statevector elements",
     ):
         dev = qml.device(device_name, mpi=True, wires=1)
+
+def test_unsupported_gate():
+    comm = MPI.COMM_WORLD
+    dev = qml.device(device_name, mpi=True, wires=4)
+    op = qml.ctrl(qml.GlobalPhase(0.1, wires=[1,2,3]), [0], control_values=[True])
+    tape = QuantumScript([op])
+    with pytest.raises(DeviceError, match="Lightning-GPU-MPI does not support Controlled GlobalPhase gates"):
+        dev.execute(tape)
+        comm.Barrier()
diff --git a/pennylane_lightning/lightning_gpu/_state_vector.py b/pennylane_lightning/lightning_gpu/_state_vector.py
index a3787d40f..e2999e398 100644
--- a/pennylane_lightning/lightning_gpu/_state_vector.py
+++ b/pennylane_lightning/lightning_gpu/_state_vector.py
@@ -323,6 +323,8 @@ def _apply_lightning(
                 operation.base, qml.GlobalPhase
             ):  # apply n-controlled gate
                 # LGPU do not support the controlled gates except for GlobalPhase
+                if self._mpi_handler.use_mpi: 
+                    raise DeviceError("Lightning-GPU-MPI does not support Controlled GlobalPhase gates.")
                 self._apply_lightning_controlled(operation)
             else:  # apply gate as a matrix
                 try:

From 2641cfc53cbc9f85a527ed4c550481188d7cbb4d Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 19:27:33 +0000
Subject: [PATCH 16/66] make format

---
 mpitests/test_device.py                            | 11 +++++++----
 pennylane_lightning/lightning_gpu/_state_vector.py |  6 ++++--
 2 files changed, 11 insertions(+), 6 deletions(-)

diff --git a/mpitests/test_device.py b/mpitests/test_device.py
index 9169a0003..666b7602b 100644
--- a/mpitests/test_device.py
+++ b/mpitests/test_device.py
@@ -17,12 +17,12 @@
 # pylint: disable=protected-access,unused-variable,missing-function-docstring,c-extension-no-member
 
 import pennylane as qml
-from pennylane import DeviceError
-from pennylane.tape import QuantumScript
 import pytest
 from conftest import LightningDevice as ld
 from conftest import device_name
 from mpi4py import MPI
+from pennylane import DeviceError
+from pennylane.tape import QuantumScript
 
 if not ld._CPP_BINARY_AVAILABLE:
     pytest.skip("No binary module found. Skipping.", allow_module_level=True)
@@ -55,11 +55,14 @@ def test_unsupported_mpi_buf_size():
     ):
         dev = qml.device(device_name, mpi=True, wires=1)
 
+
 def test_unsupported_gate():
     comm = MPI.COMM_WORLD
     dev = qml.device(device_name, mpi=True, wires=4)
-    op = qml.ctrl(qml.GlobalPhase(0.1, wires=[1,2,3]), [0], control_values=[True])
+    op = qml.ctrl(qml.GlobalPhase(0.1, wires=[1, 2, 3]), [0], control_values=[True])
     tape = QuantumScript([op])
-    with pytest.raises(DeviceError, match="Lightning-GPU-MPI does not support Controlled GlobalPhase gates"):
+    with pytest.raises(
+        DeviceError, match="Lightning-GPU-MPI does not support Controlled GlobalPhase gates"
+    ):
         dev.execute(tape)
         comm.Barrier()
diff --git a/pennylane_lightning/lightning_gpu/_state_vector.py b/pennylane_lightning/lightning_gpu/_state_vector.py
index e2999e398..c4bce5db3 100644
--- a/pennylane_lightning/lightning_gpu/_state_vector.py
+++ b/pennylane_lightning/lightning_gpu/_state_vector.py
@@ -323,8 +323,10 @@ def _apply_lightning(
                 operation.base, qml.GlobalPhase
             ):  # apply n-controlled gate
                 # LGPU do not support the controlled gates except for GlobalPhase
-                if self._mpi_handler.use_mpi: 
-                    raise DeviceError("Lightning-GPU-MPI does not support Controlled GlobalPhase gates.")
+                if self._mpi_handler.use_mpi:
+                    raise DeviceError(
+                        "Lightning-GPU-MPI does not support Controlled GlobalPhase gates."
+                    )
                 self._apply_lightning_controlled(operation)
             else:  # apply gate as a matrix
                 try:

From 56b26e28de2fbc4aabdc2c2bacbf0ded9bcc8d5b Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 20:07:01 +0000
Subject: [PATCH 17/66] add changelog

---
 .github/CHANGELOG.md | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/.github/CHANGELOG.md b/.github/CHANGELOG.md
index 1f4059f9e..5486e4dba 100644
--- a/.github/CHANGELOG.md
+++ b/.github/CHANGELOG.md
@@ -46,6 +46,9 @@
 
 ### Improvements
 
+* Optimize `GlobalPhase` and `C(GlobalPhase)` gate implementation in `lightning.gpu`.
+  [(#946)](https://github.com/PennyLaneAI/pennylane-lightning/pull/946)
+
 * Optimize the cartesian product to reduce the amount of memory necessary to set the StatePrep with LightningTensor. 
   [(#943)](https://github.com/PennyLaneAI/pennylane-lightning/pull/943)
 

From 38b762e40ddeedbea3ae8a33d66dd0880ffc2d5e Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 17 Oct 2024 20:38:15 +0000
Subject: [PATCH 18/66] Trigger CIs


From a0abe2e8d254084692588c55d6b8429b82aa71e6 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Fri, 18 Oct 2024 13:27:36 +0000
Subject: [PATCH 19/66] quick fix for the setup.py

---
 setup.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/setup.py b/setup.py
index 7e6e080e9..d6505b76d 100644
--- a/setup.py
+++ b/setup.py
@@ -157,7 +157,7 @@ def build_extension(self, ext: CMakeExtension):
         )
 
         # Ensure that catalyst shared object is copied to the build directory for pip editable install
-        if backend in ("lightning_kokkos"):
+        if platform.system() == "Linux" and backend in ("lightning_kokkos"):
             source = os.path.join(f"{extdir}", f"lib{backend}_catalyst.so")
             destination = os.path.join(os.getcwd(), "build")
             shutil.copy(source, destination)

From 2f65169b4afb8294fe943fab9ce04952b50577e1 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Fri, 18 Oct 2024 17:10:54 +0000
Subject: [PATCH 20/66] add C++ unit tests

---
 .../Test_StateVectorCudaManaged_NonParam.cpp  | 208 +++++++++++++++++-
 1 file changed, 207 insertions(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
index af864d8b0..45212573e 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
@@ -1105,4 +1105,210 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::SetIthStates",
 
         CHECK(expected_state == Pennylane::Util::approx(sv.getDataVector()));
     }
-}
\ No newline at end of file
+}
+
+TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyOperation non-param "
+                   "one-qubit with controls",
+                   "[StateVectorCudaManaged]", float, double) {
+    using PrecisionT = TestType;
+    std::mt19937 re{1337};
+    const int num_qubits = 4;
+    const auto margin = PrecisionT{1e-5};
+    const std::size_t control = GENERATE(0, 1, 2, 3);
+    const std::size_t wire = GENERATE(0, 1, 2, 3);
+    auto st0 = createRandomStateVectorData<PrecisionT>(re, num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv0(num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv1(num_qubits);
+
+    sv0.CopyHostDataToGpu(st0.data(), st0.size());
+    sv1.CopyHostDataToGpu(st0.data(), st0.size());
+
+    DYNAMIC_SECTION("N-controlled PauliX - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            sv0.applyOperation("CNOT", {control, wire});
+            sv1.applyOperation("PauliX", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+
+        if (control != 0 && wire != 0 && control != wire) {
+            sv0.applyOperation("Toffoli", {0, control, wire});
+            sv1.applyOperation("PauliX", std::vector<std::size_t>{0, control},
+                               std::vector<bool>{true, true},
+                               std::vector<std::size_t>{wire});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+
+            sv0.applyOperation("Toffoli", {control, 0, wire});
+            sv1.applyOperation("PauliX", std::vector<std::size_t>{control, 0},
+                               std::vector<bool>{true, true},
+                               std::vector<std::size_t>{wire});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled PauliY - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            sv0.applyOperation("CY", {control, wire});
+            sv1.applyOperation("PauliY", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled PauliZ - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            sv0.applyOperation("CZ", {control, wire});
+            sv1.applyOperation("PauliZ", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled Hadamard - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            const auto matrix = getHadamard<std::complex, PrecisionT>();
+
+            sv0.applyControlledMatrix(
+                matrix.data(), matrix.size(), std::vector<std::size_t>{control},
+                std::vector<bool>{true}, std::vector<std::size_t>{wire});
+            sv1.applyOperation("Hadamard", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+    DYNAMIC_SECTION("N-controlled S - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            const auto matrix = getS<std::complex, PrecisionT>();
+
+            sv0.applyControlledMatrix(
+                matrix.data(), matrix.size(), std::vector<std::size_t>{control},
+                std::vector<bool>{true}, std::vector<std::size_t>{wire});
+            sv1.applyOperation("S", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled T - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            const std::vector<std::complex<PrecisionT>> matrix =
+                getT<std::complex, PrecisionT>();
+
+            sv0.applyControlledMatrix(
+                matrix.data(), matrix.size(), std::vector<std::size_t>{control},
+                std::vector<bool>{true}, std::vector<std::size_t>{wire});
+            sv1.applyOperation("T", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+}
+
+TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyOperation non-param "
+                   "two-qubit with controls",
+                   "[StateVectorCudaManaged]", float, double) {
+    using PrecisionT = TestType;
+    std::mt19937 re{1337};
+    const int num_qubits = 4;
+    const auto margin = PrecisionT{1e-5};
+    const std::size_t control = GENERATE(0, 1, 2, 3);
+    const std::size_t wire0 = GENERATE(0, 1, 2, 3);
+    const std::size_t wire1 = GENERATE(0, 1, 2, 3);
+    auto st0 = createRandomStateVectorData<PrecisionT>(re, num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv0(num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv1(num_qubits);
+
+    sv0.CopyHostDataToGpu(st0.data(), st0.size());
+    sv1.CopyHostDataToGpu(st0.data(), st0.size());
+
+    DYNAMIC_SECTION("N-controlled SWAP - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire0 && control != wire1 && wire0 != wire1) {
+            sv0.applyOperation("CSWAP", {control, wire0, wire1});
+            sv1.applyOperation("SWAP", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire0, wire1});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled SWAP with matrix- "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire0 && control != wire1 && wire0 != wire1) {
+            const std::vector<std::complex<PrecisionT>> matrix =
+                getSWAP<std::complex, PrecisionT>();
+            sv0.applyControlledMatrix(matrix.data(), matrix.size(),
+                                      std::vector<std::size_t>{control},
+                                      std::vector<bool>{true},
+                                      std::vector<std::size_t>{wire0, wire1});
+            sv1.applyOperation("SWAP", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire0, wire1});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+}
+
+TEMPLATE_TEST_CASE("StateVectorCudaManaged::controlled Toffoli",
+                   "[StateVectorCudaManaged]", float, double) {
+    using PrecisionT = TestType;
+    std::mt19937 re{1337};
+    const int num_qubits = 6;
+    const auto margin = PrecisionT{1e-5};
+    const std::size_t control = GENERATE(0, 1, 2);
+    auto st0 = createRandomStateVectorData<PrecisionT>(re, num_qubits);
+
+    StateVectorCudaManaged<PrecisionT> sv0(num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv1(num_qubits);
+
+    sv0.CopyHostDataToGpu(st0.data(), st0.size());
+    sv1.CopyHostDataToGpu(st0.data(), st0.size());
+
+    const std::vector<std::complex<PrecisionT>> matrix =
+        getToffoli<std::complex, PrecisionT>();
+    sv0.applyControlledMatrix(
+        matrix.data(), matrix.size(), std::vector<std::size_t>{control},
+        std::vector<bool>{true}, std::vector<std::size_t>{3, 4, 5});
+    sv1.applyOperation("PauliX", std::vector<std::size_t>{control, 3, 4},
+                       std::vector<bool>{true, true, true},
+                       std::vector<std::size_t>{5});
+    REQUIRE(sv0.getDataVector() == approx(sv1.getDataVector()).margin(margin));
+}

From d303e87a6177ad2c65ae83e81b9e966386b098db Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Fri, 18 Oct 2024 17:11:17 +0000
Subject: [PATCH 21/66] Auto update version from '0.39.0-dev46' to
 '0.39.0-dev47'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index 37f7ba1bb..de59ceba0 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev46"
+__version__ = "0.39.0-dev47"

From 95662ec59486d7b1c231c17ed8430bcb45d73db0 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Fri, 18 Oct 2024 20:42:05 +0000
Subject: [PATCH 22/66] add more C++ unit tests

---
 .../Test_StateVectorCudaManaged_Param.cpp     | 396 ++++++++++++++++++
 1 file changed, 396 insertions(+)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp
index adc25eff5..ca89ed2be 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Param.cpp
@@ -1558,6 +1558,402 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyGlobalPhase",
     }
 }
 
+TEMPLATE_TEST_CASE(
+    "StateVectorCudaManaged::applyOperation param one-qubit with controls",
+    "[StateVectorCudaManaged]", float, double) {
+    using PrecisionT = TestType;
+    std::mt19937 re{1337};
+    const int num_qubits = 4;
+    const auto margin = PrecisionT{1e-5};
+    const std::size_t control = GENERATE(0, 1, 2, 3);
+    const std::size_t wire = GENERATE(0, 1, 2, 3);
+    bool inverse = GENERATE(false, true);
+    PrecisionT param = GENERATE(-0.5, 1.5);
+
+    auto st0 = createRandomStateVectorData<PrecisionT>(re, num_qubits);
+
+    StateVectorCudaManaged<PrecisionT> sv0(num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv1(num_qubits);
+
+    sv0.CopyHostDataToGpu(st0.data(), st0.size());
+    sv1.CopyHostDataToGpu(st0.data(), st0.size());
+
+    DYNAMIC_SECTION("N-controlled PhaseShift - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            sv0.applyOperation("ControlledPhaseShift", {control, wire}, inverse,
+                               {param});
+            sv1.applyOperation("PhaseShift", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire}, inverse,
+                               {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled RX - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            sv0.applyOperation("CRX", {control, wire}, inverse, {param});
+            sv1.applyOperation("RX", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire}, inverse,
+                               {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled RY - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            sv0.applyOperation("CRY", {control, wire}, inverse, {param});
+            sv1.applyOperation("RY", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire}, inverse,
+                               {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled RZ - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            sv0.applyOperation("CRZ", {control, wire}, inverse, {param});
+            sv1.applyOperation("RZ", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire}, inverse,
+                               {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled Rot - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire) {
+            const std::vector<PrecisionT> params = {
+                param, static_cast<PrecisionT>(2.0) * param,
+                static_cast<PrecisionT>(3.0) * param};
+            sv0.applyOperation("CRot", {control, wire}, inverse, params);
+            sv1.applyOperation("Rot", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire}, inverse, params);
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+}
+
+TEMPLATE_TEST_CASE(
+    "StateVectorCudaManaged::applyOperation param two-qubits with controls",
+    "[StateVectorCudaManaged]", float, double) {
+    using PrecisionT = TestType;
+    using ComplexT = std::complex<TestType>;
+    std::mt19937 re{1337};
+    const int num_qubits = 4;
+    const auto margin = PrecisionT{1e-5};
+    const std::size_t control = GENERATE(0, 1, 2, 3);
+    const std::size_t wire0 = GENERATE(0, 1, 2, 3);
+    const std::size_t wire1 = GENERATE(0, 1, 2, 3);
+    bool inverse = GENERATE(false, true);
+    PrecisionT param = GENERATE(-0.5, 1.5);
+
+    auto st0 = createRandomStateVectorData<PrecisionT>(re, num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv0(num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv1(num_qubits);
+
+    sv0.CopyHostDataToGpu(st0.data(), st0.size());
+    sv1.CopyHostDataToGpu(st0.data(), st0.size());
+
+    auto getControlledGate = [](std::vector<ComplexT> matrix) {
+        std::vector<ComplexT> cmatrix(matrix.size() * 4);
+        for (std::size_t i = 0; i < 4; i++) {
+            cmatrix[i * 8 + i] = ComplexT{1.0};
+        }
+        for (std::size_t i = 0; i < 4; i++) {
+            for (std::size_t j = 0; j < 4; j++) {
+                cmatrix[(i + 4) * 8 + j + 4] = matrix[i * 4 + j];
+            }
+        }
+        return cmatrix;
+    };
+
+    DYNAMIC_SECTION("N-controlled IsingXX - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire0 && control != wire1 && wire0 != wire1) {
+            auto matrix = getIsingXX<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1}, inverse);
+            sv1.applyOperation("IsingXX", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire0, wire1}, inverse,
+                               {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled IsingXY - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire0 && control != wire1 && wire0 != wire1) {
+            auto matrix = getIsingXY<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1}, inverse);
+            sv1.applyOperation("IsingXY", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire0, wire1}, inverse,
+                               {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled IsingYY - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire0 && control != wire1 && wire0 != wire1) {
+            auto matrix = getIsingYY<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1}, inverse);
+            sv1.applyOperation("IsingYY", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire0, wire1}, inverse,
+                               {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled IsingZZ - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire0 && control != wire1 && wire0 != wire1) {
+            auto matrix = getIsingZZ<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1}, inverse);
+            sv1.applyOperation("IsingZZ", std::vector<std::size_t>{control},
+                               std::vector<bool>{true},
+                               std::vector<std::size_t>{wire0, wire1}, inverse,
+                               {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled SingleExcitation - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire0 && control != wire1 && wire0 != wire1) {
+            auto matrix = getSingleExcitation<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1}, inverse);
+            sv1.applyOperation(
+                "SingleExcitation", std::vector<std::size_t>{control},
+                std::vector<bool>{true}, std::vector<std::size_t>{wire0, wire1},
+                inverse, {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled SingleExcitationMinus - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire0 && control != wire1 && wire0 != wire1) {
+            auto matrix =
+                getSingleExcitationMinus<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1}, inverse);
+            sv1.applyOperation(
+                "SingleExcitationMinus", std::vector<std::size_t>{control},
+                std::vector<bool>{true}, std::vector<std::size_t>{wire0, wire1},
+                inverse, {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled SingleExcitationPlus - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        if (control != wire0 && control != wire1 && wire0 != wire1) {
+            auto matrix =
+                getSingleExcitationPlus<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1}, inverse);
+            sv1.applyOperation(
+                "SingleExcitationPlus", std::vector<std::size_t>{control},
+                std::vector<bool>{true}, std::vector<std::size_t>{wire0, wire1},
+                inverse, {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+}
+
+TEMPLATE_TEST_CASE(
+    "StateVectorCudaManaged::applyOperation param four-qubits with controls",
+    "[StateVectorCudaManaged]", float, double) {
+    using PrecisionT = TestType;
+    using ComplexT = std::complex<TestType>;
+    std::mt19937 re{1337};
+    const int num_qubits = 5;
+    const auto margin = PrecisionT{1e-5};
+    const std::size_t control = GENERATE(0, 1, 2, 3, 4);
+    const std::size_t wire0 = GENERATE(0, 1, 2, 3, 4);
+    const std::size_t wire1 = GENERATE(0, 1, 2, 3, 4);
+    const std::size_t wire2 = GENERATE(0, 1, 2, 3, 4);
+    const std::size_t wire3 = GENERATE(0, 1, 2, 3, 4);
+    bool inverse = GENERATE(false, true);
+    PrecisionT param = GENERATE(-0.5, 1.5);
+
+    StateVectorCudaManaged<PrecisionT> sv0(num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv1(num_qubits);
+
+    auto st0 = createRandomStateVectorData<PrecisionT>(re, num_qubits);
+    sv0.CopyHostDataToGpu(st0.data(), st0.size());
+    sv1.CopyHostDataToGpu(st0.data(), st0.size());
+
+    auto getControlledGate = [](std::vector<ComplexT> matrix) {
+        std::vector<ComplexT> cmatrix(matrix.size() * 4);
+        for (std::size_t i = 0; i < 16; i++) {
+            cmatrix[i * 32 + i] = ComplexT{1.0};
+        }
+        for (std::size_t i = 0; i < 16; i++) {
+            for (std::size_t j = 0; j < 16; j++) {
+                cmatrix[(i + 16) * 32 + j + 16] = matrix[i * 16 + j];
+            }
+        }
+        return cmatrix;
+    };
+
+    DYNAMIC_SECTION("N-controlled DoubleExcitation - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << ", " << wire2
+                    << ", " << wire3 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        std::vector<std::size_t> wires = {control, wire0, wire1, wire2, wire3};
+        std::sort(wires.begin(), wires.end());
+        if (std::adjacent_find(wires.begin(), wires.end()) == wires.end()) {
+            auto matrix = getDoubleExcitation<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1, wire2, wire3},
+                            inverse);
+            sv1.applyOperation(
+                "DoubleExcitation", std::vector<std::size_t>{control},
+                std::vector<bool>{true},
+                std::vector<std::size_t>{wire0, wire1, wire2, wire3}, inverse,
+                {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled DoubleExcitationMinus - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << ", " << wire2
+                    << ", " << wire3 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        std::vector<std::size_t> wires = {control, wire0, wire1, wire2, wire3};
+        std::sort(wires.begin(), wires.end());
+        if (std::adjacent_find(wires.begin(), wires.end()) == wires.end()) {
+            auto matrix =
+                getDoubleExcitationMinus<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1, wire2, wire3},
+                            inverse);
+            sv1.applyOperation(
+                "DoubleExcitationMinus", std::vector<std::size_t>{control},
+                std::vector<bool>{true},
+                std::vector<std::size_t>{wire0, wire1, wire2, wire3}, inverse,
+                {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled DoubleExcitationPlus - "
+                    << "controls = {" << control << "} "
+                    << ", wires = {" << wire0 << ", " << wire1 << ", " << wire2
+                    << ", " << wire3 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        std::vector<std::size_t> wires = {control, wire0, wire1, wire2, wire3};
+        std::sort(wires.begin(), wires.end());
+        if (std::adjacent_find(wires.begin(), wires.end()) == wires.end()) {
+            auto matrix =
+                getDoubleExcitationPlus<std::complex, PrecisionT>(param);
+            std::vector<ComplexT> cmatrix = getControlledGate(matrix);
+
+            sv0.applyMatrix(cmatrix, {control, wire0, wire1, wire2, wire3},
+                            inverse);
+            sv1.applyOperation(
+                "DoubleExcitationPlus", std::vector<std::size_t>{control},
+                std::vector<bool>{true},
+                std::vector<std::size_t>{wire0, wire1, wire2, wire3}, inverse,
+                {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+
+    DYNAMIC_SECTION("N-controlled MultiRZ - "
+                    << "controls = {" << control << ", " << wire0 << ", "
+                    << wire1 << "} "
+                    << ", wires = {" << wire2 << ", " << wire3 << "} - "
+                    << PrecisionToName<PrecisionT>::value) {
+        std::vector<std::size_t> wires = {control, wire0, wire1, wire2, wire3};
+        std::sort(wires.begin(), wires.end());
+        const ComplexT e = std::exp(ComplexT{0, -0.5} * param);
+        std::vector<ComplexT> matrix(16, 0.0);
+        matrix[0] = e;
+        matrix[5] = std::conj(e);
+        matrix[10] = std::conj(e);
+        matrix[15] = e;
+        if (std::adjacent_find(wires.begin(), wires.end()) == wires.end()) {
+            sv0.applyControlledMatrix(
+                matrix.data(), matrix.size(), {control, wire0, wire1},
+                std::vector<bool>{true, false, true}, {wire2, wire3}, inverse);
+            sv1.applyOperation(
+                "MultiRZ", std::vector<std::size_t>{control, wire0, wire1},
+                std::vector<bool>{true, false, true},
+                std::vector<std::size_t>{wire2, wire3}, inverse, {param});
+            REQUIRE(sv0.getDataVector() ==
+                    approx(sv1.getDataVector()).margin(margin));
+        }
+    }
+}
+
 TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyControlledGlobalPhase",
                    "[StateVectorCudaManaged_Param]", double) {
     using ComplexT = StateVectorCudaManaged<TestType>::ComplexT;

From c035b51e40d8c1d1e6116291f40d34bbe10e85f6 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Mon, 21 Oct 2024 13:26:47 +0000
Subject: [PATCH 23/66] update the frontend

---
 pennylane_lightning/lightning_gpu/_state_vector.py | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/pennylane_lightning/lightning_gpu/_state_vector.py b/pennylane_lightning/lightning_gpu/_state_vector.py
index 93acf859b..b39e0895c 100644
--- a/pennylane_lightning/lightning_gpu/_state_vector.py
+++ b/pennylane_lightning/lightning_gpu/_state_vector.py
@@ -329,7 +329,9 @@ def _apply_lightning(
             elif method is not None:  # apply specialized gate
                 param = operation.parameters
                 method(wires, invert_param, param)
-            elif isinstance(operation, qml.ops.Controlled):  # apply n-controlled gate
+            elif (
+                isinstance(operation, qml.ops.Controlled) and not self._mpi_handler.use_mpi
+            ):  # MPI backend does not have native controlled gates support
                 self._apply_lightning_controlled(operation)
             else:  # apply gate as a matrix
                 try:

From faad0e969311275cc88d20201c0fccad9fa2dca5 Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Mon, 21 Oct 2024 13:27:09 +0000
Subject: [PATCH 24/66] Auto update version from '0.39.0-dev47' to
 '0.39.0-dev48'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index de59ceba0..6f841927f 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev47"
+__version__ = "0.39.0-dev48"

From 148b8df1fc57cf4b2aa38c12b1688e5eb4567f01 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Mon, 21 Oct 2024 13:29:01 +0000
Subject: [PATCH 25/66] update changelog

---
 .github/CHANGELOG.md | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/.github/CHANGELOG.md b/.github/CHANGELOG.md
index 1f4059f9e..d1fbb7bde 100644
--- a/.github/CHANGELOG.md
+++ b/.github/CHANGELOG.md
@@ -2,6 +2,9 @@
 
 ### New features since last release
 
+* Add native N-controlled gates support to `lightning.gpu`'s single-GPU backend.
+  [(#938)](https://github.com/PennyLaneAI/pennylane-lightning/pull/938)
+
 * Add `mid-circuit measurements` support to `lightning.gpu`'s single-GPU backend.
   [(#931)](https://github.com/PennyLaneAI/pennylane-lightning/pull/931)
 

From 7ea85d86d4647d74f78d120373f42441743d987d Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Mon, 21 Oct 2024 14:11:14 +0000
Subject: [PATCH 26/66] Auto update version from '0.39.0-dev47' to
 '0.39.0-dev48'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index de59ceba0..6f841927f 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev47"
+__version__ = "0.39.0-dev48"

From 388efae041345991c83050652ee8c202e8196515 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Mon, 21 Oct 2024 14:12:11 +0000
Subject: [PATCH 27/66] Trigger CIs


From fbae9624fc6143fcf13ce6fec4f764cfd54a1e92 Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Mon, 21 Oct 2024 14:56:44 +0000
Subject: [PATCH 28/66] Auto update version from '0.39.0-dev47' to
 '0.39.0-dev48'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index de59ceba0..6f841927f 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev47"
+__version__ = "0.39.0-dev48"

From 7ec070e0cbef1f6db7f59acdf5b57dcae3582f50 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Mon, 21 Oct 2024 17:49:31 +0000
Subject: [PATCH 29/66] update py CIs

---
 .github/workflows/tests_gpu_python.yml | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/.github/workflows/tests_gpu_python.yml b/.github/workflows/tests_gpu_python.yml
index 37a9ca95d..1db5f5a28 100644
--- a/.github/workflows/tests_gpu_python.yml
+++ b/.github/workflows/tests_gpu_python.yml
@@ -176,7 +176,7 @@ jobs:
         run: |
           rm -rf build
           PL_BACKEND=lightning_qubit python scripts/configure_pyproject_toml.py || true
-          PL_BACKEND=lightning_qubit SKIP_COMPILATION=True python -m pip install . -vv
+          PL_BACKEND=lightning_qubit python -m pip install . -vv
 
           rm -rf build
           PL_BACKEND=${{ matrix.pl_backend }} python scripts/configure_pyproject_toml.py || true

From c1a141ae3ffcdb21ccc14007e7c8d5566eaea679 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 22 Oct 2024 13:51:49 +0000
Subject: [PATCH 30/66] add more C++ unit tests

---
 .../Test_StateVectorCudaManaged_NonParam.cpp  | 40 +++++++++++++++++++
 1 file changed, 40 insertions(+)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
index 45212573e..ef6d0f1c9 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
@@ -1312,3 +1312,43 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::controlled Toffoli",
                        std::vector<std::size_t>{5});
     REQUIRE(sv0.getDataVector() == approx(sv1.getDataVector()).margin(margin));
 }
+
+TEMPLATE_TEST_CASE(
+    "StateVectorCudaManaged::controlled_gate_direct_matrix_offload",
+    "[StateVectorCudaManaged]", float, double) {
+    using PrecisionT = TestType;
+    std::mt19937 re{1337};
+    const int num_qubits = 6;
+    const auto margin = PrecisionT{1e-5};
+    const std::size_t control = GENERATE(0, 1, 2);
+    auto st0 = createRandomStateVectorData<PrecisionT>(re, num_qubits);
+
+    StateVectorCudaManaged<PrecisionT> sv0(num_qubits);
+    StateVectorCudaManaged<PrecisionT> sv1(num_qubits);
+
+    sv0.CopyHostDataToGpu(st0.data(), st0.size());
+    sv1.CopyHostDataToGpu(st0.data(), st0.size());
+
+    SECTION("Catch failures caused by unsupported named gates") {
+        PL_CHECK_THROWS_MATCHES(
+            sv0.applyOperation("paulix",
+                               std::vector<std::size_t>{control, 3, 4},
+                               std::vector<bool>{true, true, true},
+                               std::vector<std::size_t>{5}),
+            LightningException, "Currently unsupported gate: paulix");
+    }
+
+    SECTION("direct base matrix offload") {
+        const std::vector<std::complex<PrecisionT>> matrix = {
+            {0.0, 0.0}, {1.0, 0.0}, {1.0, 0.0}, {0.0, 0.0}};
+        sv0.applyControlledMatrix(matrix.data(), matrix.size(),
+                                  std::vector<std::size_t>{control, 3, 4},
+                                  std::vector<bool>{true, true, true},
+                                  std::vector<std::size_t>{5});
+        sv1.applyOperation("Paulix", std::vector<std::size_t>{control, 3, 4},
+                           std::vector<bool>{true, true, true},
+                           std::vector<std::size_t>{5}, false, {0.0}, matrix);
+        REQUIRE(sv0.getDataVector() ==
+                approx(sv1.getDataVector()).margin(margin));
+    }
+}

From b54b8ca27932b75d4cecb1d63acaae4c69404b7f Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 22 Oct 2024 14:09:21 +0000
Subject: [PATCH 31/66] add checks for ctrls and tgts

---
 .../src/simulators/lightning_gpu/StateVectorCudaManaged.hpp    | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index 330f7c67b..1c845ff51 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -368,6 +368,9 @@ class StateVectorCudaManaged
                         bool adjoint = false,
                         const std::vector<Precision> &params = {0.0},
                         const std::vector<ComplexT> &gate_matrix = {}) {
+        PL_ABORT_IF_NOT(
+            areVecsDisjoint<std::size_t>(ctrls, tgts),
+            "`controlled_wires` and `target wires` must be disjoint.");
         PL_ABORT_IF(ctrls.size() != ctrls_values.size(),
                     "`ctrls` and `ctrls_values` must have the same size.");
         std::vector<int> ctrlsInt(ctrls.size());

From b7ec1daf5a9f4468b033134ec2c180f8d7e1ef82 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 22 Oct 2024 14:53:28 +0000
Subject: [PATCH 32/66] tidy up naming

---
 .../lightning_gpu/StateVectorCudaManaged.hpp  | 25 ++++++++++---------
 .../lightning_gpu/_state_vector.py            |  7 +++---
 2 files changed, 16 insertions(+), 16 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index 2e31bbe2f..2612e8766 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -361,21 +361,22 @@ class StateVectorCudaManaged
      * @param params Optional std gate matrix if opName doesn't exist.
      */
     void applyOperation(
-        const std::string &opName, const std::vector<std::size_t> &ctrl_wires,
-        const std::vector<bool> &ctrl_values,
+        const std::string &opName,
+        const std::vector<std::size_t> &controlled_wires,
+        const std::vector<bool> &controlled_values,
         const std::vector<std::size_t> &tgt_wires, bool adjoint = false,
         const std::vector<Precision> &params = {0.0},
         [[maybe_unused]] const std::vector<ComplexT> &gate_matrix = {}) {
         PL_ABORT_IF_NOT(opName == "GlobalPhase",
                         "Only GlobalPhase gate is supported.");
-        PL_ABORT_IF(ctrl_wires.size() != ctrl_values.size(),
+        PL_ABORT_IF(controlled_wires.size() != controlled_values.size(),
                     "`ctrls` and `ctrls_values` must have the same size.");
-        std::vector<int> ctrlsInt(ctrl_wires.size());
+        std::vector<int> ctrlsInt(controlled_wires.size());
         std::vector<int> tgtsInt(tgt_wires.size());
-        std::vector<int> ctrls_valuesInt(ctrl_wires.size());
+        std::vector<int> ctrls_valuesInt(controlled_wires.size());
 
-        std::transform(ctrl_wires.begin(), ctrl_wires.end(), ctrlsInt.begin(),
-                       [&](std::size_t x) {
+        std::transform(controlled_wires.begin(), controlled_wires.end(),
+                       ctrlsInt.begin(), [&](std::size_t x) {
                            return static_cast<int>(BaseType::getNumQubits() -
                                                    1 - x);
                        });
@@ -385,7 +386,7 @@ class StateVectorCudaManaged
                                                    1 - x);
                        });
 
-        std::transform(ctrl_values.begin(), ctrl_values.end(),
+        std::transform(controlled_values.begin(), controlled_values.end(),
                        ctrls_valuesInt.begin(),
                        [&](bool x) { return static_cast<int>(x); });
         if (opName == "GlobalPhase") {
@@ -1425,10 +1426,10 @@ class StateVectorCudaManaged
     /**
      * @brief Apply parametric Pauli gates using custateVec calls.
      *
-     * @param angle Rotation angle.
      * @param pauli_words List of Pauli words representing operation.
      * @param ctrls Control wires
      * @param tgts target wires.
+     * @param params Rotation parameters.
      * @param use_adjoint Take adjoint of operation.
      */
     void applyParametricPauliGate_(const std::vector<std::string> &pauli_words,
@@ -1458,9 +1459,9 @@ class StateVectorCudaManaged
      * @brief Apply a parametric Pauli gate using custateVec calls.
      *
      * @param pauli_words List of Pauli words representing operation.
-     * @param ctrls Control wires
-     * @param ctrls_values Control values
-     * @param tgts target wires.
+     * @param ctrlsInt Control wires
+     * @param ctrls_valuesInt Control values
+     * @param tgtsInt target wires.
      * @param param Rotation angle.
      * @param use_adjoint Take adjoint of operation.
      */
diff --git a/pennylane_lightning/lightning_gpu/_state_vector.py b/pennylane_lightning/lightning_gpu/_state_vector.py
index c4bce5db3..99b9a91ef 100644
--- a/pennylane_lightning/lightning_gpu/_state_vector.py
+++ b/pennylane_lightning/lightning_gpu/_state_vector.py
@@ -238,6 +238,9 @@ def _apply_lightning_controlled(self, operation):
         Returns:
             None
         """
+        if self._mpi_handler.use_mpi:
+            raise DeviceError("Lightning-GPU-MPI does not support Controlled GlobalPhase gates.")
+
         state = self.state_vector
 
         basename = operation.base.name
@@ -323,10 +326,6 @@ def _apply_lightning(
                 operation.base, qml.GlobalPhase
             ):  # apply n-controlled gate
                 # LGPU do not support the controlled gates except for GlobalPhase
-                if self._mpi_handler.use_mpi:
-                    raise DeviceError(
-                        "Lightning-GPU-MPI does not support Controlled GlobalPhase gates."
-                    )
                 self._apply_lightning_controlled(operation)
             else:  # apply gate as a matrix
                 try:

From 4e69bb3ffceea54b4fda7ffdbcca2fe9b392ed86 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 22 Oct 2024 17:37:59 +0000
Subject: [PATCH 33/66] update

---
 .../lightning_gpu/tests/Test_StateVectorCudaManaged.cpp         | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
index c14afb4e3..519ca84ba 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
@@ -193,7 +193,7 @@ TEMPLATE_PRODUCT_TEST_CASE("StateVectorCudaManaged::applyOperations",
                                         std::vector<std::size_t>{1}, false,
                                         {0.0}, std::vector<ComplexT>{}),
             LightningException,
-            "`ctrls` and `ctrls_values` must have the "
+            "`ctrls` and `controlled_values` must have the "
             "same size."); // invalid controlled_wires
     }
 

From 454e4e004072c569f2d7e78ab7d92dc1704c3f05 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 22 Oct 2024 18:47:19 +0000
Subject: [PATCH 34/66] quick fix

---
 pennylane_lightning/lightning_gpu/_state_vector.py | 11 ++++++++---
 1 file changed, 8 insertions(+), 3 deletions(-)

diff --git a/pennylane_lightning/lightning_gpu/_state_vector.py b/pennylane_lightning/lightning_gpu/_state_vector.py
index 397e2c503..4beccdcc9 100644
--- a/pennylane_lightning/lightning_gpu/_state_vector.py
+++ b/pennylane_lightning/lightning_gpu/_state_vector.py
@@ -238,9 +238,6 @@ def _apply_lightning_controlled(self, operation):
         Returns:
             None
         """
-        if self._mpi_handler.use_mpi:
-            raise DeviceError("Lightning-GPU-MPI does not support Controlled GlobalPhase gates.")
-
         state = self.state_vector
 
         basename = operation.base.name
@@ -336,6 +333,14 @@ def _apply_lightning(
                 isinstance(operation, qml.ops.Controlled) and not self._mpi_handler.use_mpi
             ):  # MPI backend does not have native controlled gates support
                 self._apply_lightning_controlled(operation)
+            elif (
+                self._mpi_handler.use_mpi
+                and isinstance(operation, qml.ops.Controlled)
+                and isinstance(operation.base, qml.GlobalPhase)
+            ):
+                raise DeviceError(
+                    "Lightning-GPU-MPI does not support Controlled GlobalPhase gates."
+                )
             else:  # apply gate as a matrix
                 try:
                     mat = qml.matrix(operation)

From 9b566858d7f8ab314e31fdc8c9f1984fda054ed7 Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Wed, 23 Oct 2024 02:19:15 +0000
Subject: [PATCH 35/66] Auto update version from '0.39.0-dev48' to
 '0.39.0-dev49'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index 6f841927f..eba0279f9 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev48"
+__version__ = "0.39.0-dev49"

From b0300d62aa9f0c406901b57a9413ef0db3d0799d Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Wed, 23 Oct 2024 03:15:19 +0000
Subject: [PATCH 36/66] update

---
 .../lightning_gpu/StateVectorCudaManaged.hpp  | 22 ++++++++-----------
 1 file changed, 9 insertions(+), 13 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index ed464e826..e594719a5 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -472,6 +472,10 @@ class StateVectorCudaManaged
         auto ctrls_valuesInt =
             Pennylane::Util::cast_vector<bool, int>(controlled_values);
 
+        std::reverse(ctrlsInt.begin(), ctrlsInt.end());
+        std::reverse(tgtsInt.begin(), tgtsInt.end());
+        std::reverse(ctrls_valuesInt.begin(), ctrls_valuesInt.end());
+
         applyDeviceGeneralGate_(d_matrix.getData(), ctrlsInt, tgtsInt,
                                 ctrls_valuesInt, inverse);
     }
@@ -1591,19 +1595,11 @@ class StateVectorCudaManaged
                                 const std::vector<std::size_t> &ctrls,
                                 const std::vector<std::size_t> &tgts,
                                 bool use_adjoint = false) {
-        std::vector<int> ctrlsInt(ctrls.size());
-        std::vector<int> tgtsInt(tgts.size());
-
-        std::transform(
-            ctrls.begin(), ctrls.end(), ctrlsInt.begin(), [&](std::size_t x) {
-                return static_cast<int>(BaseType::getNumQubits() - 1 - x);
-            });
-        std::transform(
-            tgts.begin(), tgts.end(), tgtsInt.begin(), [&](std::size_t x) {
-                return static_cast<int>(BaseType::getNumQubits() - 1 - x);
-            });
-
-        const std::vector<int> ctrls_values(ctrls.size(), 1);
+        auto ctrlsInt = NormalizeCastIndices<std::size_t, int>(
+            ctrls, BaseType::getNumQubits());
+        auto tgtsInt = NormalizeCastIndices<std::size_t, int>(
+            tgts, BaseType::getNumQubits());
+        auto ctrls_values = std::vector<int>(ctrls.size(), 1);
 
         applyDeviceGeneralGate_(matrix, ctrlsInt, tgtsInt, ctrls_values,
                                 use_adjoint);

From 7e4a34b990391177b486ebe68dcc152d406acf3e Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Wed, 23 Oct 2024 03:31:26 +0000
Subject: [PATCH 37/66] quick fix

---
 .../src/simulators/lightning_gpu/StateVectorCudaManaged.hpp   | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index e594719a5..c4c2d27a2 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -407,6 +407,7 @@ class StateVectorCudaManaged
                                         ctrls_valuesInt, false);
             }
         } else if (par_gates_.find(opName) != par_gates_.end()) {
+            //TODO: offload to par_gates_ if available
             auto &gateMap =
                 cuGates::DynamicGateDataAccess<PrecisionT>::getInstance();
             auto &&matrix_cu = gateMap.getGateData(opName, params);
@@ -433,6 +434,9 @@ class StateVectorCudaManaged
 
                 gate_cache_.add_gate(opName, par[0], matrix_cu);
             }
+            std::reverse(ctrlsInt.begin(), ctrlsInt.end());
+            std::reverse(tgtsInt.begin(), tgtsInt.end());
+            std::reverse(ctrls_valuesInt.begin(), ctrls_valuesInt.end());
             applyDeviceGeneralGate_(
                 gate_cache_.get_gate_device_ptr(opName, par[0]), ctrlsInt,
                 tgtsInt, ctrls_valuesInt, adjoint);

From 8c07c01e313d2b479d3504f92786d2aca4652db8 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Wed, 23 Oct 2024 03:50:24 +0000
Subject: [PATCH 38/66] make format

---
 .../src/simulators/lightning_gpu/StateVectorCudaManaged.hpp     | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index c4c2d27a2..d0d0b1e5a 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -407,7 +407,7 @@ class StateVectorCudaManaged
                                         ctrls_valuesInt, false);
             }
         } else if (par_gates_.find(opName) != par_gates_.end()) {
-            //TODO: offload to par_gates_ if available
+            // TODO: offload to par_gates_ if available
             auto &gateMap =
                 cuGates::DynamicGateDataAccess<PrecisionT>::getInstance();
             auto &&matrix_cu = gateMap.getGateData(opName, params);

From e7b2ec735ce925fb14580357e39c1365e2e5be3a Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Wed, 23 Oct 2024 13:18:24 +0000
Subject: [PATCH 39/66] update lgpu and ltensor installation instructions

---
 README.rst | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/README.rst b/README.rst
index a762da889..de8b24aa0 100644
--- a/README.rst
+++ b/README.rst
@@ -220,10 +220,11 @@ To install Lightning-GPU from the package sources using the direct SDK path, Lig
     git clone https://github.com/PennyLaneAI/pennylane-lightning.git
     cd pennylane-lightning
     pip install -r requirements.txt
+    pip install custatevec-cu12
     PL_BACKEND="lightning_qubit" python scripts/configure_pyproject_toml.py
     SKIP_COMPILATION=True pip install -e . --config-settings editable_mode=compat -vv
 
-Then the `custatevec-cu12 <https://pypi.org/project/cuquantum-cu12/>`_ library can be installed and set a ``CUQUANTUM_SDK`` environment variable.
+Then a ``CUQUANTUM_SDK`` environment variable can be set.
 
 .. code-block:: bash
 
@@ -391,10 +392,11 @@ Lightning-Qubit should be installed before Lightning-Tensor (compilation is not
     git clone https://github.com/PennyLaneAI/pennylane-lightning.git
     cd pennylane-lightning
     pip install -r requirements.txt
+    pip install cutensornet-cu12
     PL_BACKEND="lightning_qubit" python scripts/configure_pyproject_toml.py
     SKIP_COMPILATION=True pip install -e . --config-settings editable_mode=compat
 
-Then the `cutensornet <https://docs.nvidia.com/cuda/cuquantum/latest/cutensornet/index.html>`_ library can be installed and set a ``CUQUANTUM_SDK`` environment variable.
+Then a ``CUQUANTUM_SDK`` environment variable can be set.
 
 .. code-block:: bash
 

From 3e87117ee783efd8e6ff5771f97ea68a480b1ce5 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Wed, 23 Oct 2024 19:19:02 +0000
Subject: [PATCH 40/66] initial commit

---
 .../lightning_gpu/StateVectorCudaMPI.hpp      |  4 +-
 .../lightning_gpu/StateVectorCudaManaged.hpp  |  4 +-
 .../core/src/utils/cuda_utils/CMakeLists.txt  |  2 +-
 .../src/utils/cuda_utils/cuGates_host.hpp     | 51 ++-----------------
 4 files changed, 12 insertions(+), 49 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
index f9c0f0116..ee391e92a 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
@@ -422,7 +422,9 @@ class StateVectorCudaMPI final
             applyDeviceMatrixGate(d_matrix.getData(), ctrls_local, tgts_local,
                                   adjoint);
         } else if (par_gates_.find(opName) != par_gates_.end()) {
-            par_gates_.at(opName)(wires, adjoint, params);
+            const std::vector<std::size_t> wires_local{wires.rbegin(),
+                                                       wires.rend()};
+            par_gates_.at(opName)(wires_local, adjoint, params);
         } else { // No offloadable function call; defer to matrix passing
             auto &&par =
                 (params.empty()) ? std::vector<Precision>{0.0} : params;
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index b01451a42..2a2932ad1 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -327,7 +327,9 @@ class StateVectorCudaManaged
             applyDeviceMatrixGate(d_matrix.getData(), ctrls_local, tgts_local,
                                   adjoint);
         } else if (par_gates_.find(opName) != par_gates_.end()) {
-            par_gates_.at(opName)(wires, adjoint, params);
+            const std::vector<std::size_t> wires_local{wires.rbegin(),
+                                                       wires.rend()};
+            par_gates_.at(opName)(wires_local, adjoint, params);
         } else { // No offloadable function call; defer to matrix passing
             auto &&par =
                 (params.empty()) ? std::vector<Precision>{0.0} : params;
diff --git a/pennylane_lightning/core/src/utils/cuda_utils/CMakeLists.txt b/pennylane_lightning/core/src/utils/cuda_utils/CMakeLists.txt
index 85ebf279e..6de99c572 100644
--- a/pennylane_lightning/core/src/utils/cuda_utils/CMakeLists.txt
+++ b/pennylane_lightning/core/src/utils/cuda_utils/CMakeLists.txt
@@ -5,7 +5,7 @@ findCUDATK(lightning_compile_options)
 
 add_library(cuda_utils INTERFACE)
 target_include_directories(cuda_utils INTERFACE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}>)
-target_link_libraries(cuda_utils INTERFACE CUDA::cudart CUDA::cublas CUDA::cusparse lightning_compile_options)
+target_link_libraries(cuda_utils INTERFACE CUDA::cudart CUDA::cublas CUDA::cusparse lightning_compile_options lightning_gates)
 
 if (BUILD_TESTS)
     enable_testing()
diff --git a/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp b/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp
index 53b712ed5..5097b6935 100644
--- a/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp
+++ b/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp
@@ -654,15 +654,8 @@ template <class CFP_t, class U = double>
 static auto getSingleExcitation(U angle) -> std::vector<CFP_t> {
     const U p2 = angle / 2;
     const CFP_t c{std::cos(p2), 0};
-    // TODO: To remove conditional compilation here in the future, current
-    // implementation will block the simultaneous installation of LGPU and
-    // cutensornet backends
-
-#ifdef _ENABLE_PLGPU
-    const CFP_t s{-std::sin(p2), 0}; // column-major
-#else
     const CFP_t s{std::sin(p2), 0}; // row-major
-#endif
+
     return {cuUtil::ONE<CFP_t>(),
             cuUtil::ZERO<CFP_t>(),
             cuUtil::ZERO<CFP_t>(),
@@ -739,14 +732,7 @@ static auto getSingleExcitationMinus(U angle) -> std::vector<CFP_t> {
     const CFP_t e =
         cuUtil::complexToCu<std::complex<U>>(std::exp(std::complex<U>(0, -p2)));
     const CFP_t c{std::cos(p2), 0};
-// TODO: To remove conditional compilation here in the future, current
-// implementation will block the simultaneous installation of LGPU and
-// cutensornet backends
-#ifdef _ENABLE_PLGPU
-    const CFP_t s{-std::sin(p2), 0}; // column-major
-#else
     const CFP_t s{std::sin(p2), 0}; // row-major
-#endif
 
     return {e,
             cuUtil::ZERO<CFP_t>(),
@@ -826,15 +812,8 @@ static auto getSingleExcitationPlus(U angle) -> std::vector<CFP_t> {
     const CFP_t e =
         cuUtil::complexToCu<std::complex<U>>(std::exp(std::complex<U>(0, p2)));
     const CFP_t c{std::cos(p2), 0};
-    // TODO: To remove conditional compilation here in the future, current
-    // implementation will block the simultaneous installation of LGPU and
-    // cutensornet backends
-
-#ifdef _ENABLE_PLGPU
-    const CFP_t s{-std::sin(p2), 0}; // column-major
-#else
     const CFP_t s{std::sin(p2), 0}; // row-major
-#endif
+
     return {e,
             cuUtil::ZERO<CFP_t>(),
             cuUtil::ZERO<CFP_t>(),
@@ -909,15 +888,8 @@ template <class CFP_t, class U = double>
 static auto getDoubleExcitation(U angle) -> std::vector<CFP_t> {
     const U p2 = angle / 2;
     const CFP_t c{std::cos(p2), 0};
-    // TODO: To remove conditional compilation here in the future, current
-    // implementation will block the simultaneous installation of LGPU and
-    // cutensornet backends
-
-#ifdef _ENABLE_PLGPU
-    const CFP_t s{-std::sin(p2), 0}; // column-major
-#else
     const CFP_t s{std::sin(p2), 0}; // row-major
-#endif
+
     std::vector<CFP_t> mat(256, cuUtil::ZERO<CFP_t>());
     mat[0] = cuUtil::ONE<CFP_t>();
     mat[17] = cuUtil::ONE<CFP_t>();
@@ -989,15 +961,8 @@ static auto getDoubleExcitationMinus(U angle) -> std::vector<CFP_t> {
     const CFP_t e =
         cuUtil::complexToCu<std::complex<U>>(std::exp(std::complex<U>(0, -p2)));
     const CFP_t c{std::cos(p2), 0};
-    // TODO: To remove conditional compilation here in the future, current
-    // implementation will block the simultaneous installation of LGPU and
-    // cutensornet backends
-
-#ifdef _ENABLE_PLGPU
-    const CFP_t s{-std::sin(p2), 0}; // column-major
-#else
     const CFP_t s{std::sin(p2), 0}; // row-major
-#endif
+
     std::vector<CFP_t> mat(256, cuUtil::ZERO<CFP_t>());
     mat[0] = e;
     mat[17] = e;
@@ -1085,14 +1050,8 @@ static auto getDoubleExcitationPlus(U angle) -> std::vector<CFP_t> {
     const CFP_t e =
         cuUtil::complexToCu<std::complex<U>>(std::exp(std::complex<U>(0, p2)));
     const CFP_t c{std::cos(p2), 0};
-    // TODO: To remove conditional compilation here in the future, current
-    // implementation will block the simultaneous installation of LGPU and
-    // cutensornet backends
-#ifdef _ENABLE_PLGPU
-    const CFP_t s{-std::sin(p2), 0}; // column-major
-#else
     const CFP_t s{std::sin(p2), 0}; // row-major
-#endif
+
     std::vector<CFP_t> mat(256, cuUtil::ZERO<CFP_t>());
     mat[0] = e;
     mat[17] = e;

From 42ab62f04c9bf7011d8544bb165d6b2f1369444e Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Wed, 23 Oct 2024 19:19:38 +0000
Subject: [PATCH 41/66] Auto update version from '0.39.0-dev48' to
 '0.39.0-dev49'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index 6f841927f..eba0279f9 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev48"
+__version__ = "0.39.0-dev49"

From f9e6a3e4544fc31a84f3978edadaf30c7389e363 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Wed, 23 Oct 2024 19:19:57 +0000
Subject: [PATCH 42/66] revert some changes

---
 README.rst | 6 ++----
 1 file changed, 2 insertions(+), 4 deletions(-)

diff --git a/README.rst b/README.rst
index de8b24aa0..a762da889 100644
--- a/README.rst
+++ b/README.rst
@@ -220,11 +220,10 @@ To install Lightning-GPU from the package sources using the direct SDK path, Lig
     git clone https://github.com/PennyLaneAI/pennylane-lightning.git
     cd pennylane-lightning
     pip install -r requirements.txt
-    pip install custatevec-cu12
     PL_BACKEND="lightning_qubit" python scripts/configure_pyproject_toml.py
     SKIP_COMPILATION=True pip install -e . --config-settings editable_mode=compat -vv
 
-Then a ``CUQUANTUM_SDK`` environment variable can be set.
+Then the `custatevec-cu12 <https://pypi.org/project/cuquantum-cu12/>`_ library can be installed and set a ``CUQUANTUM_SDK`` environment variable.
 
 .. code-block:: bash
 
@@ -392,11 +391,10 @@ Lightning-Qubit should be installed before Lightning-Tensor (compilation is not
     git clone https://github.com/PennyLaneAI/pennylane-lightning.git
     cd pennylane-lightning
     pip install -r requirements.txt
-    pip install cutensornet-cu12
     PL_BACKEND="lightning_qubit" python scripts/configure_pyproject_toml.py
     SKIP_COMPILATION=True pip install -e . --config-settings editable_mode=compat
 
-Then a ``CUQUANTUM_SDK`` environment variable can be set.
+Then the `cutensornet <https://docs.nvidia.com/cuda/cuquantum/latest/cutensornet/index.html>`_ library can be installed and set a ``CUQUANTUM_SDK`` environment variable.
 
 .. code-block:: bash
 

From a510f3ab9a15a593b439538aadcc7b6b9853b177 Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Fri, 25 Oct 2024 13:44:32 +0000
Subject: [PATCH 43/66] Auto update version from '0.39.0-dev49' to
 '0.39.0-dev50'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index eba0279f9..a9b42bc9f 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev49"
+__version__ = "0.39.0-dev50"

From f2306cc20419864f87165af4f0ee7e37ef70268a Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Fri, 25 Oct 2024 15:39:22 +0000
Subject: [PATCH 44/66] apply Joseph's comments

---
 .../lightning_gpu/StateVectorCudaManaged.hpp  |  9 ++++---
 .../tests/Test_StateVectorCudaManaged.cpp     |  2 +-
 .../lightning_gpu/lightning_gpu.py            | 25 ++++++++++++++++++-
 .../test_adjoint_jacobian_class.py            |  4 +--
 .../test_measurements_class.py                |  4 +--
 tests/test_gates.py                           |  8 +++---
 6 files changed, 39 insertions(+), 13 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index d0d0b1e5a..c643315ba 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -339,7 +339,8 @@ class StateVectorCudaManaged
 
             if (!gate_cache_.gateExists(opName, par[0]) &&
                 gate_matrix.empty()) {
-                std::string message = "Currently unsupported gate: " + opName;
+                std::string message = "Currently unsupported gate: " + opName +
+                                      "and no matrix is provided.";
                 throw LightningException(message);
             } else if (!gate_cache_.gateExists(opName, par[0])) {
                 gate_cache_.add_gate(opName, par[0], gate_matrix);
@@ -373,7 +374,8 @@ class StateVectorCudaManaged
             areVecsDisjoint<std::size_t>(controlled_wires, tgt_wires),
             "`controlled_wires` and `target wires` must be disjoint.");
         PL_ABORT_IF(controlled_wires.size() != controlled_values.size(),
-                    "`ctrls` and `controlled_values` must have the same size.");
+                    "`controlled_wires` and `controlled_values` must have the "
+                    "same size.");
         auto ctrlsInt = NormalizeCastIndices<std::size_t, int>(
             controlled_wires, BaseType::getNumQubits());
         auto tgtsInt = NormalizeCastIndices<std::size_t, int>(
@@ -422,7 +424,8 @@ class StateVectorCudaManaged
                 (params.empty()) ? std::vector<Precision>{0.0} : params;
             if (!gate_cache_.gateExists(opName, par[0]) &&
                 gate_matrix.empty()) {
-                std::string message = "Currently unsupported gate: " + opName;
+                std::string message = "Currently unsupported gate: " + opName +
+                                      " and no matrix provided.";
                 throw LightningException(message);
             } else if (!gate_cache_.gateExists(opName, par[0])) {
                 std::vector<CFP_t> matrix_cu(gate_matrix.size());
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
index 519ca84ba..c4438d0d3 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/tests/Test_StateVectorCudaManaged.cpp
@@ -193,7 +193,7 @@ TEMPLATE_PRODUCT_TEST_CASE("StateVectorCudaManaged::applyOperations",
                                         std::vector<std::size_t>{1}, false,
                                         {0.0}, std::vector<ComplexT>{}),
             LightningException,
-            "`ctrls` and `controlled_values` must have the "
+            "`controlled_wires` and `controlled_values` must have the "
             "same size."); // invalid controlled_wires
     }
 
diff --git a/pennylane_lightning/lightning_gpu/lightning_gpu.py b/pennylane_lightning/lightning_gpu/lightning_gpu.py
index d63f13949..1c02ebb8b 100644
--- a/pennylane_lightning/lightning_gpu/lightning_gpu.py
+++ b/pennylane_lightning/lightning_gpu/lightning_gpu.py
@@ -86,8 +86,31 @@
         "PauliZ",
         "MultiRZ",
         "GlobalPhase",
-        "C(GlobalPhase)",
+        "C(PauliX)",
+        "C(PauliY)",
+        "C(PauliZ)",
+        "C(Hadamard)",
+        "C(S)",
+        "C(T)",
+        "C(PhaseShift)",
+        "C(RX)",
+        "C(RY)",
+        "C(RZ)",
+        "C(Rot)",
+        "C(SWAP)",
+        "C(IsingXX)",
+        "C(IsingXY)",
+        "C(IsingYY)",
+        "C(IsingZZ)",
+        "C(SingleExcitation)",
+        "C(SingleExcitationMinus)",
+        "C(SingleExcitationPlus)",
+        "C(DoubleExcitation)",
+        "C(DoubleExcitationMinus)",
+        "C(DoubleExcitationPlus)",
         "C(MultiRZ)",
+        "C(GlobalPhase)",
+        "C(QubitUnitary)",
         "Hadamard",
         "S",
         "Adjoint(S)",
diff --git a/tests/lightning_qubit/test_adjoint_jacobian_class.py b/tests/lightning_qubit/test_adjoint_jacobian_class.py
index 3bff0e64a..45179ad45 100644
--- a/tests/lightning_qubit/test_adjoint_jacobian_class.py
+++ b/tests/lightning_qubit/test_adjoint_jacobian_class.py
@@ -166,8 +166,8 @@ def test_not_expectation_return_type(self, lightning_sv):
             self.calculate_jacobian(lightning_sv(num_wires=1), tape)
 
     @pytest.mark.skipif(
-        device_name not in ("lightning.qubit", "lightning.gpu"),
-        reason="N-controlled operations only implemented in lightning.qubit.",
+        device_name in ("lightning.kokkos"),
+        reason="N-controlled operations are not implemented in lightning.kokkos.",
     )
     @pytest.mark.parametrize("n_qubits", [1, 2, 3, 4])
     @pytest.mark.parametrize("par", [-np.pi / 7, np.pi / 5, 2 * np.pi / 3])
diff --git a/tests/lightning_qubit/test_measurements_class.py b/tests/lightning_qubit/test_measurements_class.py
index e418936c5..f24090903 100644
--- a/tests/lightning_qubit/test_measurements_class.py
+++ b/tests/lightning_qubit/test_measurements_class.py
@@ -799,8 +799,8 @@ def test_controlled_qubit_gates(self, operation, n_qubits, control_value, tol, l
                     assert np.allclose(result, expected, tol * 10)
 
     @pytest.mark.skipif(
-        device_name not in ("lightning.qubit", "lightning.tensor", "lightning.gpu"),
-        reason="N-controlled operations only implemented in lightning.qubit.",
+        device_name in ("lightning.kokkos"),
+        reason="N-controlled operations are not implemented in lightning.kokkos.",
     )
     def test_controlled_qubit_unitary_from_op(self, tol, lightning_sv):
         n_qubits = 10
diff --git a/tests/test_gates.py b/tests/test_gates.py
index d586d35ac..a19e5b23f 100644
--- a/tests/test_gates.py
+++ b/tests/test_gates.py
@@ -373,7 +373,7 @@ def test_state_prep(n_targets, tol):
 
 @pytest.mark.skipif(
     device_name not in ("lightning.qubit", "lightning.gpu"),
-    reason="N-controlled operations only implemented in lightning.qubit.",
+    reason="N-controlled operations only implemented in lightning.qubit and lightning.gpu.",
 )
 @pytest.mark.parametrize("control_value", [False, True])
 @pytest.mark.parametrize("n_qubits", list(range(2, 8)))
@@ -498,7 +498,7 @@ def circuit():
 
 @pytest.mark.skipif(
     device_name not in ("lightning.qubit", "lightning.gpu"),
-    reason="N-controlled operations only implemented in lightning.qubit.",
+    reason="N-controlled operations only implemented in lightning.qubit and lightning.gpu.",
 )
 def test_controlled_qubit_unitary_from_op(tol):
     n_qubits = 10
@@ -558,8 +558,8 @@ def test_paulirot(n_wires, n_targets, tol):
 
 
 @pytest.mark.skipif(
-    device_name not in ("lightning.qubit", "lightning.tensor", "lightning.gpu"),
-    reason="N-controlled operations only implemented in lightning.qubit.",
+    device_name in ("lightning.kokkos"),
+    reason="N-controlled operations are not implemented in lightning.kokkos.",
 )
 @pytest.mark.parametrize("control_wires", range(4))
 @pytest.mark.parametrize("target_wires", range(4))

From e6fe6380c3e95b0c87948d8140b06e27310602a0 Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Fri, 25 Oct 2024 15:40:58 +0000
Subject: [PATCH 45/66] Auto update version from '0.39.0-dev49' to
 '0.39.0-dev51'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index eba0279f9..e1f82227d 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev49"
+__version__ = "0.39.0-dev51"

From 3da92b288d433bd85f97a34785961d403c8c6775 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Fri, 25 Oct 2024 19:45:26 +0000
Subject: [PATCH 46/66] apply afredo's comments

---
 .../lightning_gpu/StateVectorCudaManaged.hpp    | 17 ++++++-----------
 tests/test_adjoint_jacobian.py                  |  4 ++--
 2 files changed, 8 insertions(+), 13 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index c643315ba..c23be1934 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -397,17 +397,12 @@ class StateVectorCudaManaged
                                              ctrls_valuesInt, tgtsInt,
                                              params.front(), adjoint);
         } else if (opName == "Rot") {
-            if (adjoint) {
-                auto rot_matrix =
-                    cuGates::getRot<CFP_t>(params[2], params[1], params[0]);
-                applyDeviceGeneralGate_(rot_matrix.data(), ctrlsInt, tgtsInt,
-                                        ctrls_valuesInt, true);
-            } else {
-                auto rot_matrix =
-                    cuGates::getRot<CFP_t>(params[0], params[1], params[2]);
-                applyDeviceGeneralGate_(rot_matrix.data(), ctrlsInt, tgtsInt,
-                                        ctrls_valuesInt, false);
-            }
+            auto rot_matrix =
+                adjoint
+                    ? cuGates::getRot<CFP_t>(params[2], params[1], params[0])
+                    : cuGates::getRot<CFP_t>(params[0], params[1], params[2]);
+            applyDeviceGeneralGate_(rot_matrix.data(), ctrlsInt, tgtsInt,
+                                    ctrls_valuesInt, adjoint);
         } else if (par_gates_.find(opName) != par_gates_.end()) {
             // TODO: offload to par_gates_ if available
             auto &gateMap =
diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py
index 8dff699b4..b8a42dd27 100644
--- a/tests/test_adjoint_jacobian.py
+++ b/tests/test_adjoint_jacobian.py
@@ -269,7 +269,7 @@ def test_Rot_gradient(self, stateprep, theta, dev):
 
     @pytest.mark.skipif(
         device_name not in ("lightning.qubit", "lightning.gpu"),
-        reason="N-controlled operations only implemented in lightning.qubit.",
+        reason="N-controlled operations only implemented in lightning.qubit and lightning.gpu.",
     )
     @pytest.mark.parametrize("n_qubits", [1, 2, 3, 4])
     @pytest.mark.parametrize("par", [-np.pi / 7, np.pi / 5, 2 * np.pi / 3])
@@ -818,7 +818,7 @@ def circuit(p):
 
     @pytest.mark.skipif(
         device_name not in ("lightning.qubit", "lightning.gpu"),
-        reason="N-controlled operations only implemented in lightning.qubit.",
+        reason="N-controlled operations only implemented in lightning.qubit and lightning.gpu.",
     )
     @pytest.mark.parametrize(
         "operation",

From c3bd58af6daeea0399b9f5b305de0031bf1bf665 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Fri, 25 Oct 2024 19:50:24 +0000
Subject: [PATCH 47/66] add todos

---
 pennylane_lightning/lightning_gpu/_state_vector.py | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/pennylane_lightning/lightning_gpu/_state_vector.py b/pennylane_lightning/lightning_gpu/_state_vector.py
index 4beccdcc9..71a212ba3 100644
--- a/pennylane_lightning/lightning_gpu/_state_vector.py
+++ b/pennylane_lightning/lightning_gpu/_state_vector.py
@@ -287,7 +287,7 @@ def _apply_lightning_midmeasure(
         if operation.reset and bool(sample):
             self.apply_operations([qml.PauliX(operation.wires)], mid_measurements=mid_measurements)
 
-    # pylint: disable=unused-argument
+    # pylint: disable=unused-argument, too-many-branches
     def _apply_lightning(
         self, operations, mid_measurements: dict = None, postselect_mode: str = None
     ):
@@ -338,6 +338,7 @@ def _apply_lightning(
                 and isinstance(operation, qml.ops.Controlled)
                 and isinstance(operation.base, qml.GlobalPhase)
             ):
+                # TODO: To move this line to the _apply_lightning_controlled method once the MPI backend supports controlled gates natively
                 raise DeviceError(
                     "Lightning-GPU-MPI does not support Controlled GlobalPhase gates."
                 )

From fb8d1027d22d821662e625acf35e129c038e7d73 Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Fri, 25 Oct 2024 19:52:09 +0000
Subject: [PATCH 48/66] Auto update version from '0.39.0-dev50' to
 '0.39.0-dev51'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index a9b42bc9f..e1f82227d 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev50"
+__version__ = "0.39.0-dev51"

From eb7e171300a161b0300f1d24c8fc2cbf3713b3da Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Mon, 28 Oct 2024 17:03:06 +0000
Subject: [PATCH 49/66] update pybind and C++ tests

---
 .../src/simulators/lightning_gpu/bindings/LGPUBindings.hpp | 2 +-
 .../gates/tests/Test_StateVectorCudaManaged_NonParam.cpp   | 7 +++++++
 2 files changed, 8 insertions(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp
index 5392631a9..2a5fb6e78 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp
@@ -84,7 +84,7 @@ void registerControlledGate(PyClass &pyclass) {
             const auto gate_name =
                 std::string(lookup(Constant::controlled_gate_names, gate_op));
             const std::string doc = "Apply the " + gate_name + " gate.";
-            auto func = [gate_name = gate_name](
+            auto func = [gate_name](
                             StateVectorT &sv,
                             const std::vector<std::size_t> &controlled_wires,
                             const std::vector<bool> &controlled_values,
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
index ef6d0f1c9..34581a1ce 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
@@ -1134,6 +1134,13 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyOperation non-param "
                                std::vector<std::size_t>{wire});
             REQUIRE(sv0.getDataVector() ==
                     approx(sv1.getDataVector()).margin(margin));
+        } else {
+            PL_REQUIRE_THROWS_MATCHES(
+                sv0.applyOperation("PauliX", std::vector<std::size_t>{control},
+                                   std::vector<bool>{true},
+                                   std::vector<std::size_t>{wire}),
+                LightningException,
+                "`controlled_wires` and `target wires` must be disjoint.");
         }
 
         if (control != 0 && wire != 0 && control != wire) {

From e3faa625d5a50abbe517b1b048c2104ddf08e778 Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Mon, 28 Oct 2024 17:03:35 +0000
Subject: [PATCH 50/66] Auto update version from '0.39.0-dev51' to
 '0.39.0-dev52'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index e1f82227d..338ef4642 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev51"
+__version__ = "0.39.0-dev52"

From 4e0d219f6ec4883922f742dc021e5bb80fc4422c Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Tue, 29 Oct 2024 13:08:36 +0000
Subject: [PATCH 51/66] Auto update version from '0.39.0-dev51' to
 '0.39.0-dev52'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index e1f82227d..338ef4642 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.39.0-dev51"
+__version__ = "0.39.0-dev52"

From 8b9f972c449ed11ed3b90b5c3be963127d161c33 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 29 Oct 2024 13:11:34 +0000
Subject: [PATCH 52/66] update py unit tests to avoid lq compilations

---
 .github/workflows/tests_gpu_python.yml | 2 +-
 tests/test_adjoint_jacobian.py         | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/.github/workflows/tests_gpu_python.yml b/.github/workflows/tests_gpu_python.yml
index 6df2e4958..241b0726e 100644
--- a/.github/workflows/tests_gpu_python.yml
+++ b/.github/workflows/tests_gpu_python.yml
@@ -183,7 +183,7 @@ jobs:
         run: |
           rm -rf build
           PL_BACKEND=lightning_qubit python scripts/configure_pyproject_toml.py || true
-          PL_BACKEND=lightning_qubit python -m pip install . -vv
+          PL_BACKEND=lightning_qubit SKIP_COMPILATION=True python -m pip install . -vv
 
           rm -rf build
           PL_BACKEND=${{ matrix.pl_backend }} python scripts/configure_pyproject_toml.py || true
diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py
index b8a42dd27..c377859ff 100644
--- a/tests/test_adjoint_jacobian.py
+++ b/tests/test_adjoint_jacobian.py
@@ -848,7 +848,7 @@ def circuit(p):
     def test_controlled_jacobian(self, par, n_qubits, control_value, operation, tol):
         """Test that the jacobian of the controlled gate matches the parameter-shift formula."""
         par = np.array([0.1234, par, 0.5678])
-        dev = qml.device("lightning.qubit", wires=n_qubits)
+        dev = qml.device(device_name, wires=n_qubits)
         dqu = qml.device("default.qubit", wires=n_qubits)
         np.random.seed(1337)
         init_state = np.random.rand(2**n_qubits) + 1.0j * np.random.rand(2**n_qubits)

From bc59f43ad760077f1ce8be0f76e01673b139eb51 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 29 Oct 2024 13:58:06 +0000
Subject: [PATCH 53/66] apply some Ali's suggestions

---
 .../lightning_gpu/StateVectorCudaManaged.hpp  |  2 +-
 .../Test_StateVectorCudaManaged_NonParam.cpp  |  2 +-
 .../lightning_gpu/lightning_gpu.py            |  2 +-
 .../lightning_gpu/lightning_gpu.toml          | 52 +++++++++----------
 4 files changed, 29 insertions(+), 29 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index c23be1934..d66e32c2e 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -372,7 +372,7 @@ class StateVectorCudaManaged
                         const std::vector<ComplexT> &gate_matrix = {}) {
         PL_ABORT_IF_NOT(
             areVecsDisjoint<std::size_t>(controlled_wires, tgt_wires),
-            "`controlled_wires` and `target wires` must be disjoint.");
+            "`controlled_wires` and target wires must be disjoint.");
         PL_ABORT_IF(controlled_wires.size() != controlled_values.size(),
                     "`controlled_wires` and `controlled_values` must have the "
                     "same size.");
diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
index 34581a1ce..42732e8c8 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_NonParam.cpp
@@ -1140,7 +1140,7 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::applyOperation non-param "
                                    std::vector<bool>{true},
                                    std::vector<std::size_t>{wire}),
                 LightningException,
-                "`controlled_wires` and `target wires` must be disjoint.");
+                "`controlled_wires` and target wires must be disjoint.");
         }
 
         if (control != 0 && wire != 0 && control != wire) {
diff --git a/pennylane_lightning/lightning_gpu/lightning_gpu.py b/pennylane_lightning/lightning_gpu/lightning_gpu.py
index 1c02ebb8b..baa1f286a 100644
--- a/pennylane_lightning/lightning_gpu/lightning_gpu.py
+++ b/pennylane_lightning/lightning_gpu/lightning_gpu.py
@@ -110,7 +110,7 @@
         "C(DoubleExcitationPlus)",
         "C(MultiRZ)",
         "C(GlobalPhase)",
-        "C(QubitUnitary)",
+        "ControlledQubitUnitary",
         "Hadamard",
         "S",
         "Adjoint(S)",
diff --git a/pennylane_lightning/lightning_gpu/lightning_gpu.toml b/pennylane_lightning/lightning_gpu/lightning_gpu.toml
index b18470da6..24b79c6d5 100644
--- a/pennylane_lightning/lightning_gpu/lightning_gpu.toml
+++ b/pennylane_lightning/lightning_gpu/lightning_gpu.toml
@@ -6,41 +6,41 @@ schema = 2
 [operators.gates.native]
 
 Identity               = { properties = [ "differentiable" ] }
-PauliX                 = { properties = [ "differentiable" ] }
-PauliY                 = { properties = [ "differentiable" ] }
-PauliZ                 = { properties = [ "differentiable" ] }
-Hadamard               = { properties = [ "differentiable" ] }
-S                      = { properties = [ "differentiable" ] }
-T                      = { properties = [ "differentiable" ] }
-PhaseShift             = { properties = [ "differentiable" ] }
-RX                     = { properties = [ "differentiable" ] }
-RY                     = { properties = [ "differentiable" ] }
-RZ                     = { properties = [ "differentiable" ] }
-Rot                    = { properties = [ "differentiable" ] }
+PauliX                 = { properties = [ "controllable", "differentiable" ] }
+PauliY                 = { properties = [ "controllable", "differentiable" ] }
+PauliZ                 = { properties = [ "controllable", "differentiable" ] }
+Hadamard               = { properties = [ "controllable", "differentiable" ] }
+S                      = { properties = [ "controllable", "differentiable" ] }
+T                      = { properties = [ "controllable", "differentiable" ] }
+PhaseShift             = { properties = [ "controllable", "differentiable" ] }
+RX                     = { properties = [ "controllable", "differentiable" ] }
+RY                     = { properties = [ "controllable", "differentiable" ] }
+RZ                     = { properties = [ "controllable", "differentiable" ] }
+Rot                    = { properties = [ "controllable", "differentiable" ] }
 CNOT                   = { properties = [ "differentiable" ] }
 CY                     = { properties = [ "differentiable" ] }
 CZ                     = { properties = [ "differentiable" ] }
-SWAP                   = { properties = [ "differentiable" ] }
+SWAP                   = { properties = [ "controllable", "differentiable" ] }
 CSWAP                  = { properties = [ "differentiable" ] }
 Toffoli                = { properties = [ "differentiable" ] }
-IsingXX                = { properties = [ "differentiable" ] }
-IsingXY                = { properties = [ "differentiable" ] }
-IsingYY                = { properties = [ "differentiable" ] }
-IsingZZ                = { properties = [ "differentiable" ] }
+IsingXX                = { properties = [ "controllable", "differentiable" ] }
+IsingXY                = { properties = [ "controllable", "differentiable" ] }
+IsingYY                = { properties = [ "controllable", "differentiable" ] }
+IsingZZ                = { properties = [ "controllable", "differentiable" ] }
 ControlledPhaseShift   = { properties = [ "differentiable" ] }
 CRX                    = { properties = [ "differentiable" ] }
 CRY                    = { properties = [ "differentiable" ] }
 CRZ                    = { properties = [ "differentiable" ] }
 CRot                   = { properties = [ "differentiable" ] }
-SingleExcitation       = { properties = [ "differentiable" ] }
-SingleExcitationPlus   = { properties = [ "differentiable" ] }
-SingleExcitationMinus  = { properties = [ "differentiable" ] }
-DoubleExcitation       = { properties = [ "differentiable" ] }
-DoubleExcitationPlus   = { properties = [ "differentiable" ] }
-DoubleExcitationMinus  = { properties = [ "differentiable" ] }
-MultiRZ                = { properties = [ "differentiable" ] }
+SingleExcitation       = { properties = [ "controllable", "differentiable" ] }
+SingleExcitationPlus   = { properties = [ "controllable", "differentiable" ] }
+SingleExcitationMinus  = { properties = [ "controllable", "differentiable" ] }
+DoubleExcitation       = { properties = [ "controllable", "differentiable" ] }
+DoubleExcitationPlus   = { properties = [ "controllable", "differentiable" ] }
+DoubleExcitationMinus  = { properties = [ "controllable", "differentiable" ] }
+MultiRZ                = { properties = [ "controllable", "differentiable" ] }
 QubitUnitary           = { properties = [ "differentiable" ] }
-GlobalPhase            = { properties = [ "differentiable" ] }
+GlobalPhase            = { properties = [ "controllable", "differentiable" ] }
 
 # Operators that should be decomposed according to the algorithm used
 # by PennyLane's device API.
@@ -58,7 +58,7 @@ MultiControlledX       = {}
 [operators.gates.matrix]
 
 BlockEncode            = {properties = [ "controllable" ]}
-GlobalPhase            = {properties = [ "controllable" ]}
+GlobalPhase            = {}
 ControlledQubitUnitary = {}
 ECR                    = {}
 SX                     = {}
@@ -90,7 +90,7 @@ Var                    = {}
 Probs                  = {}
 State                  = { condition = [ "analytic" ] }
 Sample                 = { condition = [ "finiteshots" ] }
-Counts                  = { condition = [ "finiteshots" ] }
+Counts                 = { condition = [ "finiteshots" ] }
 
 [compilation]
 # If the device is compatible with qjit

From fd474f149295d2626095b488c2e24a3ebceacfa0 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 29 Oct 2024 15:06:03 +0000
Subject: [PATCH 54/66] rebase fix

---
 .../src/simulators/lightning_gpu/StateVectorCudaManaged.hpp   | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index f10b035dc..609cdcd66 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -411,6 +411,10 @@ class StateVectorCudaManaged
                 cuGates::DynamicGateDataAccess<PrecisionT>::getInstance();
             auto &&matrix_cu = gateMap.getGateData(opName, params);
 
+            std::reverse(ctrlsInt.begin(), ctrlsInt.end());
+            std::reverse(tgtsInt.begin(), tgtsInt.end());
+            std::reverse(ctrls_valuesInt.begin(), ctrls_valuesInt.end());
+
             gate_cache_.add_gate(opName, params[0], matrix_cu);
 
             applyDeviceGeneralGate_(

From a02a236dd632375c2985a2f2ecff14f2050fbe89 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 29 Oct 2024 15:53:57 +0000
Subject: [PATCH 55/66] move wire reversion to `applyDeviceGeneralGate_`

---
 .../lightning_gpu/StateVectorCudaManaged.hpp  | 49 +++++--------------
 1 file changed, 11 insertions(+), 38 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
index 609cdcd66..d354133be 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp
@@ -314,31 +314,12 @@ class StateVectorCudaManaged
                 applyDeviceMatrixGate_(rot_matrix.data(), ctrls, tgts, false);
             }
         } else if (opName == "Matrix") {
-            DataBuffer<CFP_t, int> d_matrix{
-                gate_matrix.size(), BaseType::getDataBuffer().getDevTag(),
-                true};
-            d_matrix.CopyHostDataToGpu(gate_matrix.data(), d_matrix.getLength(),
-                                       false);
-            // ensure wire indexing correctly preserved for tensor-observables
-            const std::vector<std::size_t> ctrls_local{ctrls.rbegin(),
-                                                       ctrls.rend()};
-            const std::vector<std::size_t> tgts_local{tgts.rbegin(),
-                                                      tgts.rend()};
-            applyDeviceMatrixGate_(d_matrix.getData(), ctrls_local, tgts_local,
-                                   adjoint);
+            applyDeviceMatrixGate_(gate_matrix.data(), ctrls, tgts, adjoint);
         } else if (par_gates_.find(opName) != par_gates_.end()) {
-            const std::vector<std::size_t> wires_local{wires.rbegin(),
-                                                       wires.rend()};
-            par_gates_.at(opName)(wires_local, adjoint, params);
+            par_gates_.at(opName)(wires, adjoint, params);
         } else { // No offloadable function call; defer to matrix passing
             auto &&par =
                 (params.empty()) ? std::vector<Precision>{0.0} : params;
-            // ensure wire indexing correctly preserved for tensor-observables
-            const std::vector<std::size_t> ctrls_local{ctrls.rbegin(),
-                                                       ctrls.rend()};
-            const std::vector<std::size_t> tgts_local{tgts.rbegin(),
-                                                      tgts.rend()};
-
             if (!gate_cache_.gateExists(opName, par[0]) &&
                 gate_matrix.empty()) {
                 std::string message = "Currently unsupported gate: " + opName +
@@ -348,8 +329,8 @@ class StateVectorCudaManaged
                 gate_cache_.add_gate(opName, par[0], gate_matrix);
             }
             applyDeviceMatrixGate_(
-                gate_cache_.get_gate_device_ptr(opName, par[0]), ctrls_local,
-                tgts_local, adjoint);
+                gate_cache_.get_gate_device_ptr(opName, par[0]), ctrls, tgts,
+                adjoint);
         }
     }
 
@@ -411,10 +392,6 @@ class StateVectorCudaManaged
                 cuGates::DynamicGateDataAccess<PrecisionT>::getInstance();
             auto &&matrix_cu = gateMap.getGateData(opName, params);
 
-            std::reverse(ctrlsInt.begin(), ctrlsInt.end());
-            std::reverse(tgtsInt.begin(), tgtsInt.end());
-            std::reverse(ctrls_valuesInt.begin(), ctrls_valuesInt.end());
-
             gate_cache_.add_gate(opName, params[0], matrix_cu);
 
             applyDeviceGeneralGate_(
@@ -438,9 +415,6 @@ class StateVectorCudaManaged
 
                 gate_cache_.add_gate(opName, par[0], matrix_cu);
             }
-            std::reverse(ctrlsInt.begin(), ctrlsInt.end());
-            std::reverse(tgtsInt.begin(), tgtsInt.end());
-            std::reverse(ctrls_valuesInt.begin(), ctrls_valuesInt.end());
             applyDeviceGeneralGate_(
                 gate_cache_.get_gate_device_ptr(opName, par[0]), ctrlsInt,
                 tgtsInt, ctrls_valuesInt, adjoint);
@@ -480,10 +454,6 @@ class StateVectorCudaManaged
         auto ctrls_valuesInt =
             Pennylane::Util::cast_vector<bool, int>(controlled_values);
 
-        std::reverse(ctrlsInt.begin(), ctrlsInt.end());
-        std::reverse(tgtsInt.begin(), tgtsInt.end());
-        std::reverse(ctrls_valuesInt.begin(), ctrls_valuesInt.end());
-
         applyDeviceGeneralGate_(d_matrix.getData(), ctrlsInt, tgtsInt,
                                 ctrls_valuesInt, inverse);
     }
@@ -1626,10 +1596,9 @@ class StateVectorCudaManaged
      * @param ctrls_values Control values.
      * @param use_adjoint Use adjoint of given gate. Defaults to false.
      */
-    void applyDeviceGeneralGate_(const CFP_t *matrix,
-                                 const std::vector<int> &ctrls,
-                                 const std::vector<int> &tgts,
-                                 const std::vector<int> &ctrls_values,
+    void applyDeviceGeneralGate_(const CFP_t *matrix, std::vector<int> &ctrls,
+                                 std::vector<int> &tgts,
+                                 std::vector<int> &ctrls_values,
                                  bool use_adjoint = false) {
         void *extraWorkspace = nullptr;
         std::size_t extraWorkspaceSizeInBytes = 0;
@@ -1647,6 +1616,10 @@ class StateVectorCudaManaged
             compute_type = CUSTATEVEC_COMPUTE_32F;
         }
 
+        std::reverse(tgts.begin(), tgts.end());
+        std::reverse(ctrls.begin(), ctrls.end());
+        std::reverse(ctrls_values.begin(), ctrls_values.end());
+
         // check the size of external workspace
         PL_CUSTATEVEC_IS_SUCCESS(custatevecApplyMatrixGetWorkspaceSize(
             /* custatevecHandle_t */ handle_.get(),

From ce084c211d23da0cd075238ae67facfe8dd1aaa3 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 29 Oct 2024 15:55:10 +0000
Subject: [PATCH 56/66] update c++ unit tests

---
 ...Test_StateVectorCudaManaged_Generators.cpp | 24 +++++++++----------
 1 file changed, 12 insertions(+), 12 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp
index 1129e5a66..6b17b1973 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp
@@ -796,8 +796,8 @@ TEST_CASE("Generators::applyGeneratorSingleExcitation", "[GateGenerators]") {
     std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix{
         // clang-format off
         {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 0.0}, {0.0, -1.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 1.0}, {0.0, 0.0}, {0.0, 0.0},
+        {0.0, 0.0}, {0.0, 0.0}, {0.0, 1.0}, {0.0, 0.0},
+        {0.0, 0.0}, {0.0, -1.0}, {0.0, 0.0}, {0.0, 0.0},
         {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}
         // clang-format on
     };
@@ -878,8 +878,8 @@ TEST_CASE("Generators::applyGeneratorSingleExcitationMinus",
     std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix{
         // clang-format off
         {1.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 0.0}, {0.0,-1.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 1.0}, {0.0, 0.0}, {0.0, 0.0},
+        {0.0, 0.0}, {0.0, 0.0}, {0.0, 1.0}, {0.0, 0.0},
+        {0.0, 0.0}, {0.0, -1.0}, {0.0, 0.0}, {0.0, 0.0},
         {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {1.0, 0.0}
         // clang-format on
     };
@@ -960,8 +960,8 @@ TEST_CASE("Generators::applyGeneratorSingleExcitationPlus",
     std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix{
         // clang-format off
         {-1.0, 0.0},{0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 0.0}, {0.0,-1.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 1.0}, {0.0, 0.0}, {0.0, 0.0},
+        {0.0, 0.0}, {0.0, 0.0}, {0.0, 1.0}, {0.0, 0.0},
+        {0.0, 0.0}, {0.0, -1.0}, {0.0, 0.0}, {0.0, 0.0},
         {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {-1.0, 0.0}
         // clang-format on
     };
@@ -1063,7 +1063,7 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitation_GPU",
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, -1.0},{0, 0},{0, 0},{0, 0},
+        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 1.0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
@@ -1072,7 +1072,7 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitation_GPU",
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 1.0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
+        {0, 0},{0, 0},{0, 0},{0, -1.0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0}
@@ -1172,7 +1172,7 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitationMinus_GPU",
         {1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, -1.0},{0, 0},{0, 0},{0, 0},
+        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 1.0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
@@ -1181,7 +1181,7 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitationMinus_GPU",
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 1.0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
+        {0, 0},{0, 0},{0, 0},{0, -1.0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0}
@@ -1281,7 +1281,7 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitationPlus_GPU",
         {-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, -1.0},{0, 0},{0, 0},{0, 0},
+        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 1.0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
@@ -1290,7 +1290,7 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitationPlus_GPU",
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 1.0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
+        {0, 0},{0, 0},{0, 0},{0, -1.0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},
         {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0}

From c347b27137938edc1f7ba2299e12a36c53195f92 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 29 Oct 2024 16:33:43 +0000
Subject: [PATCH 57/66] update c++ unit tests

---
 ...Test_StateVectorCudaManaged_Generators.cpp | 124 ++++++------------
 1 file changed, 40 insertions(+), 84 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp
index 6b17b1973..d3e7de742 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp
@@ -793,14 +793,10 @@ TEST_CASE("Generators::applyGeneratorControlledPhaseShift",
 }
 
 TEST_CASE("Generators::applyGeneratorSingleExcitation", "[GateGenerators]") {
-    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix{
-        // clang-format off
-        {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 0.0}, {0.0, 1.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, -1.0}, {0.0, 0.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}
-        // clang-format on
-    };
+    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
+        16, {0.0, 0.0});
+    matrix[6] = {0.0, 1.0};
+    matrix[9] = {0.0, -1.0};
     std::mt19937 re{1337U};
 
     for (std::size_t num_qubits = 2; num_qubits <= 5; num_qubits++) {
@@ -875,14 +871,12 @@ TEST_CASE("Generators::applyGeneratorSingleExcitation", "[GateGenerators]") {
 
 TEST_CASE("Generators::applyGeneratorSingleExcitationMinus",
           "[GateGenerators]") {
-    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix{
-        // clang-format off
-        {1.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 0.0}, {0.0, 1.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, -1.0}, {0.0, 0.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {1.0, 0.0}
-        // clang-format on
-    };
+    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
+        16, {0.0, 0.0});
+    matrix[0] = {1.0, 0.0};
+    matrix[6] = {0.0, 1.0};
+    matrix[9] = {0.0, -1.0};
+    matrix[15] = {1.0, 0.0};
     std::mt19937 re{1337U};
 
     for (std::size_t num_qubits = 2; num_qubits <= 5; num_qubits++) {
@@ -957,14 +951,12 @@ TEST_CASE("Generators::applyGeneratorSingleExcitationMinus",
 
 TEST_CASE("Generators::applyGeneratorSingleExcitationPlus",
           "[GateGenerators]") {
-    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix{
-        // clang-format off
-        {-1.0, 0.0},{0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 0.0}, {0.0, 1.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, -1.0}, {0.0, 0.0}, {0.0, 0.0},
-        {0.0, 0.0}, {0.0, 0.0}, {0.0, 0.0}, {-1.0, 0.0}
-        // clang-format on
-    };
+    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
+        16, {0.0, 0.0});
+    matrix[0] = {-1.0, 0.0};
+    matrix[6] = {0.0, 1.0};
+    matrix[9] = {0.0, -1.0};
+    matrix[15] = {-1.0, 0.0};
     std::mt19937 re{1337U};
 
     for (std::size_t num_qubits = 2; num_qubits <= 5; num_qubits++) {
@@ -1058,26 +1050,10 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitation_GPU",
     */
     // clang-format on
 
-    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix{
-        // clang-format off
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 1.0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, -1.0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0}
-        // clang-format on
-    };
+    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
+        256, {0.0, 0.0});
+    matrix[60] = {0.0, 1.0};
+    matrix[195] = {0.0, -1.0};
     std::mt19937 re{1337U};
 
     for (std::size_t num_qubits = 4; num_qubits <= 8; num_qubits++) {
@@ -1167,26 +1143,16 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitation_GPU",
 
 TEST_CASE("Generators::applyGeneratorDoubleExcitationMinus_GPU",
           "[GateGenerators]") {
-    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix{
-        // clang-format off
-        {1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 1.0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, -1.0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{1.0, 0}
-        // clang-format on
-    };
+    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
+        256, {0.0, 0.0});
+    matrix[60] = {0.0, 1.0};
+    matrix[195] = {0.0, -1.0};
+    for (std::size_t i = 0; i < 16; i++) {
+        if (i != 3 && i != 12) {
+            const size_t idx = i * 17;
+            matrix[idx] = {1.0, 0.0};
+        }
+    }
     std::mt19937 re{1337U};
 
     for (std::size_t num_qubits = 4; num_qubits <= 8; num_qubits++) {
@@ -1276,26 +1242,16 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitationMinus_GPU",
 
 TEST_CASE("Generators::applyGeneratorDoubleExcitationPlus_GPU",
           "[GateGenerators]") {
-    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix{
-        // clang-format off
-        {-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 1.0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, -1.0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0},{0, 0},
-        {0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{0, 0},{-1.0, 0}
-        // clang-format on
-    };
+    std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
+        256, {0.0, 0.0});
+    matrix[60] = {0.0, 1.0};
+    matrix[195] = {0.0, -1.0};
+    for (std::size_t i = 0; i < 16; i++) {
+        if (i != 3 && i != 12) {
+            const size_t idx = i * 17;
+            matrix[idx] = {-1.0, 0.0};
+        }
+    }
     std::mt19937 re{1337U};
 
     for (std::size_t num_qubits = 4; num_qubits <= 8; num_qubits++) {

From c03e4efb4ef5d057602e2fa98b79039f4b68343a Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 29 Oct 2024 16:35:30 +0000
Subject: [PATCH 58/66] revert changes in cmake

---
 pennylane_lightning/core/src/utils/cuda_utils/CMakeLists.txt | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/src/utils/cuda_utils/CMakeLists.txt b/pennylane_lightning/core/src/utils/cuda_utils/CMakeLists.txt
index 6de99c572..85ebf279e 100644
--- a/pennylane_lightning/core/src/utils/cuda_utils/CMakeLists.txt
+++ b/pennylane_lightning/core/src/utils/cuda_utils/CMakeLists.txt
@@ -5,7 +5,7 @@ findCUDATK(lightning_compile_options)
 
 add_library(cuda_utils INTERFACE)
 target_include_directories(cuda_utils INTERFACE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}>)
-target_link_libraries(cuda_utils INTERFACE CUDA::cudart CUDA::cublas CUDA::cusparse lightning_compile_options lightning_gates)
+target_link_libraries(cuda_utils INTERFACE CUDA::cudart CUDA::cublas CUDA::cusparse lightning_compile_options)
 
 if (BUILD_TESTS)
     enable_testing()

From 40c41b15f1e1b9809fefdd38189a82eaedf97bab Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 29 Oct 2024 19:28:15 +0000
Subject: [PATCH 59/66] disable adjoint tests for N-controlled gates

---
 pennylane_lightning/lightning_gpu/lightning_gpu.py | 1 -
 tests/test_adjoint_jacobian.py                     | 4 ++--
 2 files changed, 2 insertions(+), 3 deletions(-)

diff --git a/pennylane_lightning/lightning_gpu/lightning_gpu.py b/pennylane_lightning/lightning_gpu/lightning_gpu.py
index baa1f286a..03c65e7b8 100644
--- a/pennylane_lightning/lightning_gpu/lightning_gpu.py
+++ b/pennylane_lightning/lightning_gpu/lightning_gpu.py
@@ -110,7 +110,6 @@
         "C(DoubleExcitationPlus)",
         "C(MultiRZ)",
         "C(GlobalPhase)",
-        "ControlledQubitUnitary",
         "Hadamard",
         "S",
         "Adjoint(S)",
diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py
index c377859ff..3c7d68ba8 100644
--- a/tests/test_adjoint_jacobian.py
+++ b/tests/test_adjoint_jacobian.py
@@ -817,8 +817,8 @@ def circuit(p):
             assert np.allclose(jac_ad, jac_bp, atol=tol, rtol=0)
 
     @pytest.mark.skipif(
-        device_name not in ("lightning.qubit", "lightning.gpu"),
-        reason="N-controlled operations only implemented in lightning.qubit and lightning.gpu.",
+        device_name != "lightning.qubit",
+        reason="N-controlled generator operations only implemented in lightning.qubit",
     )
     @pytest.mark.parametrize(
         "operation",

From 1dd8295e6a6606441fb42ab1b62a0ed7fbaf083b Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Tue, 29 Oct 2024 20:26:03 +0000
Subject: [PATCH 60/66] update MPI backend

---
 .../lightning_gpu/StateVectorCudaMPI.hpp      | 49 ++++++-------------
 1 file changed, 14 insertions(+), 35 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
index ee391e92a..36d169a13 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
@@ -399,41 +399,18 @@ class StateVectorCudaMPI final
             applyParametricPauliGate({opName}, ctrls, tgts, params.front(),
                                      adjoint);
         } else if (opName == "Rot" || opName == "CRot") {
-            if (adjoint) {
-                auto rot_matrix =
-                    cuGates::getRot<CFP_t>(params[2], params[1], params[0]);
-                applyDeviceMatrixGate(rot_matrix.data(), ctrls, tgts, true);
-            } else {
-                auto rot_matrix =
-                    cuGates::getRot<CFP_t>(params[0], params[1], params[2]);
-                applyDeviceMatrixGate(rot_matrix.data(), ctrls, tgts, false);
-            }
+            auto rot_matrix =
+                adjoint
+                    ? cuGates::getRot<CFP_t>(params[2], params[1], params[0])
+                    : cuGates::getRot<CFP_t>(params[0], params[1], params[2]);
+            applyDeviceMatrixGate(rot_matrix.data(), ctrls, tgts, adjoint);
         } else if (opName == "Matrix") {
-            DataBuffer<CFP_t, int> d_matrix{
-                gate_matrix.size(), BaseType::getDataBuffer().getDevTag(),
-                true};
-            d_matrix.CopyHostDataToGpu(gate_matrix.data(), d_matrix.getLength(),
-                                       false);
-            // ensure wire indexing correctly preserved for tensor-observables
-            const std::vector<std::size_t> ctrls_local{ctrls.rbegin(),
-                                                       ctrls.rend()};
-            const std::vector<std::size_t> tgts_local{tgts.rbegin(),
-                                                      tgts.rend()};
-            applyDeviceMatrixGate(d_matrix.getData(), ctrls_local, tgts_local,
-                                  adjoint);
+            applyDeviceMatrixGate(gate_matrix.data(), ctrls, tgts, adjoint);
         } else if (par_gates_.find(opName) != par_gates_.end()) {
-            const std::vector<std::size_t> wires_local{wires.rbegin(),
-                                                       wires.rend()};
-            par_gates_.at(opName)(wires_local, adjoint, params);
+            par_gates_.at(opName)(wires, adjoint, params);
         } else { // No offloadable function call; defer to matrix passing
             auto &&par =
                 (params.empty()) ? std::vector<Precision>{0.0} : params;
-            // ensure wire indexing correctly preserved for tensor-observables
-            const std::vector<std::size_t> ctrls_local{ctrls.rbegin(),
-                                                       ctrls.rend()};
-            const std::vector<std::size_t> tgts_local{tgts.rbegin(),
-                                                      tgts.rend()};
-
             if (!gate_cache_.gateExists(opName, par[0]) &&
                 gate_matrix.empty()) {
                 std::string message = "Currently unsupported gate: " + opName;
@@ -442,8 +419,8 @@ class StateVectorCudaMPI final
                 gate_cache_.add_gate(opName, par[0], gate_matrix);
             }
             applyDeviceMatrixGate(
-                gate_cache_.get_gate_device_ptr(opName, par[0]), ctrls_local,
-                tgts_local, adjoint);
+                gate_cache_.get_gate_device_ptr(opName, par[0]), ctrls, tgts,
+                adjoint);
         }
     }
 
@@ -1828,9 +1805,8 @@ class StateVectorCudaMPI final
      * @param tgts Target qubits.
      * @param use_adjoint Use adjoint of given gate.
      */
-    void applyCuSVDeviceMatrixGate(const CFP_t *matrix,
-                                   const std::vector<int> &ctrls,
-                                   const std::vector<int> &tgts,
+    void applyCuSVDeviceMatrixGate(const CFP_t *matrix, std::vector<int> &ctrls,
+                                   std::vector<int> &tgts,
                                    bool use_adjoint = false) {
         void *extraWorkspace = nullptr;
         std::size_t extraWorkspaceSizeInBytes = 0;
@@ -1847,6 +1823,9 @@ class StateVectorCudaMPI final
             data_type = CUDA_C_32F;
             compute_type = CUSTATEVEC_COMPUTE_32F;
         }
+        
+        std::reverse(tgts.begin(), tgts.end());
+        std::reverse(ctrls.begin(), ctrls.end());
 
         // check the size of external workspace
         PL_CUSTATEVEC_IS_SUCCESS(custatevecApplyMatrixGetWorkspaceSize(

From 498528f8b63cf59ee2b20a93e668d00309dbba5c Mon Sep 17 00:00:00 2001
From: ringo-but-quantum <github-ringo-but-quantum@xanadu.ai>
Date: Thu, 7 Nov 2024 15:25:06 +0000
Subject: [PATCH 61/66] Auto update version from '0.40.0-dev3' to '0.40.0-dev4'

---
 pennylane_lightning/core/_version.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/_version.py b/pennylane_lightning/core/_version.py
index e410f98f4..8432c4f82 100644
--- a/pennylane_lightning/core/_version.py
+++ b/pennylane_lightning/core/_version.py
@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.40.0-dev3"
+__version__ = "0.40.0-dev4"

From 4ff25cd5f9b824e5d7cc1fd523753eeac7dd926a Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 7 Nov 2024 18:45:53 +0000
Subject: [PATCH 62/66] make format

---
 .../core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp    | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
index 36d169a13..6fbf9d8b8 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaMPI.hpp
@@ -1823,7 +1823,7 @@ class StateVectorCudaMPI final
             data_type = CUDA_C_32F;
             compute_type = CUSTATEVEC_COMPUTE_32F;
         }
-        
+
         std::reverse(tgts.begin(), tgts.end());
         std::reverse(ctrls.begin(), ctrls.end());
 

From 6e100c3d7cde6f19848eda00fba049f3dd5f608f Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 7 Nov 2024 20:52:26 +0000
Subject: [PATCH 63/66] update generators for excitation gates

---
 .../src/utils/cuda_utils/cuGates_host.hpp     | 24 +++++++++----------
 1 file changed, 12 insertions(+), 12 deletions(-)

diff --git a/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp b/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp
index 5097b6935..56702d527 100644
--- a/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp
+++ b/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp
@@ -705,9 +705,9 @@ static constexpr auto getGeneratorSingleExcitation() -> std::vector<CFP_t> {
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
-        cuUtil::IMAG<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        -cuUtil::IMAG<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), -cuUtil::IMAG<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(), cuUtil::IMAG<CFP_t>(),
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
@@ -785,9 +785,9 @@ static constexpr auto getGeneratorSingleExcitationMinus()
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
-        cuUtil::IMAG<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        -cuUtil::IMAG<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), -cuUtil::IMAG<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(), cuUtil::IMAG<CFP_t>(),
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
@@ -864,9 +864,9 @@ static constexpr auto getGeneratorSingleExcitationPlus() -> std::vector<CFP_t> {
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
-        cuUtil::IMAG<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        -cuUtil::IMAG<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), -cuUtil::IMAG<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(), cuUtil::IMAG<CFP_t>(),
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
         cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
@@ -939,8 +939,8 @@ static auto getDoubleExcitation(const std::vector<U> &params)
 template <class CFP_t, class U = double>
 static constexpr auto getGeneratorDoubleExcitation() -> std::vector<CFP_t> {
     std::vector<CFP_t> mat(256, cuUtil::ZERO<CFP_t>());
-    mat[60] = cuUtil::IMAG<CFP_t>();
-    mat[195] = -cuUtil::IMAG<CFP_t>();
+    mat[60] = -cuUtil::IMAG<CFP_t>();
+    mat[195] = cuUtil::IMAG<CFP_t>();
     return mat;
 }
 
@@ -1017,7 +1017,7 @@ static constexpr auto getGeneratorDoubleExcitationMinus()
     mat[0] = cuUtil::ONE<CFP_t>();
     mat[17] = cuUtil::ONE<CFP_t>();
     mat[34] = cuUtil::ONE<CFP_t>();
-    mat[60] = cuUtil::IMAG<CFP_t>();
+    mat[60] = -cuUtil::IMAG<CFP_t>();
     mat[68] = cuUtil::ONE<CFP_t>();
     mat[85] = cuUtil::ONE<CFP_t>();
     mat[102] = cuUtil::ONE<CFP_t>();
@@ -1026,7 +1026,7 @@ static constexpr auto getGeneratorDoubleExcitationMinus()
     mat[153] = cuUtil::ONE<CFP_t>();
     mat[170] = cuUtil::ONE<CFP_t>();
     mat[187] = cuUtil::ONE<CFP_t>();
-    mat[195] = -cuUtil::IMAG<CFP_t>();
+    mat[195] = cuUtil::IMAG<CFP_t>();
     mat[221] = cuUtil::ONE<CFP_t>();
     mat[238] = cuUtil::ONE<CFP_t>();
     mat[255] = cuUtil::ONE<CFP_t>();
@@ -1105,7 +1105,7 @@ static constexpr auto getGeneratorDoubleExcitationPlus() -> std::vector<CFP_t> {
     mat[0] = -cuUtil::ONE<CFP_t>();
     mat[17] = -cuUtil::ONE<CFP_t>();
     mat[34] = -cuUtil::ONE<CFP_t>();
-    mat[60] = cuUtil::IMAG<CFP_t>();
+    mat[60] = -cuUtil::IMAG<CFP_t>();
     mat[68] = -cuUtil::ONE<CFP_t>();
     mat[85] = -cuUtil::ONE<CFP_t>();
     mat[102] = -cuUtil::ONE<CFP_t>();
@@ -1114,7 +1114,7 @@ static constexpr auto getGeneratorDoubleExcitationPlus() -> std::vector<CFP_t> {
     mat[153] = -cuUtil::ONE<CFP_t>();
     mat[170] = -cuUtil::ONE<CFP_t>();
     mat[187] = -cuUtil::ONE<CFP_t>();
-    mat[195] = -cuUtil::IMAG<CFP_t>();
+    mat[195] = cuUtil::IMAG<CFP_t>();
     mat[221] = -cuUtil::ONE<CFP_t>();
     mat[238] = -cuUtil::ONE<CFP_t>();
     mat[255] = -cuUtil::ONE<CFP_t>();

From 8cb35b03d3f2fadcceec2c3cb2b1755c9a626d1b Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 7 Nov 2024 21:03:21 +0000
Subject: [PATCH 64/66] add changelog

---
 .github/CHANGELOG.md | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/.github/CHANGELOG.md b/.github/CHANGELOG.md
index 423c10879..a3d437563 100644
--- a/.github/CHANGELOG.md
+++ b/.github/CHANGELOG.md
@@ -9,6 +9,9 @@
 
 ### Improvements
 
+* Unify excitation gates memory layout to row-major for both LGPU and LT.
+  [(#959)](https://github.com/PennyLaneAI/pennylane-lightning/pull/959)
+
 * Update the `lightning.kokkos` CUDA backend for compatibility with Catalyst.
   [(#942)](https://github.com/PennyLaneAI/pennylane-lightning/pull/942)
 

From b8ac6bf441c2727c87b8bcc01425b406f127fd67 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 7 Nov 2024 21:12:43 +0000
Subject: [PATCH 65/66] update generator tests

---
 ...Test_StateVectorCudaManaged_Generators.cpp | 24 +++++++++----------
 1 file changed, 12 insertions(+), 12 deletions(-)

diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp
index d3e7de742..a5aba04eb 100644
--- a/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp
+++ b/pennylane_lightning/core/src/simulators/lightning_gpu/gates/tests/Test_StateVectorCudaManaged_Generators.cpp
@@ -795,8 +795,8 @@ TEST_CASE("Generators::applyGeneratorControlledPhaseShift",
 TEST_CASE("Generators::applyGeneratorSingleExcitation", "[GateGenerators]") {
     std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
         16, {0.0, 0.0});
-    matrix[6] = {0.0, 1.0};
-    matrix[9] = {0.0, -1.0};
+    matrix[6] = {0.0, -1.0};
+    matrix[9] = {0.0, 1.0};
     std::mt19937 re{1337U};
 
     for (std::size_t num_qubits = 2; num_qubits <= 5; num_qubits++) {
@@ -874,8 +874,8 @@ TEST_CASE("Generators::applyGeneratorSingleExcitationMinus",
     std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
         16, {0.0, 0.0});
     matrix[0] = {1.0, 0.0};
-    matrix[6] = {0.0, 1.0};
-    matrix[9] = {0.0, -1.0};
+    matrix[6] = {0.0, -1.0};
+    matrix[9] = {0.0, 1.0};
     matrix[15] = {1.0, 0.0};
     std::mt19937 re{1337U};
 
@@ -954,8 +954,8 @@ TEST_CASE("Generators::applyGeneratorSingleExcitationPlus",
     std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
         16, {0.0, 0.0});
     matrix[0] = {-1.0, 0.0};
-    matrix[6] = {0.0, 1.0};
-    matrix[9] = {0.0, -1.0};
+    matrix[6] = {0.0, -1.0};
+    matrix[9] = {0.0, 1.0};
     matrix[15] = {-1.0, 0.0};
     std::mt19937 re{1337U};
 
@@ -1052,8 +1052,8 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitation_GPU",
 
     std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
         256, {0.0, 0.0});
-    matrix[60] = {0.0, 1.0};
-    matrix[195] = {0.0, -1.0};
+    matrix[60] = {0.0, -1.0};
+    matrix[195] = {0.0, 1.0};
     std::mt19937 re{1337U};
 
     for (std::size_t num_qubits = 4; num_qubits <= 8; num_qubits++) {
@@ -1145,8 +1145,8 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitationMinus_GPU",
           "[GateGenerators]") {
     std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
         256, {0.0, 0.0});
-    matrix[60] = {0.0, 1.0};
-    matrix[195] = {0.0, -1.0};
+    matrix[60] = {0.0, -1.0};
+    matrix[195] = {0.0, 1.0};
     for (std::size_t i = 0; i < 16; i++) {
         if (i != 3 && i != 12) {
             const size_t idx = i * 17;
@@ -1244,8 +1244,8 @@ TEST_CASE("Generators::applyGeneratorDoubleExcitationPlus_GPU",
           "[GateGenerators]") {
     std::vector<typename StateVectorCudaManaged<double>::CFP_t> matrix(
         256, {0.0, 0.0});
-    matrix[60] = {0.0, 1.0};
-    matrix[195] = {0.0, -1.0};
+    matrix[60] = {0.0, -1.0};
+    matrix[195] = {0.0, 1.0};
     for (std::size_t i = 0; i < 16; i++) {
         if (i != 3 && i != 12) {
             const size_t idx = i * 17;

From 832eb8e5ed49fbc7bbc369ddb0284fb2085ae647 Mon Sep 17 00:00:00 2001
From: Shuli Shu <08cnbj@gmail.com>
Date: Thu, 7 Nov 2024 21:16:45 +0000
Subject: [PATCH 66/66] make format

---
 .../src/utils/cuda_utils/cuGates_host.hpp     | 42 +++++++++----------
 1 file changed, 21 insertions(+), 21 deletions(-)

diff --git a/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp b/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp
index 56702d527..4d09555d7 100644
--- a/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp
+++ b/pennylane_lightning/core/src/utils/cuda_utils/cuGates_host.hpp
@@ -701,17 +701,17 @@ static auto getSingleExcitation(const std::vector<U> &params)
 template <class CFP_t, class U = double>
 static constexpr auto getGeneratorSingleExcitation() -> std::vector<CFP_t> {
     return {
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
         -cuUtil::IMAG<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), cuUtil::IMAG<CFP_t>(),
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::IMAG<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
     };
 }
 
@@ -781,17 +781,17 @@ template <class CFP_t, class U = double>
 static constexpr auto getGeneratorSingleExcitationMinus()
     -> std::vector<CFP_t> {
     return {
-        cuUtil::ONE<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        cuUtil::ONE<CFP_t>(),   cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
         -cuUtil::IMAG<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), cuUtil::IMAG<CFP_t>(),
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::IMAG<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
-        cuUtil::ZERO<CFP_t>(), cuUtil::ONE<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ONE<CFP_t>(),
     };
 }
 
@@ -860,17 +860,17 @@ static auto getSingleExcitationPlus(const std::vector<U> &params)
 template <class CFP_t, class U = double>
 static constexpr auto getGeneratorSingleExcitationPlus() -> std::vector<CFP_t> {
     return {
-        -cuUtil::ONE<CFP_t>(), cuUtil::ZERO<CFP_t>(),
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        -cuUtil::ONE<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
         -cuUtil::IMAG<CFP_t>(), cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), cuUtil::IMAG<CFP_t>(),
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::IMAG<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
 
-        cuUtil::ZERO<CFP_t>(), cuUtil::ZERO<CFP_t>(),
-        cuUtil::ZERO<CFP_t>(), -cuUtil::ONE<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  cuUtil::ZERO<CFP_t>(),
+        cuUtil::ZERO<CFP_t>(),  -cuUtil::ONE<CFP_t>(),
     };
 }