Add GlobalPhase to all devices

pennylane_lightning/core/_serialize.py

@@ -394,3 +394,18 @@ def get_wires(operation, single_op):
             controlled_wires,
             controlled_values,
         ), uses_stateprep
+
+
+def global_phase_diagonal(par, wires, controls, control_values):
+    """Returns the diagonal of a C(GlobalPhase) operator."""
+    diag = np.ones(2 ** len(wires), dtype=np.complex128)
+    controls = np.array(controls)
+    control_values = np.array(control_values)
+    ind = np.argsort(controls)
+    controls = controls[ind[-1::-1]]
+    control_values = control_values[ind[-1::-1]]
+    idx = np.arange(2 ** len(wires), dtype=np.int64).reshape([2 for _ in wires])
+    for c, w in zip(control_values, controls):
+        idx = np.take(idx, np.array(int(c)), w)
+    diag[idx.ravel()] = np.exp(-1j * par)
+    return diag

pennylane_lightning/core/_version.py

@@ -16,4 +16,4 @@
    Version number (major.minor.patch[-label])
 """
 
-__version__ = "0.34.0-dev20"
+__version__ = "0.34.0-dev21"

pennylane_lightning/core/src/gates/Constant.hpp

@@ -26,7 +26,8 @@ namespace Pennylane::Gates::Constant {
 /**
  * @brief List of multi-qubit gates
  */
-[[maybe_unused]] constexpr std::array multi_qubit_gates{GateOperation::MultiRZ};
+[[maybe_unused]] constexpr std::array multi_qubit_gates{
+    GateOperation::MultiRZ, GateOperation::GlobalPhase};
 /**
  * @brief List of multi-qubit generators
  */
@@ -78,7 +79,8 @@ using GateView = typename std::pair<GateOperation, std::string_view>;
     GateView{GateOperation::DoubleExcitation, "DoubleExcitation"},
     GateView{GateOperation::DoubleExcitationMinus, "DoubleExcitationMinus"},
     GateView{GateOperation::DoubleExcitationPlus, "DoubleExcitationPlus"},
-    GateView{GateOperation::MultiRZ, "MultiRZ"}};
+    GateView{GateOperation::MultiRZ, "MultiRZ"},
+    GateView{GateOperation::GlobalPhase, "GlobalPhase"}};
 
 using CGateView = typename std::pair<ControlledGateOperation, std::string_view>;
 [[maybe_unused]] constexpr std::array controlled_gate_names = {
@@ -336,6 +338,7 @@ using GateNParams = typename std::pair<GateOperation, size_t>;
     GateNParams{GateOperation::DoubleExcitationPlus, 1},
     GateNParams{GateOperation::CSWAP, 0},
     GateNParams{GateOperation::MultiRZ, 1},
+    GateNParams{GateOperation::GlobalPhase, 1},
 };
 
 /**

pennylane_lightning/core/src/gates/GateOperation.hpp

@@ -65,6 +65,7 @@ enum class GateOperation : uint32_t {
     DoubleExcitationPlus,
     /* Multi-qubit gates */
     MultiRZ,
+    GlobalPhase,
     /* END (placeholder) */
     END
 };

pennylane_lightning/core/src/simulators/lightning_gpu/StateVectorCudaManaged.hpp

@@ -65,6 +65,24 @@
                                const size_t index, bool async,
                                cudaStream_t stream_id);
 
+extern void globalPhaseStateVector_CUDA(cuComplex *sv, size_t num_sv,
+                                        cuComplex phase,
+                                        size_t thread_per_block,
+                                        cudaStream_t stream_id);
+extern void globalPhaseStateVector_CUDA(cuDoubleComplex *sv, size_t num_sv,
+                                        cuDoubleComplex phase,
+                                        size_t thread_per_block,
+                                        cudaStream_t stream_id);
+
+extern void cGlobalPhaseStateVector_CUDA(cuComplex *sv, size_t num_sv,
+                                         cuComplex *phase,
+                                         size_t thread_per_block,
+                                         cudaStream_t stream_id);
+extern void cGlobalPhaseStateVector_CUDA(cuDoubleComplex *sv, size_t num_sv,
+                                         cuDoubleComplex *phase,
+                                         size_t thread_per_block,
+                                         cudaStream_t stream_id);
+
 /**
  * @brief Managed memory CUDA state-vector class using custateVec backed
  * gate-calls.
@@ -195,6 +213,44 @@
                             thread_per_block, stream_id);
     }
 
+    /**
+     * @brief Multiplies the state-vector by a global phase.
+     *
+     * @param adjoint Indicates whether to use adjoint of gate.
+     * @param param Complex phase generator.
+     */
+    template <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 =
+            (adjoint) ? std::exp(std::complex<Precision>{0, param})
+                      : std::exp(std::complex<Precision>{0, -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 <size_t thread_per_block = 256>
+    void cGlobalPhaseStateVector(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(),
+                                     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
@@ -232,10 +288,10 @@
      * @param params Optional parameter list for parametric gates.
      * @param gate_matrix Gate data (in row-major format).
      */
-    void applyOperation(
-        const std::string &opName, const std::vector<size_t> &wires,
-        bool adjoint = false, const std::vector<Precision> &params = {0.0},
-        [[maybe_unused]] const std::vector<CFP_t> &gate_matrix = {}) {
+    void applyOperation(const std::string &opName,
+                        const std::vector<size_t> &wires, bool adjoint = false,
+                        const std::vector<Precision> &params = {0.0},
+                        const std::vector<CFP_t> &gate_matrix = {}) {
         const auto ctrl_offset = (BaseType::getCtrlMap().find(opName) !=
                                   BaseType::getCtrlMap().end())
                                      ? BaseType::getCtrlMap().at(opName)
@@ -246,6 +302,10 @@
                                             wires.end()};
         if (opName == "Identity") {
             return;
+        } else if (opName == "C(GlobalPhase)") {
+            cGlobalPhaseStateVector(gate_matrix);
+        } else if (opName == "GlobalPhase") {
+            globalPhaseStateVector(adjoint, params[0]);
         } else if (native_gates_.find(opName) != native_gates_.end()) {
             applyParametricPauliGate({opName}, ctrls, tgts, params.front(),
                                      adjoint);

pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp

@@ -155,18 +155,23 @@ void registerBackendClassSpecificBindings(PyClass &pyclass) {
             "apply",
             [](StateVectorT &sv, const std::string &str,
                const std::vector<size_t> &wires, bool inv,
-               [[maybe_unused]] const std::vector<std::vector<ParamT>> &params,
-               [[maybe_unused]] const np_arr_c &gate_matrix) {
+               const std::vector<std::vector<ParamT>> &params,
+               const np_arr_c &gate_matrix) {
                 const auto m_buffer = gate_matrix.request();
                 std::vector<CFP_t> matrix_cu;
                 if (m_buffer.size) {
                     const auto m_ptr = static_cast<const CFP_t *>(m_buffer.ptr);
                     matrix_cu =
                         std::vector<CFP_t>{m_ptr, m_ptr + m_buffer.size};
                 }
-
-                sv.applyOperation(str, wires, inv, std::vector<ParamT>{},
-                                  matrix_cu);
+                if (params.empty()) {
+                    sv.applyOperation(str, wires, inv, std::vector<ParamT>{},
+                                      matrix_cu);
+                } else {
+                    PL_ABORT_IF(params.size() != 1,
+                                "params should be a List[List[float]].")
+                    sv.applyOperation(str, wires, inv, params[0], matrix_cu);
+                }
             },
             "Apply operation via the gate matrix");
 }

pennylane_lightning/core/src/simulators/lightning_gpu/initSV.cu

@@ -14,6 +14,16 @@
 #include "cuError.hpp"
 #include <cuComplex.h>
 
+namespace {
+template <typename ComplexT>
+__device__ inline ComplexT ComplexMul(ComplexT a, ComplexT b) {
+    ComplexT c;
+    c.x = a.x * b.x - a.y * b.y;
+    c.y = a.x * b.y + a.y * b.x;
+    return c;
+}
+} // namespace
+
 namespace Pennylane::LightningGPU {
 
 /**
@@ -69,6 +79,7 @@ __global__ void setStateVectorkernel(GPUDataT *sv, index_type num_indices,
         sv[indices[i]] = value[i];
     }
 }
+
 /**
  * @brief The CUDA kernel call wrapper.
  *
@@ -95,6 +106,91 @@ void setStateVector_CUDA_call(GPUDataT *sv, index_type &num_indices,
                                                 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] = ComplexMul(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, size_t thread_per_block,
+                                      cudaStream_t stream_id) {
+    auto dv = std::div(static_cast<long>(num_sv), thread_per_block);
+    size_t num_blocks = dv.quot + (dv.rem == 0 ? 0 : 1);
+    const 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>
+__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) {
+        sv[i] = ComplexMul(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>
+void cGlobalPhaseStateVector_CUDA_call(GPUDataT *sv, index_type num_sv,
+                                       GPUDataT *phase, size_t thread_per_block,
+                                       cudaStream_t stream_id) {
+    auto dv = std::div(static_cast<long>(num_sv), thread_per_block);
+    size_t num_blocks = dv.quot + (dv.rem == 0 ? 0 : 1);
+    const 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>
+        <<<gridSize, blockSize, 0, stream_id>>>(sv, num_sv, phase);
+    PL_CUDA_IS_SUCCESS(cudaGetLastError());
+}
+
 /**
  * @brief CUDA runtime API call wrapper.
  *
@@ -140,4 +236,31 @@ void setBasisState_CUDA(cuDoubleComplex *sv, cuDoubleComplex &value,
     setBasisState_CUDA_call(sv, value, index, async, stream_id);
 }
 
+void globalPhaseStateVector_CUDA(cuComplex *sv, size_t num_sv, cuComplex phase,
+                                 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, size_t num_sv,
+                                 cuDoubleComplex phase, 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, size_t num_sv,
+                                  cuComplex *phase, size_t thread_per_block,
+                                  cudaStream_t stream_id) {
+    cGlobalPhaseStateVector_CUDA_call(sv, num_sv, phase, thread_per_block,
+                                      stream_id);
+}
+void cGlobalPhaseStateVector_CUDA(cuDoubleComplex *sv, size_t num_sv,
+                                  cuDoubleComplex *phase,
+                                  size_t thread_per_block,
+                                  cudaStream_t stream_id) {
+    cGlobalPhaseStateVector_CUDA_call(sv, num_sv, phase, thread_per_block,
+                                      stream_id);
+}
+
 } // namespace Pennylane::LightningGPU