From beaa17aaad0a241a0ed2171da2d2b77a98a920c7 Mon Sep 17 00:00:00 2001 From: Shuli <08cnbj@gmail.com> Date: Fri, 20 Oct 2023 15:44:45 +0000 Subject: [PATCH 01/43] Init commit --- .../core/src/observables/Observables.hpp | 101 ++++++++++++ .../observables/ObservablesGPU.cpp | 3 + .../observables/ObservablesGPU.hpp | 86 ++++++++++ .../observables/ObservablesGPUMPI.cpp | 3 + .../observables/ObservablesGPUMPI.hpp | 156 ++++++++++++++++++ 5 files changed, 349 insertions(+) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index 42227f62e5..58770c7af1 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -414,4 +414,105 @@ class HamiltonianBase : public Observable { } }; +/** + * @brief Sparse representation of SparseHamiltonian + * + * @tparam T Floating-point precision. + */ +template +class SparseHamiltonianBase : public Observable { + public: + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + // cuSparse required index type + using IdxT = + typename std::conditional::value, + int32_t, int64_t>::type; + + protected: + std::vector data_; + std::vector indices_; + std::vector offsets_; + std::vector wires_; + + private: + [[nodiscard]] bool + isEqual(const Observable &other) const override { + const auto &other_cast = + static_cast &>(other); + + if (data_ != other_cast.data_ || indices_ != other_cast.indices_ || + offsets_ != other_cast.offsets_) { + return false; + } + + return true; + } + + public: + /** + * @brief Create a SparseHamiltonianBase from data, indices and offsets in + * CSR format. + * + * @param arg1 Arguments to construct data + * @param arg2 Arguments to construct indices + * @param arg3 Arguments to construct offsets + * @param arg4 Arguments to construct wires + */ + template > + SparseHamiltonianBase(T1 &&arg1, T2 &&arg2, T3 &&arg3, T4 &&arg4) + : data_{std::forward(arg1)}, indices_{std::forward(arg2)}, + offsets_{std::forward(arg3)}, wires_{std::forward(arg4)} { + PL_ASSERT(data_.size() == indices_.size()); + } + + /** + * @brief Convenient wrapper for the constructor as the constructor does not + * convert the std::shared_ptr with a derived class correctly. + * + * This function is useful as std::make_shared does not handle + * brace-enclosed initializer list correctly. + * + * @param arg1 Argument to construct data + * @param arg2 Argument to construct indices + * @param arg3 Argument to construct ofsets + * @param arg4 Argument to construct wires + */ + static auto create(std::initializer_list arg1, + std::initializer_list arg2, + std::initializer_list arg3, + std::initializer_list arg4) + -> std::shared_ptr> { + return std::shared_ptr>( + new SparseHamiltonianBase{ + std::move(arg1), std::move(arg2), std::move(arg3), + std::move(arg4)}); + } + + [[nodiscard]] auto getObsName() const -> std::string override { + using Pennylane::Util::operator<<; + std::ostringstream ss; + ss << "SparseHamiltonian: {\n'data' : \n"; + for (const auto &d : data_) + //Note that for LGPU backend, ComplexT is std::complex as of 0.33 release + //Need to revisit it once we set ComplexT as cuComplex later + ss <<"{" << d.real() <<", "<< d.real() <<"},"<< "\n"; + ss << ",\n'indices' : \n"; + for (const auto &i : indices_) + ss << i <<; + ss << ",\n'offsets' : \n"; + for (const auto &o : offsets_) + ss << o; + ss << "\n}"; + return ss.str(); + } + /** + * @brief Get the wires the observable applies to. + */ + [[nodiscard]] auto getWires() const -> std::vector { + return wires_; + }; +}; + } // namespace Pennylane::Observables \ No newline at end of file diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.cpp index b186caa39e..c4f2ca82d4 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.cpp @@ -28,3 +28,6 @@ template class Observables::TensorProdObs>; template class Observables::Hamiltonian>; template class Observables::Hamiltonian>; + +template class Observables::SparseHamiltonian>; +template class Observables::SparseHamiltonian>; diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp index c92e7dac5d..c4042a938c 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp @@ -209,4 +209,90 @@ class Hamiltonian final : public HamiltonianBase { } }; +/** + * @brief Sparse representation of Hamiltonian + * + */ +template +class SparseHamiltonian final : public SparseHamiltonianBase { + public: + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + // cuSparse required index type + using IdxT = + typename std::conditional::value, + int32_t, int64_t>::type; + + private: + using BaseType = SparseHamiltonianBase; + + public: + /** + * @brief Create a SparseHamiltonian from data, indices and offsets in CSR + * format. + * + * @param data Arguments to construct data + * @param indices Arguments to construct indices + * @param offsets Arguments to construct offsets + * @param wires Arguments to construct wires + */ + template + explicit SparseHamiltonian(T1 &&data, T2 &&indices, T3 &&offsets, + T4 &&wires) + : BaseType{data, indices, offsets, wires} {} + + /** + * @brief Convenient wrapper for the constructor as the constructor does not + * convert the std::shared_ptr with a derived class correctly. + * + * This function is useful as std::make_shared does not handle + * brace-enclosed initializer list correctly. + * + * @param data Argument to construct data + * @param indices Argument to construct indices + * @param offsets Argument to construct ofsets + * @param wires Argument to construct wires + */ + static auto create(std::initializer_list data, + std::initializer_list indices, + std::initializer_list offsets, + std::initializer_list wires) + -> std::shared_ptr> { + return std::shared_ptr>( + new SparseHamiltonian{ + std::move(data), std::move(indices), std::move(offsets), + std::move(wires)}); + } + + /** + * @brief Updates the statevector SV:->SV', where SV' = a*H*SV, and where H + * is a sparse Hamiltonian. + * + */ + void applyInPlace(StateVectorT &sv) const override { + PL_ABORT_IF_NOT(this->wires_.size() == sv.getNumQubits(), + "SparseH wire count does not match state-vector size"); + using CFP_t = typename StateVectorT::CFP_t; + + const std::size_t nIndexBits = sv.getNumQubits(); + const std::size_t length = std::size_t{1} << nIndexBits; + + auto device_id = sv.getDataBuffer().getDevTag().getDeviceID(); + auto stream_id = sv.getDataBuffer().getDevTag().getStreamID(); + + cusparseHandle_t handle = sv.getCusparseHandle(); + + std::unique_ptr> d_sv_prime = + std::make_unique>(length, device_id, stream_id, + true); + + SparseMV_cuSparse( + this->offsets_.data(), this->offsets_.size(), this->indices_.data(), + this->data_.data(), this->data_.size(), sv.getData(), + d_sv_prime->getData(), device_id, stream_id, handle); + + sv.updateData(std::move(d_sv_prime)); + } +}; + } // namespace Pennylane::LightningGPU::Observables diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.cpp index 9b1f776e0f..ae9ac9100a 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.cpp @@ -28,3 +28,6 @@ template class Observables::TensorProdObsMPI>; template class Observables::HamiltonianMPI>; template class Observables::HamiltonianMPI>; + +template class Observables::SparseHamiltonianMPI>; +template class Observables::SparseHamiltonianMPI>; diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp index 28c8a88287..ed349cf233 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp @@ -214,4 +214,160 @@ class HamiltonianMPI final : public HamiltonianBase { } }; +/** + * @brief Sparse representation of Hamiltonian + * + */ +template +class SparseHamiltonianMPI final : public SparseHamiltonianBase { + public: + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + // cuSparse required index type + using IdxT = + typename std::conditional::value, + int32_t, int64_t>::type; + + private: + using BaseType = SparseHamiltonianBase; + + public: + /** + * @brief Create a SparseHamiltonianMPI from data, indices and offsets in + * CSR format. + * + * @param data Arguments to construct data + * @param indices Arguments to construct indices + * @param offsets Arguments to construct offsets + * @param wires Arguments to construct wires + */ + template + explicit SparseHamiltonianMPI(T1 &&data, T2 &&indices, T3 &&offsets, + T4 &&wires) + : BaseType{data, indices, offsets, wires} {} + + /** + * @brief Convenient wrapper for the constructor as the constructor does not + * convert the std::shared_ptr with a derived class correctly. + * + * This function is useful as std::make_shared does not handle + * brace-enclosed initializer list correctly. + * + * @param data Argument to construct data + * @param indices Argument to construct indices + * @param offsets Argument to construct ofsets + * @param wires Argument to construct wires + */ + static auto create(std::initializer_list data, + std::initializer_list indices, + std::initializer_list offsets, + std::initializer_list wires) + -> std::shared_ptr> { + return std::shared_ptr>( + new SparseHamiltonianMPI{ + std::move(data), std::move(indices), std::move(offsets), + std::move(wires)}); + } + + /** + * @brief Updates the statevector SV:->SV', where SV' = a*H*SV, and where H + * is a sparse Hamiltonian. + * + */ + void applyInPlace(StateVectorT &sv) const override { + auto mpi_manager = sv.getMPIManager(); + if (mpi_manager.getRank() == 0) { + PL_ABORT_IF_NOT( + this->wires_.size() == sv.getTotalNumQubits(), + "SparseH wire count does not match state-vector size"); + } + using CFP_t = typename StateVectorT::CFP_t; + + // Distribute sparse matrix across multi-nodes/multi-gpus + size_t num_rows = size_t{1} << sv.getTotalNumQubits(); + size_t local_num_rows = size_t{1} << sv.getNumLocalQubits(); + + std::vector>> csrmatrix_blocks; + + if (mpi_manager.getRank() == 0) { + csrmatrix_blocks = splitCSRMatrix( + mpi_manager, num_rows, this->offsets_.data(), + this->indices_.data(), this->data_.data()); + } + mpi_manager.Barrier(); + + std::vector> localCSRMatVector; + for (size_t i = 0; i < mpi_manager.getSize(); i++) { + auto localCSRMat = scatterCSRMatrix( + mpi_manager, csrmatrix_blocks[i], local_num_rows, 0); + localCSRMatVector.push_back(localCSRMat); + } + + mpi_manager.Barrier(); + + auto device_id = sv.getDataBuffer().getDevTag().getDeviceID(); + auto stream_id = sv.getDataBuffer().getDevTag().getStreamID(); + cusparseHandle_t handle = sv.getCusparseHandle(); + + const size_t length_local = size_t{1} << sv.getNumLocalQubits(); + + std::unique_ptr> d_sv_prime = + std::make_unique>(length_local, device_id, + stream_id, true); + std::unique_ptr> d_tmp = + std::make_unique>(length_local, device_id, + stream_id, true); + d_sv_prime->zeroInit(); + PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); + mpi_manager.Barrier(); + + for (size_t i = 0; i < mpi_manager.getSize(); i++) { + size_t color = 0; + auto &localCSRMatrix = localCSRMatVector[i]; + + if (localCSRMatrix.getValues().size() != 0) { + color = 1; + SparseMV_cuSparse( + localCSRMatrix.getCsrOffsets().data(), + localCSRMatrix.getCsrOffsets().size(), + localCSRMatrix.getColumns().data(), + localCSRMatrix.getValues().data(), + localCSRMatrix.getValues().size(), sv.getData(), + d_sv_prime->getData(), device_id, stream_id, handle); + } + PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); + mpi_manager.Barrier(); + + if (mpi_manager.getRank() == i) { + color = 1; + if (localCSRMatrix.getValues().size() == 0) { + d_tmp->zeroInit(); + } + } + PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); + mpi_manager.Barrier(); + + auto new_mpi_manager = + mpi_manager.split(color, mpi_manager.getRank()); + int reduce_root_rank = -1; + + if (mpi_manager.getRank() == i) { + reduce_root_rank = new_mpi_manager.getRank(); + } + + mpi_manager.template Bcast(reduce_root_rank, i); + + if (new_mpi_manager.getComm() != MPI_COMM_NULL) { + new_mpi_manager.template Reduce( + d_tmp->getData(), d_sv_prime->getData(), length_local, + reduce_root_rank, "sum"); + } + PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); + mpi_manager.Barrier(); + } + sv.CopyGpuDataToGpuIn(d_sv_prime->getData(), d_sv_prime->getLength()); + mpi_manager.Barrier(); + } +}; + } // namespace Pennylane::LightningGPU::Observables From 97d3f853375bc31f268107dc5f014bfb79c8a317 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Fri, 20 Oct 2023 16:00:25 +0000 Subject: [PATCH 02/43] Fix std::endl; --- .../core/src/observables/Observables.hpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index 58770c7af1..d4e9273723 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -495,12 +495,14 @@ class SparseHamiltonianBase : public Observable { std::ostringstream ss; ss << "SparseHamiltonian: {\n'data' : \n"; for (const auto &d : data_) - //Note that for LGPU backend, ComplexT is std::complex as of 0.33 release - //Need to revisit it once we set ComplexT as cuComplex later - ss <<"{" << d.real() <<", "<< d.real() <<"},"<< "\n"; + // Note that for LGPU backend, ComplexT is std::complex as of 0.33 + // release Need to revisit it once we set ComplexT as cuComplex + // later + ss << "{" << d.real() << ", " << d.real() << "}," + << "\n"; ss << ",\n'indices' : \n"; for (const auto &i : indices_) - ss << i <<; + ss << i << std::endl; ss << ",\n'offsets' : \n"; for (const auto &o : offsets_) ss << o; From 705f5497b4874a46167c0395fde1a3c3a0f2a4cf Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Fri, 20 Oct 2023 16:59:56 +0000 Subject: [PATCH 03/43] Use more generic indices in base std::size_t. --- .../core/src/observables/Observables.hpp | 14 ++-- .../observables/ObservablesKokkos.cpp | 3 + .../observables/ObservablesKokkos.hpp | 75 +++++++++++++++++++ 3 files changed, 83 insertions(+), 9 deletions(-) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index d4e9273723..b56b69020c 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -424,15 +424,11 @@ class SparseHamiltonianBase : public Observable { public: using PrecisionT = typename StateVectorT::PrecisionT; using ComplexT = typename StateVectorT::ComplexT; - // cuSparse required index type - using IdxT = - typename std::conditional::value, - int32_t, int64_t>::type; protected: std::vector data_; - std::vector indices_; - std::vector offsets_; + std::vector indices_; + std::vector offsets_; std::vector wires_; private: @@ -480,8 +476,8 @@ class SparseHamiltonianBase : public Observable { * @param arg4 Argument to construct wires */ static auto create(std::initializer_list arg1, - std::initializer_list arg2, - std::initializer_list arg3, + std::initializer_list arg2, + std::initializer_list arg3, std::initializer_list arg4) -> std::shared_ptr> { return std::shared_ptr>( @@ -502,7 +498,7 @@ class SparseHamiltonianBase : public Observable { << "\n"; ss << ",\n'indices' : \n"; for (const auto &i : indices_) - ss << i << std::endl; + ss << i; ss << ",\n'offsets' : \n"; for (const auto &o : offsets_) ss << o; diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.cpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.cpp index d90f3e6019..66192b934a 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.cpp @@ -28,3 +28,6 @@ template class Observables::TensorProdObs>; template class Observables::Hamiltonian>; template class Observables::Hamiltonian>; + +template class Observables::SparseHamiltonian>; +template class Observables::SparseHamiltonian>; diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp index a0371df7f2..860be835f6 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp @@ -31,6 +31,7 @@ namespace { using namespace Pennylane::Util; using namespace Pennylane::Observables; using Pennylane::LightningKokkos::StateVectorKokkos; +using Pennylane::LightningKokkos::Util::SparseMV_Kokkos; } // namespace /// @endcond @@ -199,6 +200,80 @@ class Hamiltonian final : public HamiltonianBase { } }; +/** + * @brief Sparse representation of Hamiltonian + * + */ +template +class SparseHamiltonian final : public SparseHamiltonianBase { + public: + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + + private: + using BaseType = SparseHamiltonianBase; + + public: + /** + * @brief Create a SparseHamiltonian from data, indices and offsets in CSR + * format. + * + * @param data Arguments to construct data + * @param indices Arguments to construct indices + * @param offsets Arguments to construct offsets + * @param wires Arguments to construct wires + */ + template + explicit SparseHamiltonian(T1 &&data, T2 &&indices, T3 &&offsets, + T4 &&wires) + : BaseType{data, indices, offsets, wires} {} + + /** + * @brief Convenient wrapper for the constructor as the constructor does not + * convert the std::shared_ptr with a derived class correctly. + * + * This function is useful as std::make_shared does not handle + * brace-enclosed initializer list correctly. + * + * @param data Argument to construct data + * @param indices Argument to construct indices + * @param offsets Argument to construct ofsets + * @param wires Argument to construct wires + */ + static auto create(std::initializer_list data, + std::initializer_list indices, + std::initializer_list offsets, + std::initializer_list wires) + -> std::shared_ptr> { + return std::shared_ptr>( + new SparseHamiltonian{ + std::move(data), std::move(indices), std::move(offsets), + std::move(wires)}); + } + + /** + * @brief Updates the statevector SV:->SV', where SV' = a*H*SV, and where H + * is a sparse Hamiltonian. + * + */ + void applyInPlace(StateVectorT &sv) const override { + PL_ABORT_IF_NOT(this->wires_.size() == sv.getNumQubits(), + "SparseH wire count does not match state-vector size"); + StateVectorT d_sv_prime(sv.getNumQubits()); + + SparseMV_Kokkos( + sv.getView(), d_sv_prime.getView(), this->indices_.data(), + this->indices_.size(), this->offsets_.data(), this->data_.data(), + this->data_.size()); + // (Kokkos::View x, Kokkos::View y, + // const size_t *row_map, const size_t + // row_map_size, const size_t *entries_ptr, const + // ComplexT *values_ptr, const size_t numNNZ) + + sv.updateData(d_sv_prime); + } +}; + /// @cond DEV namespace detail { using Pennylane::LightningKokkos::Util::axpy_Kokkos; From 343c62a1c667dca9142d8402ba387580f59460d9 Mon Sep 17 00:00:00 2001 From: Shuli <08cnbj@gmail.com> Date: Fri, 20 Oct 2023 19:02:25 +0000 Subject: [PATCH 04/43] add pybind layer --- .../core/src/bindings/Bindings.hpp | 47 +++++++++++++++++++ .../core/src/bindings/BindingsMPI.hpp | 47 +++++++++++++++++++ 2 files changed, 94 insertions(+) diff --git a/pennylane_lightning/core/src/bindings/Bindings.hpp b/pennylane_lightning/core/src/bindings/Bindings.hpp index 30c994a719..a91df3ece6 100644 --- a/pennylane_lightning/core/src/bindings/Bindings.hpp +++ b/pennylane_lightning/core/src/bindings/Bindings.hpp @@ -299,6 +299,10 @@ template void registerObservables(py::module_ &m) { using np_arr_c = py::array_t, py::array::c_style>; using np_arr_r = py::array_t; + using np_arr_sparse_ind = typename std::conditional< + std::is_same::value, + py::array_t, + py::array_t>::type; std::string class_name; @@ -406,6 +410,49 @@ template void registerObservables(py::module_ &m) { return self == other_cast; }, "Compare two observables"); +#ifdef _ENABLE_PLGPU + class_name = "SparseHamiltonianC" + bitsize; + using SpIDX = typename SparseHamiltonian::IdxT; + py::class_, + std::shared_ptr>, + Observable>(m, class_name.c_str(), + py::module_local()) + .def(py::init([](const np_arr_c &data, const np_arr_sparse_ind &indices, + const np_arr_sparse_ind &offsets, + const std::vector &wires) { + const py::buffer_info buffer_data = data.request(); + const auto *data_ptr = + static_cast(buffer_data.ptr); + + const py::buffer_info buffer_indices = indices.request(); + const auto *indices_ptr = static_cast(buffer_indices.ptr); + + const py::buffer_info buffer_offsets = offsets.request(); + const auto *offsets_ptr = static_cast(buffer_offsets.ptr); + + return SparseHamiltonian{ + std::vector( + {data_ptr, data_ptr + data.size()}), + std::vector({indices_ptr, indices_ptr + indices.size()}), + std::vector({offsets_ptr, offsets_ptr + offsets.size()}), + wires}; + })) + .def("__repr__", &SparseHamiltonian::getObsName) + .def("get_wires", &SparseHamiltonian::getWires, + "Get wires of observables") + .def( + "__eq__", + [](const SparseHamiltonianGPU &self, + py::handle other) -> bool { + if (!py::isinstance>(other)) { + return false; + } + auto other_cast = + other.cast>(); + return self == other_cast; + }, + "Compare two observables"); +#endif } /** diff --git a/pennylane_lightning/core/src/bindings/BindingsMPI.hpp b/pennylane_lightning/core/src/bindings/BindingsMPI.hpp index 6c828b0ff3..ae9da8b2b1 100644 --- a/pennylane_lightning/core/src/bindings/BindingsMPI.hpp +++ b/pennylane_lightning/core/src/bindings/BindingsMPI.hpp @@ -82,6 +82,10 @@ template void registerObservablesMPI(py::module_ &m) { using np_arr_c = py::array_t, py::array::c_style>; using np_arr_r = py::array_t; + using np_arr_sparse_ind = typename std::conditional< + std::is_same::value, + py::array_t, + py::array_t>::type; std::string class_name; @@ -191,6 +195,49 @@ template void registerObservablesMPI(py::module_ &m) { return self == other_cast; }, "Compare two observables"); +#ifdef _ENABLE_PLGPU + class_name = "SparseHamiltonianMPIC" + bitsize; + using SpIDX = typename SparseHamiltonianMPI::IdxT; + py::class_, + std::shared_ptr>, + Observable>(m, class_name.c_str(), + py::module_local()) + .def(py::init([](const np_arr_c &data, const np_arr_sparse_ind &indices, + const np_arr_sparse_ind &offsets, + const std::vector &wires) { + const py::buffer_info buffer_data = data.request(); + const auto *data_ptr = + static_cast(buffer_data.ptr); + + const py::buffer_info buffer_indices = indices.request(); + const auto *indices_ptr = static_cast(buffer_indices.ptr); + + const py::buffer_info buffer_offsets = offsets.request(); + const auto *offsets_ptr = static_cast(buffer_offsets.ptr); + + return SparseHamiltonianMPI{ + std::vector( + {data_ptr, data_ptr + data.size()}), + std::vector({indices_ptr, indices_ptr + indices.size()}), + std::vector({offsets_ptr, offsets_ptr + offsets.size()}), + wires}; + })) + .def("__repr__", &SparseHamiltonianMPI::getObsName) + .def("get_wires", &SparseHamiltonianMPI::getWires, + "Get wires of observables") + .def( + "__eq__", + [](const SparseHamiltonianMPI &self, + py::handle other) -> bool { + if (!py::isinstance>(other)) { + return false; + } + auto other_cast = + other.cast>(); + return self == other_cast; + }, + "Compare two observables"); +#endif } /** From 1c0bda2c825cc8288a5b645fb6971520733debd1 Mon Sep 17 00:00:00 2001 From: Shuli <08cnbj@gmail.com> Date: Fri, 20 Oct 2023 19:19:08 +0000 Subject: [PATCH 05/43] add python layer --- pennylane_lightning/core/_serialize.py | 40 +++++++++++++++++++ .../core/src/bindings/Bindings.hpp | 2 +- 2 files changed, 41 insertions(+), 1 deletion(-) diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index 80d6bd9805..9c461f62af 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -52,6 +52,7 @@ class QuantumScriptSerializer: # pylint: disable=import-outside-toplevel, too-many-instance-attributes def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False): self.use_csingle = use_csingle + self.device_name = device_name if device_name == "lightning.qubit": try: import pennylane_lightning.lightning_qubit_ops as lightning_ops @@ -85,6 +86,11 @@ def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False self.hamiltonian_c64 = lightning_ops.observables.HamiltonianC64 self.hamiltonian_c128 = lightning_ops.observables.HamiltonianC128 + if self.device_name == "lightning.gpu": + self.sparsehamiltonian_c64 = lightning_ops.observables.SparseHamiltonianC64 + self.sparsehamiltonian_c128 = lightning_ops.observables.SparseHamiltonianC128 + + self.use_mpi = False if use_mpi: @@ -98,6 +104,11 @@ def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False self.tensor_prod_obsmpi_c128 = lightning_ops.observablesMPI.TensorProdObsMPIC128 self.hamiltonianmpi_c64 = lightning_ops.observablesMPI.HamiltonianMPIC64 self.hamiltonianmpi_c128 = lightning_ops.observablesMPI.HamiltonianMPIC128 + + if self.device_name == "lightning.gpu": + self.sparsehamiltonianmpi_c64 = lightning_ops.observablesMPI.SparseHamiltonianMPIC64 + self.sparsehamiltonianmpi_c128 = lightning_ops.observablesMPI.SparseHamiltoniaMPInC128 + self.mpi_manager = lightning_ops.MPIManager @property def ctype(self): @@ -142,6 +153,13 @@ def hamiltonian_obs(self): if self.use_mpi: return self.hamiltonianmpi_c64 if self.use_csingle else self.hamiltonianmpi_c128 return self.hamiltonian_c64 if self.use_csingle else self.hamiltonian_c128 + + @property + def sparsehamiltonian_obs(self): + """Sparse Hamiltonian observable matching ``use_csingle`` precision.""" + if self.use_mpi: + return self.sparsehamiltonianmpi_c64 if self.use_csingle else self.sparsehamiltonianmpi_c128 + return self.sparsehamiltonian_c64 if self.use_csingle else self.sparsehamiltonian_c128 def _named_obs(self, observable, wires_map: dict): """Serializes a Named observable""" @@ -166,6 +184,26 @@ def _hamiltonian(self, observable, wires_map: dict): coeffs = np.array(unwrap(observable.coeffs)).astype(self.rtype) terms = [self._ob(t, wires_map) for t in observable.ops] return self.hamiltonian_obs(coeffs, terms) + + def _sparsehamiltonian(self, observable, wires_map: dict): + wires = [] + wires_list = observable.wires.tolist() + wires.extend([wires_map[w] for w in wires_list]) + if use_mpi: + Hmat = qml.Hamiltonian([1.0], [qml.Identity(0)]).sparse_matrix() + H_sparse = qml.SparseHamiltonian(Hmat, wires=range(1)) + spm = H_sparse.sparse_matrix() + # Only root 0 needs the overall sparsematrix data + if self.mpi_manager().getRank() == 0: + spm = observable.sparse_matrix() + self.mpi_manager().Barrier() + else: + spm = observable.sparse_matrix() + data = np.array(spm.data).astype(ctype) + indices = np.array(spm.indices).astype(rtype) + offsets = np.array(spm.indptr).astype(rtype) + + return self.sparsehamiltonian_obs(data, indices, offsets, wires) def _pauli_word(self, observable, wires_map: dict): """Serialize a :class:`pennylane.pauli.PauliWord` into a Named or Tensor observable.""" @@ -194,6 +232,8 @@ def _ob(self, observable, wires_map): return self._tensor_ob(observable, wires_map) if observable.name == "Hamiltonian": return self._hamiltonian(observable, wires_map) + if observable.name == "SparseHamiltonian": + return self._sparsehamiltonian(observable, wires_map) if isinstance(observable, (PauliX, PauliY, PauliZ, Identity, Hadamard)): return self._named_obs(observable, wires_map) if observable._pauli_rep is not None: diff --git a/pennylane_lightning/core/src/bindings/Bindings.hpp b/pennylane_lightning/core/src/bindings/Bindings.hpp index a91df3ece6..501f4905bd 100644 --- a/pennylane_lightning/core/src/bindings/Bindings.hpp +++ b/pennylane_lightning/core/src/bindings/Bindings.hpp @@ -442,7 +442,7 @@ template void registerObservables(py::module_ &m) { "Get wires of observables") .def( "__eq__", - [](const SparseHamiltonianGPU &self, + [](const SparseHamiltonian &self, py::handle other) -> bool { if (!py::isinstance>(other)) { return false; From 255f6d5bbb7593fe1cc2f35ffad0287f62afe607 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Fri, 20 Oct 2023 19:34:15 +0000 Subject: [PATCH 06/43] Quick and dirty spham bindings. --- .../core/src/bindings/Bindings.hpp | 34 +++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/pennylane_lightning/core/src/bindings/Bindings.hpp b/pennylane_lightning/core/src/bindings/Bindings.hpp index 30c994a719..e2989202ca 100644 --- a/pennylane_lightning/core/src/bindings/Bindings.hpp +++ b/pennylane_lightning/core/src/bindings/Bindings.hpp @@ -406,6 +406,40 @@ template void registerObservables(py::module_ &m) { return self == other_cast; }, "Compare two observables"); + +#if _ENABLE_PLKOKKOS == 1 + class_name = "SparseHamiltonianC" + bitsize; + py::class_, + std::shared_ptr>, + Observable>(m, class_name.c_str(), + py::module_local()) + .def(py::init([](const np_arr_c &data, + const std::vector &indices, + const std::vector &indptr, + const std::vector &wires) { + using ComplexT = typename StateVectorT::ComplexT; + const py::buffer_info buffer_data = data.request(); + const auto *data_ptr = static_cast(buffer_data.ptr); + + return SparseHamiltonian{ + std::vector({data_ptr, data_ptr + data.size()}), + indices, indptr, wires}; + })) + .def("__repr__", &SparseHamiltonian::getObsName) + .def("get_wires", &SparseHamiltonian::getWires, + "Get wires of observables") + .def( + "__eq__", + [](const SparseHamiltonian &self, + py::handle other) -> bool { + if (!py::isinstance>(other)) { + return false; + } + auto other_cast = other.cast>(); + return self == other_cast; + }, + "Compare two observables"); +#endif } /** From 25b90c84a0b6e41514dba2df0979c7e2bae8beb2 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Fri, 20 Oct 2023 20:09:05 +0000 Subject: [PATCH 07/43] Add sparse_ham serialization. --- pennylane_lightning/core/_serialize.py | 33 ++++++++++++++++++++++++++ tests/test_serialize.py | 6 +++++ 2 files changed, 39 insertions(+) diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index 12981faf92..a592239731 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -84,6 +84,9 @@ def __init__(self, device_name, use_csingle: bool = False): self.tensor_prod_obs_c128 = lightning_ops.observables.TensorProdObsC128 self.hamiltonian_c64 = lightning_ops.observables.HamiltonianC64 self.hamiltonian_c128 = lightning_ops.observables.HamiltonianC128 + if device_name == "lightning.kokkos": + self.sparse_hamiltonian_c64 = lightning_ops.observables.SparseHamiltonianC64 + self.sparse_hamiltonian_c128 = lightning_ops.observables.SparseHamiltonianC128 @property def ctype(self): @@ -115,6 +118,11 @@ def hamiltonian_obs(self): """Hamiltonian observable matching ``use_csingle`` precision.""" return self.hamiltonian_c64 if self.use_csingle else self.hamiltonian_c128 + @property + def sparse_hamiltonian_obs(self): + """SparseHamiltonian observable matching ``use_csingle`` precision.""" + return self.sparse_hamiltonian_c64 if self.use_csingle else self.sparse_hamiltonian_c128 + def _named_obs(self, observable, wires_map: dict): """Serializes a Named observable""" wires = [wires_map[w] for w in observable.wires] @@ -139,6 +147,28 @@ def _hamiltonian(self, observable, wires_map: dict): terms = [self._ob(t, wires_map) for t in observable.ops] return self.hamiltonian_obs(coeffs, terms) + def _sparse_hamiltonian(self, observable, wires_map: dict): + """Serialize an observable (Sparse Hamiltonian) + + Args: + observable (Observable): the input observable (Sparse Hamiltonian) + wire_map (dict): a dictionary mapping input wires to the device's backend wires + + Returns: + sparsehamiltonian_obs (SparseHamiltonianKokkos_C64 or SparseHamiltonianKokkos_C128): A Sparse Hamiltonian observable object compatible with the C++ backend + """ + + spm = observable.sparse_matrix() + data = np.array(spm.data).astype(self.ctype) + indices = np.array(spm.indices).astype(np.int) + offsets = np.array(spm.indptr).astype(np.int) + + wires = [] + wires_list = observable.wires.tolist() + wires.extend([wires_map[w] for w in wires_list]) + + return self.sparse_hamiltonian_obs(data, indices, offsets, wires) + def _pauli_word(self, observable, wires_map: dict): """Serialize a :class:`pennylane.pauli.PauliWord` into a Named or Tensor observable.""" if len(observable) == 1: @@ -162,10 +192,13 @@ def _pauli_sentence(self, observable, wires_map: dict): # pylint: disable=protected-access def _ob(self, observable, wires_map): """Serialize a :class:`pennylane.operation.Observable` into an Observable.""" + print(observable.__class__) if isinstance(observable, Tensor): return self._tensor_ob(observable, wires_map) if observable.name == "Hamiltonian": return self._hamiltonian(observable, wires_map) + if observable.name == "SparseHamiltonian": + return self._sparse_hamiltonian(observable, wires_map) if isinstance(observable, (PauliX, PauliY, PauliZ, Identity, Hadamard)): return self._named_obs(observable, wires_map) if observable._pauli_rep is not None: diff --git a/tests/test_serialize.py b/tests/test_serialize.py index ab6df3c26e..a1e2d74323 100644 --- a/tests/test_serialize.py +++ b/tests/test_serialize.py @@ -34,6 +34,8 @@ TensorProdObsC128, HamiltonianC64, HamiltonianC128, + SparseHamiltonianC64, + SparseHamiltonianC128, ) elif device_name == "lightning.gpu": from pennylane_lightning.lightning_gpu_ops.observables import ( @@ -92,6 +94,10 @@ def test_wrong_device_name(): (qml.Projector([0], wires=0), HermitianObsC128), (qml.Hamiltonian([1], [qml.PauliZ(0)]), HamiltonianC128), (qml.sum(qml.Hadamard(0), qml.PauliX(1)), HermitianObsC128), + ( + qml.SparseHamiltonian(qml.Hamiltonian([1], [qml.PauliZ(0)]).sparse_matrix(), wires=[0]), + SparseHamiltonianC128, + ), ], ) def test_obs_returns_expected_type(obs, obs_type): From e815e6f81cce11639a5270239c869d432b16461c Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Fri, 20 Oct 2023 20:35:29 +0000 Subject: [PATCH 08/43] Add sparse_ham tests in tests/test_adjoint_jacobian.py' --- pennylane_lightning/core/_serialize.py | 7 ++- tests/test_adjoint_jacobian.py | 60 ++++++++++++++++++++++++-- 2 files changed, 59 insertions(+), 8 deletions(-) diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index a592239731..eba82f1d80 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -155,13 +155,13 @@ def _sparse_hamiltonian(self, observable, wires_map: dict): wire_map (dict): a dictionary mapping input wires to the device's backend wires Returns: - sparsehamiltonian_obs (SparseHamiltonianKokkos_C64 or SparseHamiltonianKokkos_C128): A Sparse Hamiltonian observable object compatible with the C++ backend + sparse_hamiltonian_obs (SparseHamiltonianC64 or SparseHamiltonianC128): A Sparse Hamiltonian observable object compatible with the C++ backend """ spm = observable.sparse_matrix() data = np.array(spm.data).astype(self.ctype) - indices = np.array(spm.indices).astype(np.int) - offsets = np.array(spm.indptr).astype(np.int) + indices = np.array(spm.indices).astype(np.int64) + offsets = np.array(spm.indptr).astype(np.int64) wires = [] wires_list = observable.wires.tolist() @@ -192,7 +192,6 @@ def _pauli_sentence(self, observable, wires_map: dict): # pylint: disable=protected-access def _ob(self, observable, wires_map): """Serialize a :class:`pennylane.operation.Observable` into an Observable.""" - print(observable.__class__) if isinstance(observable, Tensor): return self._tensor_ob(observable, wires_map) if observable.name == "Hamiltonian": diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py index c8d65dd093..f65a5f1928 100644 --- a/tests/test_adjoint_jacobian.py +++ b/tests/test_adjoint_jacobian.py @@ -25,6 +25,8 @@ from pennylane import QNode, qnode from pennylane import qchem + + I, X, Y, Z = ( np.eye(2), qml.PauliX.compute_matrix(), @@ -856,7 +858,6 @@ def f(params1, params2): assert np.allclose(grad_adjoint, grad_fd, atol=tol) - def circuit_ansatz(params, wires): """Circuit ansatz containing all the parametrized gates""" qml.QubitStateVector(unitary_group.rvs(2**4, random_state=0)[0], wires=wires) @@ -943,6 +944,60 @@ def circuit(params): assert np.allclose(qml.grad(circuit_ld)(params), qml.grad(circuit_dq)(params), tol) +custom_wires = ["alice", 3.14, -1, 0] + + +@pytest.mark.skipif(not ld._CPP_BINARY_AVAILABLE, reason="Lightning binary required") +@pytest.mark.skipif(device_name != "lightning.kokkos", reason="SparseHamiltonian only supported by Lightning-Kokkos") +@pytest.mark.parametrize( + "returns", + [ + qml.SparseHamiltonian( + qml.Hamiltonian( + [0.1], + [qml.PauliX(wires=custom_wires[0]) @ qml.PauliZ(wires=custom_wires[1])], + ).sparse_matrix(custom_wires), + wires=custom_wires, + ), + qml.SparseHamiltonian( + qml.Hamiltonian( + [2.0], + [qml.PauliX(wires=custom_wires[2]) @ qml.PauliZ(wires=custom_wires[0])], + ).sparse_matrix(custom_wires), + wires=custom_wires, + ), + qml.SparseHamiltonian( + qml.Hamiltonian( + [1.1], + [qml.PauliX(wires=custom_wires[0]) @ qml.PauliZ(wires=custom_wires[2])], + ).sparse_matrix(custom_wires), + wires=custom_wires, + ), + ], +) +def test_adjoint_SparseHamiltonian(returns): + """Integration tests that compare to default.qubit for a large circuit containing parametrized + operations and when using custom wire labels""" + + dev_kokkos = qml.device("lightning.kokkos", wires=custom_wires) + dev_default = qml.device("default.qubit", wires=custom_wires) + + def circuit(params): + circuit_ansatz(params, wires=custom_wires) + return qml.expval(returns) + + n_params = 30 + np.random.seed(1337) + params = np.random.rand(n_params) + + qnode_kokkos = qml.QNode(circuit, dev_kokkos, diff_method="adjoint") + qnode_default = qml.QNode(circuit, dev_default, diff_method="parameter-shift") + + j_kokkos = qml.jacobian(qnode_kokkos)(params) + j_default = qml.jacobian(qnode_default)(params) + + assert np.allclose(j_kokkos, j_default) + @pytest.mark.parametrize( "returns", @@ -1024,9 +1079,6 @@ def casted_to_array_batched(params): assert np.allclose(j_def, j_lightning_batched) -custom_wires = ["alice", 3.14, -1, 0] - - @pytest.mark.parametrize( "returns", [ From 7d8f0d93c261e7729044533a1a1fb6a0d39f94b1 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Fri, 20 Oct 2023 20:58:15 +0000 Subject: [PATCH 09/43] Bug fix sparse product. --- .../lightning_kokkos/observables/ObservablesKokkos.hpp | 8 ++------ tests/test_adjoint_jacobian.py | 7 +++++-- 2 files changed, 7 insertions(+), 8 deletions(-) diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp index 860be835f6..cfbe9bd171 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp @@ -262,13 +262,9 @@ class SparseHamiltonian final : public SparseHamiltonianBase { StateVectorT d_sv_prime(sv.getNumQubits()); SparseMV_Kokkos( - sv.getView(), d_sv_prime.getView(), this->indices_.data(), - this->indices_.size(), this->offsets_.data(), this->data_.data(), + sv.getView(), d_sv_prime.getView(), this->offsets_.data(), + this->offsets_.size(), this->indices_.data(), this->data_.data(), this->data_.size()); - // (Kokkos::View x, Kokkos::View y, - // const size_t *row_map, const size_t - // row_map_size, const size_t *entries_ptr, const - // ComplexT *values_ptr, const size_t numNNZ) sv.updateData(d_sv_prime); } diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py index f65a5f1928..4a947b7c8d 100644 --- a/tests/test_adjoint_jacobian.py +++ b/tests/test_adjoint_jacobian.py @@ -26,7 +26,6 @@ from pennylane import qchem - I, X, Y, Z = ( np.eye(2), qml.PauliX.compute_matrix(), @@ -858,6 +857,7 @@ def f(params1, params2): assert np.allclose(grad_adjoint, grad_fd, atol=tol) + def circuit_ansatz(params, wires): """Circuit ansatz containing all the parametrized gates""" qml.QubitStateVector(unitary_group.rvs(2**4, random_state=0)[0], wires=wires) @@ -944,11 +944,14 @@ def circuit(params): assert np.allclose(qml.grad(circuit_ld)(params), qml.grad(circuit_dq)(params), tol) + custom_wires = ["alice", 3.14, -1, 0] @pytest.mark.skipif(not ld._CPP_BINARY_AVAILABLE, reason="Lightning binary required") -@pytest.mark.skipif(device_name != "lightning.kokkos", reason="SparseHamiltonian only supported by Lightning-Kokkos") +@pytest.mark.skipif( + device_name != "lightning.kokkos", reason="SparseHamiltonian only supported by Lightning-Kokkos" +) @pytest.mark.parametrize( "returns", [ From 71ebe4b58218857a2df5f5b9ca47011a1fd03666 Mon Sep 17 00:00:00 2001 From: Shuli <08cnbj@gmail.com> Date: Sat, 21 Oct 2023 04:14:30 +0000 Subject: [PATCH 10/43] add sparseH --- pennylane_lightning/core/_serialize.py | 25 ++++--- .../core/src/bindings/Bindings.hpp | 11 +-- .../core/src/bindings/BindingsMPI.hpp | 11 ++- .../lightning_gpu/bindings/LGPUBindings.hpp | 23 ++++-- .../bindings/LGPUBindingsMPI.hpp | 19 +++-- .../measurements/MeasurementsGPU.hpp | 14 ++-- .../measurements/MeasurementsGPUMPI.hpp | 15 ++-- .../Test_StateVectorCudaManaged_Expval.cpp | 34 +++++---- .../mpi/Test_StateVectorCudaMPI_Expval.cpp | 34 +++++---- .../observables/ObservablesGPU.hpp | 5 +- .../observables/ObservablesGPUMPI.hpp | 28 +++++-- .../observables/tests/Test_ObservablesGPU.cpp | 28 +++++++ .../tests/mpi/Test_ObservablesGPUMPI.cpp | 75 +++++++++++++++++++ .../lightning_gpu/utils/LinearAlg.hpp | 49 ++++++------ .../lightning_gpu/utils/MPILinearAlg.hpp | 12 +-- .../utils/tests/Test_LinearAlgebra.cpp | 16 ++-- .../utils/tests/mpi/Test_LinearAlgebraMPI.cpp | 15 ++-- .../core/src/utils/TestHelpers.hpp | 28 +++---- tests/test_adjoint_jacobian.py | 10 ++- 19 files changed, 302 insertions(+), 150 deletions(-) diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index 9c461f62af..9cb603ba7e 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -90,7 +90,6 @@ def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False self.sparsehamiltonian_c64 = lightning_ops.observables.SparseHamiltonianC64 self.sparsehamiltonian_c128 = lightning_ops.observables.SparseHamiltonianC128 - self.use_mpi = False if use_mpi: @@ -104,10 +103,12 @@ def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False self.tensor_prod_obsmpi_c128 = lightning_ops.observablesMPI.TensorProdObsMPIC128 self.hamiltonianmpi_c64 = lightning_ops.observablesMPI.HamiltonianMPIC64 self.hamiltonianmpi_c128 = lightning_ops.observablesMPI.HamiltonianMPIC128 - + if self.device_name == "lightning.gpu": self.sparsehamiltonianmpi_c64 = lightning_ops.observablesMPI.SparseHamiltonianMPIC64 - self.sparsehamiltonianmpi_c128 = lightning_ops.observablesMPI.SparseHamiltoniaMPInC128 + self.sparsehamiltonianmpi_c128 = ( + lightning_ops.observablesMPI.SparseHamiltonianMPIC128 + ) self.mpi_manager = lightning_ops.MPIManager @property @@ -153,12 +154,16 @@ def hamiltonian_obs(self): if self.use_mpi: return self.hamiltonianmpi_c64 if self.use_csingle else self.hamiltonianmpi_c128 return self.hamiltonian_c64 if self.use_csingle else self.hamiltonian_c128 - + @property def sparsehamiltonian_obs(self): """Sparse Hamiltonian observable matching ``use_csingle`` precision.""" if self.use_mpi: - return self.sparsehamiltonianmpi_c64 if self.use_csingle else self.sparsehamiltonianmpi_c128 + return ( + self.sparsehamiltonianmpi_c64 + if self.use_csingle + else self.sparsehamiltonianmpi_c128 + ) return self.sparsehamiltonian_c64 if self.use_csingle else self.sparsehamiltonian_c128 def _named_obs(self, observable, wires_map: dict): @@ -184,12 +189,12 @@ def _hamiltonian(self, observable, wires_map: dict): coeffs = np.array(unwrap(observable.coeffs)).astype(self.rtype) terms = [self._ob(t, wires_map) for t in observable.ops] return self.hamiltonian_obs(coeffs, terms) - + def _sparsehamiltonian(self, observable, wires_map: dict): wires = [] wires_list = observable.wires.tolist() wires.extend([wires_map[w] for w in wires_list]) - if use_mpi: + if self.use_mpi: Hmat = qml.Hamiltonian([1.0], [qml.Identity(0)]).sparse_matrix() H_sparse = qml.SparseHamiltonian(Hmat, wires=range(1)) spm = H_sparse.sparse_matrix() @@ -199,9 +204,9 @@ def _sparsehamiltonian(self, observable, wires_map: dict): self.mpi_manager().Barrier() else: spm = observable.sparse_matrix() - data = np.array(spm.data).astype(ctype) - indices = np.array(spm.indices).astype(rtype) - offsets = np.array(spm.indptr).astype(rtype) + data = np.array(spm.data).astype(self.ctype) + indices = np.array(spm.indices).astype(self.rtype) + offsets = np.array(spm.indptr).astype(self.rtype) return self.sparsehamiltonian_obs(data, indices, offsets, wires) diff --git a/pennylane_lightning/core/src/bindings/Bindings.hpp b/pennylane_lightning/core/src/bindings/Bindings.hpp index 501f4905bd..1c8fa5ba6c 100644 --- a/pennylane_lightning/core/src/bindings/Bindings.hpp +++ b/pennylane_lightning/core/src/bindings/Bindings.hpp @@ -416,13 +416,12 @@ template void registerObservables(py::module_ &m) { py::class_, std::shared_ptr>, Observable>(m, class_name.c_str(), - py::module_local()) + py::module_local()) .def(py::init([](const np_arr_c &data, const np_arr_sparse_ind &indices, const np_arr_sparse_ind &offsets, const std::vector &wires) { const py::buffer_info buffer_data = data.request(); - const auto *data_ptr = - static_cast(buffer_data.ptr); + const auto *data_ptr = static_cast(buffer_data.ptr); const py::buffer_info buffer_indices = indices.request(); const auto *indices_ptr = static_cast(buffer_indices.ptr); @@ -431,8 +430,7 @@ template void registerObservables(py::module_ &m) { const auto *offsets_ptr = static_cast(buffer_offsets.ptr); return SparseHamiltonian{ - std::vector( - {data_ptr, data_ptr + data.size()}), + std::vector({data_ptr, data_ptr + data.size()}), std::vector({indices_ptr, indices_ptr + indices.size()}), std::vector({offsets_ptr, offsets_ptr + offsets.size()}), wires}; @@ -447,8 +445,7 @@ template void registerObservables(py::module_ &m) { if (!py::isinstance>(other)) { return false; } - auto other_cast = - other.cast>(); + auto other_cast = other.cast>(); return self == other_cast; }, "Compare two observables"); diff --git a/pennylane_lightning/core/src/bindings/BindingsMPI.hpp b/pennylane_lightning/core/src/bindings/BindingsMPI.hpp index ae9da8b2b1..41276afe5d 100644 --- a/pennylane_lightning/core/src/bindings/BindingsMPI.hpp +++ b/pennylane_lightning/core/src/bindings/BindingsMPI.hpp @@ -201,13 +201,12 @@ template void registerObservablesMPI(py::module_ &m) { py::class_, std::shared_ptr>, Observable>(m, class_name.c_str(), - py::module_local()) + py::module_local()) .def(py::init([](const np_arr_c &data, const np_arr_sparse_ind &indices, const np_arr_sparse_ind &offsets, const std::vector &wires) { const py::buffer_info buffer_data = data.request(); - const auto *data_ptr = - static_cast(buffer_data.ptr); + const auto *data_ptr = static_cast(buffer_data.ptr); const py::buffer_info buffer_indices = indices.request(); const auto *indices_ptr = static_cast(buffer_indices.ptr); @@ -216,8 +215,7 @@ template void registerObservablesMPI(py::module_ &m) { const auto *offsets_ptr = static_cast(buffer_offsets.ptr); return SparseHamiltonianMPI{ - std::vector( - {data_ptr, data_ptr + data.size()}), + std::vector({data_ptr, data_ptr + data.size()}), std::vector({indices_ptr, indices_ptr + indices.size()}), std::vector({offsets_ptr, offsets_ptr + offsets.size()}), wires}; @@ -229,7 +227,8 @@ template void registerObservablesMPI(py::module_ &m) { "__eq__", [](const SparseHamiltonianMPI &self, py::handle other) -> bool { - if (!py::isinstance>(other)) { + if (!py::isinstance>( + other)) { return false; } auto other_cast = 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 e713ea2eef..3f4e81b015 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp @@ -188,10 +188,13 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { using np_arr_c = py::array_t, py::array::c_style | py::array::forcecast>; - using sparse_index_type = std::size_t; - using np_arr_sparse_ind = - py::array_t; + using sparse_index_type = + typename std::conditional::value, int32_t, + int64_t>::type; + using np_arr_sparse_ind = typename std::conditional< + std::is_same::value, + py::array_t, + py::array_t>::type; pyclass .def("expval", @@ -205,10 +208,14 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { const np_arr_sparse_ind &entries, const np_arr_c &values) { return M.expval( static_cast(row_map.request().ptr), - static_cast(row_map.request().size), + static_cast( + row_map.request() + .size), // int64_t is required by cusparse static_cast(entries.request().ptr), static_cast(values.request().ptr), - static_cast(values.request().size)); + static_cast( + values.request() + .size)); // int64_t is required by cusparse }, "Expected value of a sparse Hamiltonian.") .def( @@ -249,10 +256,10 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { const np_arr_sparse_ind &entries, const np_arr_c &values) { return M.var( static_cast(row_map.request().ptr), - static_cast(row_map.request().size), + static_cast(row_map.request().size), static_cast(entries.request().ptr), static_cast(values.request().ptr), - static_cast(values.request().size)); + static_cast(values.request().size)); }, "Variance of a sparse Hamiltonian."); } diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindingsMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindingsMPI.hpp index 40568fd16b..25bee731a4 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindingsMPI.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindingsMPI.hpp @@ -191,10 +191,13 @@ void registerBackendSpecificMeasurementsMPI(PyClass &pyclass) { using np_arr_c = py::array_t, py::array::c_style | py::array::forcecast>; - using sparse_index_type = std::size_t; - using np_arr_sparse_ind = - py::array_t; + using sparse_index_type = + typename std::conditional::value, int32_t, + int64_t>::type; + using np_arr_sparse_ind = typename std::conditional< + std::is_same::value, + py::array_t, + py::array_t>::type; pyclass .def("expval", @@ -209,10 +212,10 @@ void registerBackendSpecificMeasurementsMPI(PyClass &pyclass) { const np_arr_sparse_ind &entries, const np_arr_c &values) { return M.expval( static_cast(row_map.request().ptr), - static_cast(row_map.request().size), + static_cast(row_map.request().size), static_cast(entries.request().ptr), static_cast(values.request().ptr), - static_cast(values.request().size)); + static_cast(values.request().size)); }, "Expected value of a sparse Hamiltonian.") .def( @@ -254,10 +257,10 @@ void registerBackendSpecificMeasurementsMPI(PyClass &pyclass) { const np_arr_sparse_ind &entries, const np_arr_c &values) { return M.var( static_cast(row_map.request().ptr), - static_cast(row_map.request().size), + static_cast(row_map.request().size), static_cast(entries.request().ptr), static_cast(values.request().ptr), - static_cast(values.request().size)); + static_cast(values.request().size)); }, "Variance of a sparse Hamiltonian."); } diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/MeasurementsGPU.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/MeasurementsGPU.hpp index af2ae38ec2..8ca6eacc69 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/MeasurementsGPU.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/MeasurementsGPU.hpp @@ -259,10 +259,10 @@ class Measurements final * @return auto Expectation value. */ template - auto expval(const index_type *csrOffsets_ptr, - const index_type csrOffsets_size, const index_type *columns_ptr, + auto expval(const index_type *csrOffsets_ptr, const int64_t csrOffsets_size, + const index_type *columns_ptr, const std::complex *values_ptr, - const index_type numNNZ) -> PrecisionT { + const int64_t numNNZ) -> PrecisionT { const std::size_t nIndexBits = this->_statevector.getNumQubits(); const std::size_t length = std::size_t{1} << nIndexBits; @@ -580,10 +580,10 @@ class Measurements final * @return Floating point with the variance of the sparse Hamiltonian. */ template - PrecisionT - var(const index_type *csrOffsets_ptr, const index_type csrOffsets_size, - const index_type *columns_ptr, - const std::complex *values_ptr, const index_type numNNZ) { + PrecisionT var(const index_type *csrOffsets_ptr, + const int64_t csrOffsets_size, const index_type *columns_ptr, + const std::complex *values_ptr, + const int64_t numNNZ) { PL_ABORT_IF( (this->_statevector.getLength() != (size_t(csrOffsets_size) - 1)), "Statevector and Hamiltonian have incompatible sizes."); diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/MeasurementsGPUMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/MeasurementsGPUMPI.hpp index f96e2bc217..ff101654df 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/MeasurementsGPUMPI.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/MeasurementsGPUMPI.hpp @@ -47,6 +47,7 @@ using namespace Pennylane; using namespace Pennylane::Measures; using namespace Pennylane::Observables; using namespace Pennylane::LightningGPU::Observables; +using namespace Pennylane::LightningGPU::MPI; namespace cuUtil = Pennylane::LightningGPU::Util; using Pennylane::LightningGPU::StateVectorCudaManaged; using namespace Pennylane::Util; @@ -366,10 +367,10 @@ class MeasurementsMPI final * @return auto Expectation value. */ template - auto expval(const index_type *csrOffsets_ptr, - const index_type csrOffsets_size, const index_type *columns_ptr, + auto expval(const index_type *csrOffsets_ptr, const int64_t csrOffsets_size, + const index_type *columns_ptr, const std::complex *values_ptr, - const index_type numNNZ) -> PrecisionT { + const int64_t numNNZ) -> PrecisionT { if (mpi_manager_.getRank() == 0) { PL_ABORT_IF_NOT( static_cast(csrOffsets_size - 1) == @@ -657,10 +658,10 @@ class MeasurementsMPI final * @return Floating point with the variance of the sparse Hamiltonian. */ template - PrecisionT - var(const index_type *csrOffsets_ptr, const index_type csrOffsets_size, - const index_type *columns_ptr, - const std::complex *values_ptr, const index_type numNNZ) { + PrecisionT var(const index_type *csrOffsets_ptr, + const int64_t csrOffsets_size, const index_type *columns_ptr, + const std::complex *values_ptr, + const int64_t numNNZ) { if (mpi_manager_.getRank() == 0) { PL_ABORT_IF_NOT( static_cast(csrOffsets_size - 1) == diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/tests/Test_StateVectorCudaManaged_Expval.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/tests/Test_StateVectorCudaManaged_Expval.cpp index 82b4d41e93..36f1f1f128 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/tests/Test_StateVectorCudaManaged_Expval.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/tests/Test_StateVectorCudaManaged_Expval.cpp @@ -336,6 +336,8 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::Hamiltonian_expval_Sparse", "[StateVectorCudaManaged_Expval]", float, double) { using StateVectorT = StateVectorCudaManaged; using ComplexT = StateVectorT::ComplexT; + using IdxT = typename std::conditional::value, + int32_t, int64_t>::type; SECTION("Sparse expval") { std::vector init_state{{0.0, 0.0}, {0.0, 0.1}, {0.1, 0.1}, @@ -344,17 +346,18 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::Hamiltonian_expval_Sparse", StateVectorT sv{init_state.data(), init_state.size()}; auto m = Measurements(sv); - std::vector index_ptr = {0, 2, 4, 6, 8, 10, 12, 14, 16}; - std::vector indices = {0, 3, 1, 2, 1, 2, 0, 3, - 4, 7, 5, 6, 5, 6, 4, 7}; + std::vector index_ptr = {0, 2, 4, 6, 8, 10, 12, 14, 16}; + std::vector indices = {0, 3, 1, 2, 1, 2, 0, 3, + 4, 7, 5, 6, 5, 6, 4, 7}; std::vector values = { {3.1415, 0.0}, {0.0, -3.1415}, {3.1415, 0.0}, {0.0, 3.1415}, {0.0, -3.1415}, {3.1415, 0.0}, {0.0, 3.1415}, {3.1415, 0.0}, {3.1415, 0.0}, {0.0, -3.1415}, {3.1415, 0.0}, {0.0, 3.1415}, {0.0, -3.1415}, {3.1415, 0.0}, {0.0, 3.1415}, {3.1415, 0.0}}; - auto result = m.expval(index_ptr.data(), index_ptr.size(), - indices.data(), values.data(), values.size()); + auto result = m.expval( + index_ptr.data(), static_cast(index_ptr.size()), + indices.data(), values.data(), static_cast(values.size())); auto expected = TestType(3.1415); CHECK(expected == Approx(result).epsilon(1e-7)); } @@ -372,22 +375,23 @@ TEMPLATE_TEST_CASE("StateVectorCudaManaged::Hamiltonian_expval_Sparse", // measurements. Measurements Measurer(sv); const size_t num_qubits = 3; - const size_t data_size = Pennylane::Util::exp2(num_qubits); + size_t data_size = Pennylane::Util::exp2(num_qubits); - std::vector row_map; - std::vector entries; + std::vector row_map; + std::vector entries; std::vector values; - write_CSR_vectors(row_map, entries, values, data_size); + write_CSR_vectors(row_map, entries, values, + static_cast(data_size)); - PrecisionT exp_values = - Measurer.expval(row_map.data(), row_map.size(), entries.data(), - values.data(), values.size()); + PrecisionT exp_values = Measurer.expval( + row_map.data(), static_cast(row_map.size()), + entries.data(), values.data(), static_cast(values.size())); PrecisionT exp_values_ref = 0.5930885; REQUIRE(exp_values == Approx(exp_values_ref).margin(1e-6)); - PrecisionT var_values = - Measurer.var(row_map.data(), row_map.size(), entries.data(), - values.data(), values.size()); + PrecisionT var_values = Measurer.var( + row_map.data(), static_cast(row_map.size()), + entries.data(), values.data(), static_cast(values.size())); PrecisionT var_values_ref = 2.4624654; REQUIRE(var_values == Approx(var_values_ref).margin(1e-6)); } diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/tests/mpi/Test_StateVectorCudaMPI_Expval.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/tests/mpi/Test_StateVectorCudaMPI_Expval.cpp index 97b8e0b151..1ef5d75970 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/tests/mpi/Test_StateVectorCudaMPI_Expval.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/measurements/tests/mpi/Test_StateVectorCudaMPI_Expval.cpp @@ -387,6 +387,8 @@ TEMPLATE_TEST_CASE("StateVectorCudaMPI::Hamiltonian_expval_Sparse", "[StateVectorCudaMPI_Expval]", float, double) { using StateVectorT = StateVectorCudaMPI; using ComplexT = StateVectorT::ComplexT; + using IdxT = typename std::conditional::value, + int32_t, int64_t>::type; MPIManager mpi_manager(MPI_COMM_WORLD); @@ -415,17 +417,18 @@ TEMPLATE_TEST_CASE("StateVectorCudaMPI::Hamiltonian_expval_Sparse", sv.CopyHostDataToGpu(local_init_sv.data(), local_init_sv.size(), false); auto m = MeasurementsMPI(sv); - std::vector index_ptr = {0, 2, 4, 6, 8, 10, 12, 14, 16}; - std::vector indices = {0, 3, 1, 2, 1, 2, 0, 3, - 4, 7, 5, 6, 5, 6, 4, 7}; + std::vector index_ptr = {0, 2, 4, 6, 8, 10, 12, 14, 16}; + std::vector indices = {0, 3, 1, 2, 1, 2, 0, 3, + 4, 7, 5, 6, 5, 6, 4, 7}; std::vector values = { {3.1415, 0.0}, {0.0, -3.1415}, {3.1415, 0.0}, {0.0, 3.1415}, {0.0, -3.1415}, {3.1415, 0.0}, {0.0, 3.1415}, {3.1415, 0.0}, {3.1415, 0.0}, {0.0, -3.1415}, {3.1415, 0.0}, {0.0, 3.1415}, {0.0, -3.1415}, {3.1415, 0.0}, {0.0, 3.1415}, {3.1415, 0.0}}; - auto result = m.expval(index_ptr.data(), index_ptr.size(), - indices.data(), values.data(), values.size()); + auto result = m.expval( + index_ptr.data(), static_cast(index_ptr.size()), + indices.data(), values.data(), static_cast(values.size())); auto expected = TestType(3.1415); CHECK(expected == Approx(result).epsilon(1e-7)); } @@ -448,24 +451,25 @@ TEMPLATE_TEST_CASE("StateVectorCudaMPI::Hamiltonian_expval_Sparse", // This object attaches to the statevector allowing several // measurements. MeasurementsMPI Measurer(sv); - const size_t data_size = Pennylane::Util::exp2(num_qubits); + size_t data_size = Pennylane::Util::exp2(num_qubits); - std::vector row_map; - std::vector entries; + std::vector row_map; + std::vector entries; std::vector values; - write_CSR_vectors(row_map, entries, values, data_size); + write_CSR_vectors(row_map, entries, values, + static_cast(data_size)); - PrecisionT exp_values = - Measurer.expval(row_map.data(), row_map.size(), entries.data(), - values.data(), values.size()); + PrecisionT exp_values = Measurer.expval( + row_map.data(), static_cast(row_map.size()), + entries.data(), values.data(), static_cast(values.size())); PrecisionT exp_values_ref = 0.5930885; REQUIRE(exp_values == Approx(exp_values_ref).margin(1e-6)); mpi_manager.Barrier(); - PrecisionT var_values = - Measurer.var(row_map.data(), row_map.size(), entries.data(), - values.data(), values.size()); + PrecisionT var_values = Measurer.var( + row_map.data(), static_cast(row_map.size()), + entries.data(), values.data(), static_cast(values.size())); PrecisionT var_values_ref = 2.4624654; REQUIRE(var_values == Approx(var_values_ref).margin(1e-6)); } diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp index c4042a938c..cf7898340e 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp @@ -287,8 +287,9 @@ class SparseHamiltonian final : public SparseHamiltonianBase { true); SparseMV_cuSparse( - this->offsets_.data(), this->offsets_.size(), this->indices_.data(), - this->data_.data(), this->data_.size(), sv.getData(), + this->offsets_.data(), static_cast(this->offsets_.size()), + this->indices_.data(), this->data_.data(), + static_cast(this->data_.size()), sv.getData(), d_sv_prime->getData(), device_id, stream_id, handle); sv.updateData(std::move(d_sv_prime)); diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp index ed349cf233..86a5dcfdde 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp @@ -21,6 +21,7 @@ #include "Constant.hpp" #include "ConstantUtil.hpp" // lookup #include "LinearAlg.hpp" +#include "MPILinearAlg.hpp" #include "Observables.hpp" #include "StateVectorCudaMPI.hpp" #include "Util.hpp" @@ -283,6 +284,7 @@ class SparseHamiltonianMPI final : public SparseHamiltonianBase { } using CFP_t = typename StateVectorT::CFP_t; + /* // Distribute sparse matrix across multi-nodes/multi-gpus size_t num_rows = size_t{1} << sv.getTotalNumQubits(); size_t local_num_rows = size_t{1} << sv.getNumLocalQubits(); @@ -304,23 +306,33 @@ class SparseHamiltonianMPI final : public SparseHamiltonianBase { } mpi_manager.Barrier(); + */ auto device_id = sv.getDataBuffer().getDevTag().getDeviceID(); auto stream_id = sv.getDataBuffer().getDevTag().getStreamID(); - cusparseHandle_t handle = sv.getCusparseHandle(); + // cusparseHandle_t handle = sv.getCusparseHandle(); const size_t length_local = size_t{1} << sv.getNumLocalQubits(); std::unique_ptr> d_sv_prime = std::make_unique>(length_local, device_id, stream_id, true); - std::unique_ptr> d_tmp = - std::make_unique>(length_local, device_id, - stream_id, true); d_sv_prime->zeroInit(); PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); mpi_manager.Barrier(); + cuUtil::SparseMV_cuSparseMPI( + mpi_manager, length_local, this->offsets_.data(), + static_cast(this->offsets_.size()), this->indices_.data(), + this->data_.data(), const_cast(sv.getData()), + d_sv_prime->getData(), device_id, stream_id, + sv.getCusparseHandle()); + + /* + std::unique_ptr> d_tmp = + std::make_unique>(length_local, device_id, + stream_id, true); + for (size_t i = 0; i < mpi_manager.getSize(); i++) { size_t color = 0; auto &localCSRMatrix = localCSRMatVector[i]; @@ -329,11 +341,11 @@ class SparseHamiltonianMPI final : public SparseHamiltonianBase { color = 1; SparseMV_cuSparse( localCSRMatrix.getCsrOffsets().data(), - localCSRMatrix.getCsrOffsets().size(), + static_cast(localCSRMatrix.getCsrOffsets().size()), localCSRMatrix.getColumns().data(), localCSRMatrix.getValues().data(), - localCSRMatrix.getValues().size(), sv.getData(), - d_sv_prime->getData(), device_id, stream_id, handle); + static_cast(localCSRMatrix.getValues().size()), + sv.getData(), d_sv_prime->getData(), device_id, stream_id, handle); } PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); mpi_manager.Barrier(); @@ -365,6 +377,8 @@ class SparseHamiltonianMPI final : public SparseHamiltonianBase { PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); mpi_manager.Barrier(); } + */ + sv.CopyGpuDataToGpuIn(d_sv_prime->getData(), d_sv_prime->getLength()); mpi_manager.Barrier(); } diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/Test_ObservablesGPU.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/Test_ObservablesGPU.cpp index 0d8bd7d388..816b285695 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/Test_ObservablesGPU.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/Test_ObservablesGPU.cpp @@ -257,3 +257,31 @@ TEMPLATE_PRODUCT_TEST_CASE("Hamiltonian::ApplyInPlace", "[Observables]", } } } + +TEMPLATE_PRODUCT_TEST_CASE("SparseHamiltonian::ApplyInPlace", "[Observables]", + (StateVectorCudaManaged), (float, double)) { + using StateVectorT = TestType; + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + + const std::size_t num_qubits = 3; + std::mt19937 re{1337}; + + auto sparseH = SparseHamiltonian::create( + {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, + {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, {0, 1, 2}); + + auto init_state = createRandomStateVectorData(re, num_qubits); + + StateVectorT state_vector(init_state.data(), init_state.size()); + + sparseH->applyInPlace(state_vector); + + std::reverse(init_state.begin(), init_state.end()); + + REQUIRE(isApproxEqual(state_vector.getDataVector().data(), + state_vector.getDataVector().size(), + init_state.data(), init_state.size())); +} \ No newline at end of file diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/mpi/Test_ObservablesGPUMPI.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/mpi/Test_ObservablesGPUMPI.cpp index e757fb4b84..b960536bba 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/mpi/Test_ObservablesGPUMPI.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/mpi/Test_ObservablesGPUMPI.cpp @@ -287,4 +287,79 @@ TEMPLATE_PRODUCT_TEST_CASE("Observables::HermitianHasherMPI", "[Observables]", CHECK(ham_1->getObsName() == res1.str()); CHECK(ham_2->getObsName() == res2.str()); } +} + +TEMPLATE_PRODUCT_TEST_CASE("SparseHamiltonian::ApplyInPlace", "[Observables]", + (StateVectorCudaMPI), (float, double)) { + using StateVectorT = TestType; + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + MPIManager mpi_manager(MPI_COMM_WORLD); + + const std::size_t num_qubits = 3; + std::mt19937 re{1337}; + + auto sparseH = SparseHamiltonianMPI::create( + {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, + {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, {0, 1, 2}); + + size_t mpi_buffersize = 1; + size_t nGlobalIndexBits = + std::bit_width(static_cast(mpi_manager.getSize())) - 1; + size_t nLocalIndexBits = num_qubits - nGlobalIndexBits; + size_t subSvLength = 1 << nLocalIndexBits; + size_t svLength = 1 << num_qubits; + + mpi_manager.Barrier(); + std::vector expected_sv(svLength); + std::vector local_state(subSvLength); + + auto init_state = createRandomStateVectorData(re, num_qubits); + + mpi_manager.Scatter(init_state.data(), local_state.data(), subSvLength, 0); + mpi_manager.Barrier(); + + int nDevices = 0; + cudaGetDeviceCount(&nDevices); + int deviceId = mpi_manager.getRank() % nDevices; + cudaSetDevice(deviceId); + DevTag dt_local(deviceId, 0); + mpi_manager.Barrier(); + + StateVectorT sv_mpi(mpi_manager, dt_local, mpi_buffersize, nGlobalIndexBits, + nLocalIndexBits); + + sv_mpi.CopyHostDataToGpu(local_state, false); + + sparseH->applyInPlace(sv_mpi); + + std::reverse(init_state.begin(), init_state.end()); + mpi_manager.Scatter(init_state.data(), expected_sv.data(), subSvLength, 0); + mpi_manager.Barrier(); + + CHECK(sv_mpi.getDataVector()[0].real() == + Approx(expected_sv[0].real()).epsilon(1e-5)); + CHECK(sv_mpi.getDataVector()[1].real() == + Approx(expected_sv[1].real()).epsilon(1e-5)); + CHECK(sv_mpi.getDataVector()[2].real() == + Approx(expected_sv[2].real()).epsilon(1e-5)); + CHECK(sv_mpi.getDataVector()[3].real() == + Approx(expected_sv[3].real()).epsilon(1e-5)); + + CHECK(sv_mpi.getDataVector()[0].imag() == + Approx(expected_sv[0].imag()).epsilon(1e-5)); + CHECK(sv_mpi.getDataVector()[1].imag() == + Approx(expected_sv[1].imag()).epsilon(1e-5)); + CHECK(sv_mpi.getDataVector()[2].imag() == + Approx(expected_sv[2].imag()).epsilon(1e-5)); + CHECK(sv_mpi.getDataVector()[3].imag() == + Approx(expected_sv[3].imag()).epsilon(1e-5)); + + /* + REQUIRE(isApproxEqual(sv_mpi.getDataVector().data(), + sv_mpi.getDataVector().size(), + expected_sv.data(), expected_sv.size())); + */ } \ No newline at end of file diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/utils/LinearAlg.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/utils/LinearAlg.hpp index f70d4ea9f2..1827bfffa8 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/utils/LinearAlg.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/utils/LinearAlg.hpp @@ -307,20 +307,20 @@ inline SharedCusparseHandle make_shared_cusparse_handle() { * @param handle cuSparse handle. */ template -inline void SparseMV_cuSparse(const index_type *csrOffsets_ptr, - const index_type csrOffsets_size, - const index_type *columns_ptr, - const std::complex *values_ptr, - const index_type numNNZ, CFP_t *X, CFP_t *Y, - DevTypeID device_id, cudaStream_t stream_id, - cusparseHandle_t handle) { - const int64_t num_rows = static_cast( +inline void +SparseMV_cuSparse(const index_type *csrOffsets_ptr, + const int64_t csrOffsets_size, const index_type *columns_ptr, + const std::complex *values_ptr, + const int64_t numNNZ, CFP_t *X, CFP_t *Y, DevTypeID device_id, + cudaStream_t stream_id, cusparseHandle_t handle) { + + const int64_t num_rows = csrOffsets_size - - 1); // int64_t is required for num_rows by cusparseCreateCsr - const int64_t num_cols = static_cast( - num_rows); // int64_t is required for num_cols by cusparseCreateCsr - const int64_t nnz = static_cast( - numNNZ); // int64_t is required for nnz by cusparseCreateCsr + 1; // int64_t is required for num_rows by cusparseCreateCsr + const int64_t num_cols = + num_rows; // int64_t is required for num_cols by cusparseCreateCsr + const int64_t nnz = + numNNZ; // int64_t is required for nnz by cusparseCreateCsr const CFP_t alpha = {1.0, 0.0}; const CFP_t beta = {0.0, 0.0}; @@ -338,13 +338,15 @@ inline void SparseMV_cuSparse(const index_type *csrOffsets_ptr, d_values.CopyHostDataToGpu(values_ptr, d_values.getLength(), false); cudaDataType_t data_type; - cusparseIndexType_t compute_type = CUSPARSE_INDEX_64I; + cusparseIndexType_t compute_type; if constexpr (std::is_same_v || std::is_same_v) { data_type = CUDA_C_64F; + compute_type = CUSPARSE_INDEX_64I; } else { data_type = CUDA_C_32F; + compute_type = CUSPARSE_INDEX_32I; } // CUSPARSE APIs @@ -394,8 +396,7 @@ inline void SparseMV_cuSparse(const index_type *csrOffsets_ptr, /* cusparseSpMVAlg_t */ CUSPARSE_SPMV_ALG_DEFAULT, /* size_t* */ &bufferSize)); - DataBuffer dBuffer{bufferSize, device_id, stream_id, - true}; + DataBuffer dBuffer{bufferSize, device_id, stream_id, true}; // execute SpMV PL_CUSPARSE_IS_SUCCESS(cusparseSpMV( @@ -439,19 +440,19 @@ inline void SparseMV_cuSparse(const index_type *csrOffsets_ptr, */ template inline void SparseMV_cuSparse(const index_type *csrOffsets_ptr, - const index_type csrOffsets_size, + const int64_t csrOffsets_size, const index_type *columns_ptr, const std::complex *values_ptr, - const index_type numNNZ, const CFP_t *X, CFP_t *Y, + const int64_t numNNZ, const CFP_t *X, CFP_t *Y, DevTypeID device_id, cudaStream_t stream_id, cusparseHandle_t handle) { - const int64_t num_rows = static_cast( + const int64_t num_rows = csrOffsets_size - - 1); // int64_t is required for num_rows by cusparseCreateCsr - const int64_t num_cols = static_cast( - num_rows); // int64_t is required for num_cols by cusparseCreateCsr - const int64_t nnz = static_cast( - numNNZ); // int64_t is required for nnz by cusparseCreateCsr + 1; // int64_t is required for num_rows by cusparseCreateCsr + const int64_t num_cols = + num_rows; // int64_t is required for num_cols by cusparseCreateCsr + const int64_t nnz = + numNNZ; // int64_t is required for nnz by cusparseCreateCsr const CFP_t alpha = {1.0, 0.0}; const CFP_t beta = {0.0, 0.0}; diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/utils/MPILinearAlg.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/utils/MPILinearAlg.hpp index 4b2e905b8e..cd2afd426b 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/utils/MPILinearAlg.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/utils/MPILinearAlg.hpp @@ -45,15 +45,15 @@ namespace Pennylane::LightningGPU::Util { template inline void SparseMV_cuSparseMPI( MPIManager &mpi_manager, const size_t &length_local, - const index_type *csrOffsets_ptr, const index_type csrOffsets_size, + const index_type *csrOffsets_ptr, const int64_t csrOffsets_size, const index_type *columns_ptr, const std::complex *values_ptr, CFP_t *X, CFP_t *Y, DevTypeID device_id, cudaStream_t stream_id, cusparseHandle_t handle) { std::vector>> csrmatrix_blocks; if (mpi_manager.getRank() == 0) { csrmatrix_blocks = splitCSRMatrix( - mpi_manager, csrOffsets_size - 1, csrOffsets_ptr, columns_ptr, - values_ptr); + mpi_manager, static_cast(csrOffsets_size - 1), + csrOffsets_ptr, columns_ptr, values_ptr); } mpi_manager.Barrier(); @@ -79,11 +79,11 @@ inline void SparseMV_cuSparseMPI( color = 1; SparseMV_cuSparse( localCSRMatrix.getCsrOffsets().data(), - localCSRMatrix.getCsrOffsets().size(), + static_cast(localCSRMatrix.getCsrOffsets().size()), localCSRMatrix.getColumns().data(), localCSRMatrix.getValues().data(), - localCSRMatrix.getValues().size(), X, d_res_per_block.getData(), - device_id, stream_id, handle); + static_cast(localCSRMatrix.getValues().size()), X, + d_res_per_block.getData(), device_id, stream_id, handle); } PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/utils/tests/Test_LinearAlgebra.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/utils/tests/Test_LinearAlgebra.cpp index cd1dd3937a..a2b35d0742 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/utils/tests/Test_LinearAlgebra.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/utils/tests/Test_LinearAlgebra.cpp @@ -40,6 +40,8 @@ TEMPLATE_TEST_CASE("Linear Algebra::SparseMV", "[Linear Algebra]", float, using StateVectorT = StateVectorCudaManaged; using ComplexT = StateVectorT::ComplexT; using CFP_t = StateVectorT::CFP_t; + using IdxT = typename std::conditional::value, + int32_t, int64_t>::type; std::size_t num_qubits = 3; std::size_t data_size = exp2(num_qubits); @@ -52,9 +54,9 @@ TEMPLATE_TEST_CASE("Linear Algebra::SparseMV", "[Linear Algebra]", float, {0.2, -0.1}, {-0.1, 0.2}, {0.2, 0.1}, {0.1, 0.2}, {0.7, -0.2}, {-0.1, 0.6}, {0.6, 0.1}, {0.2, 0.7}}; - std::vector indptr = {0, 2, 4, 6, 8, 10, 12, 14, 16}; - std::vector indices = {0, 3, 1, 2, 1, 2, 0, 3, - 4, 7, 5, 6, 5, 6, 4, 7}; + std::vector indptr = {0, 2, 4, 6, 8, 10, 12, 14, 16}; + std::vector indices = {0, 3, 1, 2, 1, 2, 0, 3, + 4, 7, 5, 6, 5, 6, 4, 7}; std::vector values = { {1.0, 0.0}, {0.0, -1.0}, {1.0, 0.0}, {0.0, 1.0}, {0.0, -1.0}, {1.0, 0.0}, {0.0, 1.0}, {1.0, 0.0}, @@ -69,10 +71,10 @@ TEMPLATE_TEST_CASE("Linear Algebra::SparseMV", "[Linear Algebra]", float, SECTION("Testing sparse matrix vector product:") { std::vector result(data_size); - cuUtil::SparseMV_cuSparse( - indptr.data(), indptr.size(), indices.data(), values.data(), - values.size(), sv_x.getData(), sv_y.getData(), - sv_x.getDataBuffer().getDevTag().getDeviceID(), + cuUtil::SparseMV_cuSparse( + indptr.data(), static_cast(indptr.size()), indices.data(), + values.data(), static_cast(values.size()), sv_x.getData(), + sv_y.getData(), sv_x.getDataBuffer().getDevTag().getDeviceID(), sv_x.getDataBuffer().getDevTag().getStreamID(), sv_x.getCusparseHandle()); diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/utils/tests/mpi/Test_LinearAlgebraMPI.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/utils/tests/mpi/Test_LinearAlgebraMPI.cpp index 4a3892b4fd..41ef00c3d9 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/utils/tests/mpi/Test_LinearAlgebraMPI.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/utils/tests/mpi/Test_LinearAlgebraMPI.cpp @@ -40,6 +40,8 @@ TEMPLATE_TEST_CASE("Linear Algebra::SparseMV", "[Linear Algebra]", float, using StateVectorT = StateVectorCudaMPI; using ComplexT = StateVectorT::ComplexT; using CFP_t = StateVectorT::CFP_t; + using IdxT = typename std::conditional::value, + int32_t, int64_t>::type; MPIManager mpi_manager(MPI_COMM_WORLD); @@ -53,9 +55,9 @@ TEMPLATE_TEST_CASE("Linear Algebra::SparseMV", "[Linear Algebra]", float, {0.1, 0.2}, {0.7, -0.2}, {-0.1, 0.6}, {0.6, 0.1}, {0.2, 0.7}}; - std::vector indptr = {0, 2, 4, 6, 8, 10, 12, 14, 16}; - std::vector indices = {0, 3, 1, 2, 1, 2, 0, 3, - 4, 7, 5, 6, 5, 6, 4, 7}; + std::vector indptr = {0, 2, 4, 6, 8, 10, 12, 14, 16}; + std::vector indices = {0, 3, 1, 2, 1, 2, 0, 3, + 4, 7, 5, 6, 5, 6, 4, 7}; std::vector values = { {1.0, 0.0}, {0.0, -1.0}, {1.0, 0.0}, {0.0, 1.0}, {0.0, -1.0}, {1.0, 0.0}, {0.0, 1.0}, {1.0, 0.0}, @@ -93,9 +95,10 @@ TEMPLATE_TEST_CASE("Linear Algebra::SparseMV", "[Linear Algebra]", float, nGlobalIndexBits, nLocalIndexBits); sv_x.CopyHostDataToGpu(local_state, false); - cuUtil::SparseMV_cuSparseMPI( - mpi_manager, sv_x.getLength(), indptr.data(), indptr.size(), - indices.data(), values.data(), sv_x.getData(), sv_y.getData(), + cuUtil::SparseMV_cuSparseMPI( + mpi_manager, sv_x.getLength(), indptr.data(), + static_cast(indptr.size()), indices.data(), values.data(), + sv_x.getData(), sv_y.getData(), sv_x.getDataBuffer().getDevTag().getDeviceID(), sv_x.getDataBuffer().getDevTag().getStreamID(), sv_x.getCusparseHandle()); diff --git a/pennylane_lightning/core/src/utils/TestHelpers.hpp b/pennylane_lightning/core/src/utils/TestHelpers.hpp index 605e584bcb..af1a46618a 100644 --- a/pennylane_lightning/core/src/utils/TestHelpers.hpp +++ b/pennylane_lightning/core/src/utils/TestHelpers.hpp @@ -429,7 +429,8 @@ void write_CSR_vectors(std::vector &row_map, const ComplexT SC_ONE = 1.0; row_map.resize(numRows + 1); - for (IndexT rowIdx = 1; rowIdx < (IndexT)row_map.size(); ++rowIdx) { + for (IndexT rowIdx = 1; rowIdx < static_cast(row_map.size()); + ++rowIdx) { row_map[rowIdx] = row_map[rowIdx - 1] + 3; }; const IndexT numNNZ = row_map[numRows]; @@ -437,6 +438,7 @@ void write_CSR_vectors(std::vector &row_map, entries.resize(numNNZ); values.resize(numNNZ); for (IndexT rowIdx = 0; rowIdx < numRows; ++rowIdx) { + size_t idx = row_map[rowIdx]; if (rowIdx == 0) { entries[0] = rowIdx; entries[1] = rowIdx + 1; @@ -446,21 +448,21 @@ void write_CSR_vectors(std::vector &row_map, values[1] = -SC_ONE; values[2] = -SC_ONE; } else if (rowIdx == numRows - 1) { - entries[row_map[rowIdx]] = 0; - entries[row_map[rowIdx] + 1] = rowIdx - 1; - entries[row_map[rowIdx] + 2] = rowIdx; + entries[idx] = 0; + entries[idx + 1] = rowIdx - 1; + entries[idx + 2] = rowIdx; - values[row_map[rowIdx]] = -SC_ONE; - values[row_map[rowIdx] + 1] = -SC_ONE; - values[row_map[rowIdx] + 2] = SC_ONE; + values[idx] = -SC_ONE; + values[idx + 1] = -SC_ONE; + values[idx + 2] = SC_ONE; } else { - entries[row_map[rowIdx]] = rowIdx - 1; - entries[row_map[rowIdx] + 1] = rowIdx; - entries[row_map[rowIdx] + 2] = rowIdx + 1; + entries[idx] = rowIdx - 1; + entries[idx + 1] = rowIdx; + entries[idx + 2] = rowIdx + 1; - values[row_map[rowIdx]] = -SC_ONE; - values[row_map[rowIdx] + 1] = SC_ONE; - values[row_map[rowIdx] + 2] = -SC_ONE; + values[idx] = -SC_ONE; + values[idx + 1] = SC_ONE; + values[idx + 2] = -SC_ONE; } } }; diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py index c8d65dd093..69fef4ed11 100644 --- a/tests/test_adjoint_jacobian.py +++ b/tests/test_adjoint_jacobian.py @@ -892,7 +892,10 @@ def circuit_ansatz(params, wires): qml.RX(params[29], wires=wires[1]) -@pytest.mark.skipif(not ld._CPP_BINARY_AVAILABLE, reason="Lightning binary required") +@pytest.mark.skipif( + device_name != "lightning.gpu" or not ld._CPP_BINARY_AVAILABLE, + reason="Lightning binary required", +) def test_tape_qchem(tol): """Tests the circuit Ansatz with a QChem Hamiltonian produces correct results""" @@ -915,7 +918,10 @@ def circuit(params): assert np.allclose(qml.grad(circuit_ld)(params), qml.grad(circuit_dq)(params), tol) -@pytest.mark.skipif(not ld._CPP_BINARY_AVAILABLE, reason="Lightning binary required") +@pytest.mark.skipif( + device_name != "lightning.gpu" or not ld._CPP_BINARY_AVAILABLE, + reason="Lightning binary required", +) def test_tape_qchem_sparse(tol): """Tests the circuit Ansatz with a QChem Hamiltonian produces correct results""" From c06593d11db622a9cc1b2dcf42a7ee03edd8280f Mon Sep 17 00:00:00 2001 From: Shuli <08cnbj@gmail.com> Date: Sat, 21 Oct 2023 04:19:40 +0000 Subject: [PATCH 11/43] Trigger CI From 1cbc50fde92e1199ffb3ae9ad5cbe6634729e52e Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 13:47:52 +0000 Subject: [PATCH 12/43] Fix python bindings LGPU idxT --- pennylane_lightning/core/src/bindings/Bindings.hpp | 14 +++++++++----- .../lightning_gpu/observables/ObservablesGPU.hpp | 11 +++++++---- 2 files changed, 16 insertions(+), 9 deletions(-) diff --git a/pennylane_lightning/core/src/bindings/Bindings.hpp b/pennylane_lightning/core/src/bindings/Bindings.hpp index d352b1d93c..bbd9f8c687 100644 --- a/pennylane_lightning/core/src/bindings/Bindings.hpp +++ b/pennylane_lightning/core/src/bindings/Bindings.hpp @@ -413,7 +413,7 @@ template void registerObservables(py::module_ &m) { #ifdef _ENABLE_PLGPU class_name = "SparseHamiltonianC" + bitsize; - using SpIDX = typename SparseHamiltonian::IdxT; + // using SpIDX = typename SparseHamiltonian::IdxT; py::class_, std::shared_ptr>, Observable>(m, class_name.c_str(), @@ -425,15 +425,19 @@ template void registerObservables(py::module_ &m) { const auto *data_ptr = static_cast(buffer_data.ptr); const py::buffer_info buffer_indices = indices.request(); - const auto *indices_ptr = static_cast(buffer_indices.ptr); + const auto *indices_ptr = + static_cast(buffer_indices.ptr); const py::buffer_info buffer_offsets = offsets.request(); - const auto *offsets_ptr = static_cast(buffer_offsets.ptr); + const auto *offsets_ptr = + static_cast(buffer_offsets.ptr); return SparseHamiltonian{ std::vector({data_ptr, data_ptr + data.size()}), - std::vector({indices_ptr, indices_ptr + indices.size()}), - std::vector({offsets_ptr, offsets_ptr + offsets.size()}), + std::vector( + {indices_ptr, indices_ptr + indices.size()}), + std::vector( + {offsets_ptr, offsets_ptr + offsets.size()}), wires}; })) .def("__repr__", &SparseHamiltonian::getObsName) diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp index c4042a938c..20425c4c2b 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp @@ -254,8 +254,8 @@ class SparseHamiltonian final : public SparseHamiltonianBase { * @param wires Argument to construct wires */ static auto create(std::initializer_list data, - std::initializer_list indices, - std::initializer_list offsets, + std::initializer_list indices, + std::initializer_list offsets, std::initializer_list wires) -> std::shared_ptr> { return std::shared_ptr>( @@ -286,11 +286,14 @@ class SparseHamiltonian final : public SparseHamiltonianBase { std::make_unique>(length, device_id, stream_id, true); + std::vector offsets_cp(this->offsets_.begin(), + this->offsets_.end()); + std::vector indices_cp(this->indices_.begin(), + this->indices_.end()); SparseMV_cuSparse( - this->offsets_.data(), this->offsets_.size(), this->indices_.data(), + offsets_cp.data(), offsets_cp.size(), indices_cp.data(), this->data_.data(), this->data_.size(), sv.getData(), d_sv_prime->getData(), device_id, stream_id, handle); - sv.updateData(std::move(d_sv_prime)); } }; From 6956e7828ec973326ec4b1bff1322519553676f5 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 14:40:41 +0000 Subject: [PATCH 13/43] Fix serial tests and update changelog. --- .github/CHANGELOG.md | 3 + pennylane_lightning/core/_serialize.py | 5 +- .../observables/ObservablesGPU.hpp | 2 +- .../observables/ObservablesKokkos.hpp | 2 +- tests/test_serialize.py | 57 ++++++++++--------- 5 files changed, 40 insertions(+), 29 deletions(-) diff --git a/.github/CHANGELOG.md b/.github/CHANGELOG.md index 2288658d22..9ef71c78df 100644 --- a/.github/CHANGELOG.md +++ b/.github/CHANGELOG.md @@ -2,6 +2,9 @@ ### New features since last release +* Add `SparseHamiltonian` support for Lightning-Kokkos. + [(#527)] (https://github.com/PennyLaneAI/pennylane-lightning/pull/527) + * Integrate python/pybind layer of distributed Lightning-GPU into the Lightning monorepo with python unit tests. [(#518)] (https://github.com/PennyLaneAI/pennylane-lightning/pull/518) diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index 28e230f5ee..49d46768a7 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -254,7 +254,10 @@ def _ob(self, observable, wires_map): return self._tensor_ob(observable, wires_map) if observable.name == "Hamiltonian": return self._hamiltonian(observable, wires_map) - if observable.name == "SparseHamiltonian": + if ( + self.device_name in ["lightning.gpu", "lightning.kokkos"] + and observable.name == "SparseHamiltonian" + ): return self._sparse_hamiltonian(observable, wires_map) if isinstance(observable, (PauliX, PauliY, PauliZ, Identity, Hadamard)): return self._named_obs(observable, wires_map) diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp index bbab9a466a..6d0d0a94e7 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp @@ -222,7 +222,7 @@ class SparseHamiltonian final : public SparseHamiltonianBase { using PrecisionT = typename StateVectorT::PrecisionT; using ComplexT = typename StateVectorT::ComplexT; // cuSparse required index type - using IdxT = BaseType::IdxT; + using IdxT = typename BaseType::IdxT; /** * @brief Create a SparseHamiltonian from data, indices and offsets in CSR diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp index b84c0c800a..c3fae6b3ea 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp @@ -212,7 +212,7 @@ class SparseHamiltonian final : public SparseHamiltonianBase { public: using PrecisionT = typename StateVectorT::PrecisionT; using ComplexT = typename StateVectorT::ComplexT; - using IdxT = BaseType::IdxT; + using IdxT = typename BaseType::IdxT; /** * @brief Create a SparseHamiltonian from data, indices and offsets in CSR diff --git a/tests/test_serialize.py b/tests/test_serialize.py index a1e2d74323..92938b3406 100644 --- a/tests/test_serialize.py +++ b/tests/test_serialize.py @@ -47,6 +47,8 @@ TensorProdObsC128, HamiltonianC64, HamiltonianC128, + SparseHamiltonianC64, + SparseHamiltonianC128, ) else: from pennylane_lightning.lightning_qubit_ops.observables import ( @@ -68,38 +70,41 @@ def test_wrong_device_name(): QuantumScriptSerializer("thunder.qubit") -@pytest.mark.parametrize( - "obs,obs_type", - [ - (qml.PauliZ(0), NamedObsC128), - (qml.PauliZ(0) @ qml.PauliZ(1), TensorProdObsC128), - (qml.Hadamard(0), NamedObsC128), - (qml.Hermitian(np.eye(2), wires=0), HermitianObsC128), - ( - qml.PauliZ(0) @ qml.Hadamard(1) @ (0.1 * (qml.PauliZ(2) + qml.PauliX(3))), - HamiltonianC128, - ), - ( - ( - qml.Hermitian(np.eye(2), wires=0) - @ qml.Hermitian(np.eye(2), wires=1) - @ qml.Projector([0], wires=2) - ), - TensorProdObsC128, - ), +obs_list = [ + (qml.PauliZ(0), NamedObsC128), + (qml.PauliZ(0) @ qml.PauliZ(1), TensorProdObsC128), + (qml.Hadamard(0), NamedObsC128), + (qml.Hermitian(np.eye(2), wires=0), HermitianObsC128), + ( + qml.PauliZ(0) @ qml.Hadamard(1) @ (0.1 * (qml.PauliZ(2) + qml.PauliX(3))), + HamiltonianC128, + ), + ( ( - qml.PauliZ(0) @ qml.Hermitian(np.eye(2), wires=1) @ qml.Projector([0], wires=2), - TensorProdObsC128, + qml.Hermitian(np.eye(2), wires=0) + @ qml.Hermitian(np.eye(2), wires=1) + @ qml.Projector([0], wires=2) ), - (qml.Projector([0], wires=0), HermitianObsC128), - (qml.Hamiltonian([1], [qml.PauliZ(0)]), HamiltonianC128), - (qml.sum(qml.Hadamard(0), qml.PauliX(1)), HermitianObsC128), + TensorProdObsC128, + ), + ( + qml.PauliZ(0) @ qml.Hermitian(np.eye(2), wires=1) @ qml.Projector([0], wires=2), + TensorProdObsC128, + ), + (qml.Projector([0], wires=0), HermitianObsC128), + (qml.Hamiltonian([1], [qml.PauliZ(0)]), HamiltonianC128), + (qml.sum(qml.Hadamard(0), qml.PauliX(1)), HermitianObsC128), +] +if device_name in ["lightning.gpu", "lightning.kokkos"]: + obs_list += [ ( qml.SparseHamiltonian(qml.Hamiltonian([1], [qml.PauliZ(0)]).sparse_matrix(), wires=[0]), SparseHamiltonianC128, ), - ], -) + ] + + +@pytest.mark.parametrize("obs,obs_type", obs_list) def test_obs_returns_expected_type(obs, obs_type): """Tests that observables get serialized to the expected type.""" assert isinstance( From b256ead746bbbaa1240596e47e9d52a2485e15b0 Mon Sep 17 00:00:00 2001 From: Shuli <08cnbj@gmail.com> Date: Mon, 23 Oct 2023 14:52:49 +0000 Subject: [PATCH 14/43] add more unit tests for sparseH base class --- pennylane_lightning/core/_serialize.py | 2 +- .../core/src/observables/Observables.hpp | 12 +- .../observables/tests/Test_Observables.cpp | 39 ++++++ .../tests/mpi/Test_ObservablesMPI.cpp | 123 +++++++++--------- .../tests/mpi/Test_ObservablesGPUMPI.cpp | 27 +--- 5 files changed, 118 insertions(+), 85 deletions(-) diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index 9cb603ba7e..eb05f41709 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -237,7 +237,7 @@ def _ob(self, observable, wires_map): return self._tensor_ob(observable, wires_map) if observable.name == "Hamiltonian": return self._hamiltonian(observable, wires_map) - if observable.name == "SparseHamiltonian": + if observable.name == "SparseHamiltonian" and self.device_name != "lightning.qubit": return self._sparsehamiltonian(observable, wires_map) if isinstance(observable, (PauliX, PauliY, PauliZ, Identity, Hadamard)): return self._named_obs(observable, wires_map) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index d4e9273723..ae062cc531 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -425,9 +425,13 @@ class SparseHamiltonianBase : public Observable { using PrecisionT = typename StateVectorT::PrecisionT; using ComplexT = typename StateVectorT::ComplexT; // cuSparse required index type +#ifdef _ENABLE_PLGPU using IdxT = typename std::conditional::value, int32_t, int64_t>::type; +#else + using IdxT = std::size_t; +#endif protected: std::vector data_; @@ -490,6 +494,12 @@ class SparseHamiltonianBase : public Observable { std::move(arg4)}); } + void applyInPlace([[maybe_unused]] StateVectorT &sv) const override { + PL_ABORT("For SparseHamiltonian Observables, the applyInPlace method " + "must be " + "defined at the backend level."); + } + [[nodiscard]] auto getObsName() const -> std::string override { using Pennylane::Util::operator<<; std::ostringstream ss; @@ -498,7 +508,7 @@ class SparseHamiltonianBase : public Observable { // Note that for LGPU backend, ComplexT is std::complex as of 0.33 // release Need to revisit it once we set ComplexT as cuComplex // later - ss << "{" << d.real() << ", " << d.real() << "}," + ss << "{" << d.real() << ", " << d.imag() << "}," << "\n"; ss << ",\n'indices' : \n"; for (const auto &i : indices_) diff --git a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp index 2c4e3b60da..b6ff4758d6 100644 --- a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp +++ b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp @@ -464,3 +464,42 @@ TEST_CASE("Methods implemented in the HamiltonianBase class", testHamiltonianBase(); } } + +template void testSparseHamiltonianBase() { + if constexpr (!std::is_same_v) { + using StateVectorT = typename TypeList::Type; + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + + const std::size_t num_qubits = 3; + std::mt19937 re{1337}; + + DYNAMIC_SECTION("applyInPlace must fail - " + << StateVectorToName::name) { + + auto sparseH = SparseHamiltonianBase::create( + {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, + {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, + {0, 1, 2}); + + auto init_state = + createRandomStateVectorData(re, num_qubits); + + StateVectorT state_vector(init_state.data(), init_state.size()); + + REQUIRE_THROWS_AS(sparseH->applyInPlace(state_vector), + LightningException); + } + + testSparseHamiltonianBase(); + } +} + +TEST_CASE("Methods implemented in the SparseHamiltonianBase class", + "[SparseHamiltonianBase]") { + if constexpr (BACKEND_FOUND) { + testSparseHamiltonianBase(); + } +} \ No newline at end of file diff --git a/pennylane_lightning/core/src/observables/tests/mpi/Test_ObservablesMPI.cpp b/pennylane_lightning/core/src/observables/tests/mpi/Test_ObservablesMPI.cpp index dee6e37532..c7a0e1923e 100644 --- a/pennylane_lightning/core/src/observables/tests/mpi/Test_ObservablesMPI.cpp +++ b/pennylane_lightning/core/src/observables/tests/mpi/Test_ObservablesMPI.cpp @@ -243,51 +243,6 @@ template void testTensorProdObsBase() { REQUIRE(ob1 != ob4); REQUIRE(ob1 != ob5); } - - /* - DYNAMIC_SECTION("Tensor product applies to a statevector correctly" - << StateVectorMPIToName::name) { - using VectorT = TestVector; - - auto obs = TensorProdObsT{ - std::make_shared("PauliX", std::vector{0}), - std::make_shared("PauliX", std::vector{2}), - }; - - SECTION("Test using |1+0>") { - VectorT st_data = - createProductState("1+0"); - - StateVectorT state_vector(st_data.data(), st_data.size()); - - obs.applyInPlace(state_vector); - - VectorT expected = - createProductState("0+1"); - - REQUIRE(isApproxEqual(state_vector.getDataVector().data(), - state_vector.getDataVector().size(), - expected.data(), expected.size())); - } - - SECTION("Test using |+-01>") { - VectorT st_data = - createProductState("+-01"); - - StateVectorT state_vector(st_data.data(), st_data.size()); - - obs.applyInPlace(state_vector); - - VectorT expected = - createProductState("+-11"); - - REQUIRE(isApproxEqual(state_vector.getDataVector().data(), - state_vector.getDataVector().size(), - expected.data(), expected.size())); - } - } - */ - testTensorProdObsBase(); } } @@ -422,20 +377,6 @@ template void testHamiltonianBase() { REQUIRE(ham1->getWires() == std::vector{0, 5, 9}); } - - /* - DYNAMIC_SECTION("applyInPlace must fail - " - << StateVectorMPIToName::name) { - auto ham = - HamiltonianT::create({PrecisionT{1.0}, h, h}, {zz, x1, x2}); - auto st_data = createZeroState(2); - - StateVectorT state_vector(st_data.data(), st_data.size()); - - REQUIRE_THROWS_AS(ham->applyInPlace(state_vector), - LightningException); - } - */ } testHamiltonianBase(); } @@ -447,3 +388,67 @@ TEST_CASE("Methods implemented in the HamiltonianBase class", testHamiltonianBase(); } } + +template void testSparseHamiltonianBase() { + if constexpr (!std::is_same_v) { + using StateVectorT = typename TypeList::Type; + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + + const std::size_t num_qubits = 3; + std::mt19937 re{1337}; + + MPIManager mpi_manager(MPI_COMM_WORLD); + + size_t mpi_buffersize = 1; + size_t nGlobalIndexBits = + std::bit_width(static_cast(mpi_manager.getSize())) - 1; + size_t nLocalIndexBits = num_qubits - nGlobalIndexBits; + size_t subSvLength = 1 << nLocalIndexBits; + + int nDevices = 0; + cudaGetDeviceCount(&nDevices); + int deviceId = mpi_manager.getRank() % nDevices; + cudaSetDevice(deviceId); + DevTag dt_local(deviceId, 0); + mpi_manager.Barrier(); + + std::vector expected_sv(subSvLength); + std::vector local_state(subSvLength); + + auto init_state = + createRandomStateVectorData(re, num_qubits); + + mpi_manager.Scatter(init_state.data(), local_state.data(), subSvLength, + 0); + mpi_manager.Barrier(); + + DYNAMIC_SECTION("applyInPlace must fail - " + << StateVectorMPIToName::name) { + + auto sparseH = SparseHamiltonianBase::create( + {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, + {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, + {0, 1, 2}); + + StateVectorT sv_mpi(mpi_manager, dt_local, mpi_buffersize, + nGlobalIndexBits, nLocalIndexBits); + + sv_mpi.CopyHostDataToGpu(local_state, false); + + REQUIRE_THROWS_AS(sparseH->applyInPlace(sv_mpi), + LightningException); + } + + testSparseHamiltonianBase(); + } +} + +TEST_CASE("Methods implemented in the SparseHamiltonianBase class", + "[SparseHamiltonianBase]") { + if constexpr (BACKEND_FOUND) { + testSparseHamiltonianBase(); + } +} diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/mpi/Test_ObservablesGPUMPI.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/mpi/Test_ObservablesGPUMPI.cpp index b960536bba..c28bcbee1e 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/mpi/Test_ObservablesGPUMPI.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/mpi/Test_ObservablesGPUMPI.cpp @@ -310,10 +310,9 @@ TEMPLATE_PRODUCT_TEST_CASE("SparseHamiltonian::ApplyInPlace", "[Observables]", std::bit_width(static_cast(mpi_manager.getSize())) - 1; size_t nLocalIndexBits = num_qubits - nGlobalIndexBits; size_t subSvLength = 1 << nLocalIndexBits; - size_t svLength = 1 << num_qubits; mpi_manager.Barrier(); - std::vector expected_sv(svLength); + std::vector expected_sv(subSvLength); std::vector local_state(subSvLength); auto init_state = createRandomStateVectorData(re, num_qubits); @@ -339,27 +338,7 @@ TEMPLATE_PRODUCT_TEST_CASE("SparseHamiltonian::ApplyInPlace", "[Observables]", mpi_manager.Scatter(init_state.data(), expected_sv.data(), subSvLength, 0); mpi_manager.Barrier(); - CHECK(sv_mpi.getDataVector()[0].real() == - Approx(expected_sv[0].real()).epsilon(1e-5)); - CHECK(sv_mpi.getDataVector()[1].real() == - Approx(expected_sv[1].real()).epsilon(1e-5)); - CHECK(sv_mpi.getDataVector()[2].real() == - Approx(expected_sv[2].real()).epsilon(1e-5)); - CHECK(sv_mpi.getDataVector()[3].real() == - Approx(expected_sv[3].real()).epsilon(1e-5)); - - CHECK(sv_mpi.getDataVector()[0].imag() == - Approx(expected_sv[0].imag()).epsilon(1e-5)); - CHECK(sv_mpi.getDataVector()[1].imag() == - Approx(expected_sv[1].imag()).epsilon(1e-5)); - CHECK(sv_mpi.getDataVector()[2].imag() == - Approx(expected_sv[2].imag()).epsilon(1e-5)); - CHECK(sv_mpi.getDataVector()[3].imag() == - Approx(expected_sv[3].imag()).epsilon(1e-5)); - - /* REQUIRE(isApproxEqual(sv_mpi.getDataVector().data(), - sv_mpi.getDataVector().size(), - expected_sv.data(), expected_sv.size())); - */ + sv_mpi.getDataVector().size(), expected_sv.data(), + expected_sv.size())); } \ No newline at end of file From 5813cf4a295d5c9cce03a16b9433119ed7641aba Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 14:54:17 +0000 Subject: [PATCH 15/43] Fix tidy & sparse adjoint test device name. --- pennylane_lightning/core/src/observables/Observables.hpp | 6 ++++-- tests/test_adjoint_jacobian.py | 2 +- 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index c44dfec97e..b44e0475b7 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -503,11 +503,13 @@ class SparseHamiltonianBase : public Observable { ss << "{" << d.real() << ", " << d.real() << "}," << "\n"; ss << ",\n'indices' : \n"; - for (const auto &i : indices_) + for (const auto &i : indices_) { ss << i; + } ss << ",\n'offsets' : \n"; - for (const auto &o : offsets_) + for (const auto &o : offsets_) { ss << o; + } ss << "\n}"; return ss.str(); } diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py index c8ba86150f..c9b338cdc3 100644 --- a/tests/test_adjoint_jacobian.py +++ b/tests/test_adjoint_jacobian.py @@ -990,7 +990,7 @@ def test_adjoint_SparseHamiltonian(returns): """Integration tests that compare to default.qubit for a large circuit containing parametrized operations and when using custom wire labels""" - dev_kokkos = qml.device("lightning.kokkos", wires=custom_wires) + dev_kokkos = qml.device(device_name, wires=custom_wires) dev_default = qml.device("default.qubit", wires=custom_wires) def circuit(params): From cc66546ef5b628b41f0f91eeaab892db5d8c2c7b Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 15:49:36 +0000 Subject: [PATCH 16/43] Fix tidy warning for sparse_ham. --- .../core/src/observables/Observables.hpp | 16 +++++---- .../observables/tests/Test_Observables.cpp | 33 ++++++++++++++----- .../lightning_kokkos/StateVectorKokkos.hpp | 19 ++++++++--- 3 files changed, 48 insertions(+), 20 deletions(-) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index 3ab1bb7fc9..d7cd3a1f8f 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -414,6 +414,7 @@ class HamiltonianBase : public Observable { } }; +// NOLINTBEGIN /** * @brief Sparse representation of SparseHamiltonian * @@ -480,7 +481,7 @@ class SparseHamiltonianBase : public Observable { * * @param arg1 Argument to construct data * @param arg2 Argument to construct indices - * @param arg3 Argument to construct ofsets + * @param arg3 Argument to construct offsets * @param arg4 Argument to construct wires */ static auto create(std::initializer_list arg1, @@ -504,19 +505,19 @@ class SparseHamiltonianBase : public Observable { using Pennylane::Util::operator<<; std::ostringstream ss; ss << "SparseHamiltonian: {\n'data' : \n"; - for (const auto &d : data_) - // Note that for LGPU backend, ComplexT is std::complex as of 0.33 + for (const auto &d : data_) { // Note that for LGPU backend, ComplexT is + // std::complex as of 0.33 // release Need to revisit it once we set ComplexT as cuComplex // later - ss << "{" << d.real() << ", " << d.imag() << "}," - << "\n"; + ss << "{" << d.real() << ", " << d.imag() << "}, "; + } ss << ",\n'indices' : \n"; for (const auto &i : indices_) { - ss << i; + ss << i << ", "; } ss << ",\n'offsets' : \n"; for (const auto &o : offsets_) { - ss << o; + ss << o << ", "; } ss << "\n}"; return ss.str(); @@ -528,5 +529,6 @@ class SparseHamiltonianBase : public Observable { return wires_; }; }; +// NOLINTEND } // namespace Pennylane::Observables \ No newline at end of file diff --git a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp index b6ff4758d6..12f25ef667 100644 --- a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp +++ b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp @@ -41,6 +41,7 @@ using Pennylane::Util::TestVector; #ifdef _ENABLE_PLQUBIT constexpr bool BACKEND_FOUND = true; +constexpr bool SPARSE_HAM_SUPPORTED = false; #include "TestHelpersStateVectors.hpp" // TestStateVectorBackends, StateVectorToName @@ -52,6 +53,7 @@ using namespace Pennylane::LightningQubit::Util; #elif _ENABLE_PLKOKKOS == 1 constexpr bool BACKEND_FOUND = true; +constexpr bool SPARSE_HAM_SUPPORTED = true; #include "TestHelpersStateVectors.hpp" // TestStateVectorBackends, StateVectorToName @@ -63,6 +65,7 @@ using namespace Pennylane::LightningKokkos::Util; #elif _ENABLE_PLGPU == 1 constexpr bool BACKEND_FOUND = true; +constexpr bool SPARSE_HAM_SUPPORTED = true; #include "TestHelpersStateVectors.hpp" @@ -74,6 +77,7 @@ using namespace Pennylane::LightningGPU::Util; #else constexpr bool BACKEND_FOUND = false; +constexpr bool SPARSE_HAM_SUPPORTED = false; using TestStateVectorBackends = Pennylane::Util::TypeList; template struct StateVectorToName {}; @@ -474,15 +478,28 @@ template void testSparseHamiltonianBase() { const std::size_t num_qubits = 3; std::mt19937 re{1337}; - DYNAMIC_SECTION("applyInPlace must fail - " + auto sparseH = SparseHamiltonianBase::create( + {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, + {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, {0, 1, 2}); + + DYNAMIC_SECTION("getObsName - " << StateVectorToName::name) { + REQUIRE(sparseH->getObsName() == + "SparseHamiltonian: {\n" + "'data' : \n" + "{1, 0}, {1, 0}, {1, 0}, {1, 0}, {1, 0}, {1, 0}, {1, 0}, " + "{1, 0}, ,\n" + "'indices' : \n" + "7, 6, 5, 4, 3, 2, 1, 0, ,\n" + "'offsets' : \n" + "0, 1, 2, 3, 4, 5, 6, 7, 8, \n" + "}"); + } - auto sparseH = SparseHamiltonianBase::create( - {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, - ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, - ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, - {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, - {0, 1, 2}); + DYNAMIC_SECTION("applyInPlace must fail - " + << StateVectorToName::name) { auto init_state = createRandomStateVectorData(re, num_qubits); @@ -499,7 +516,7 @@ template void testSparseHamiltonianBase() { TEST_CASE("Methods implemented in the SparseHamiltonianBase class", "[SparseHamiltonianBase]") { - if constexpr (BACKEND_FOUND) { + if constexpr (SPARSE_HAM_SUPPORTED) { testSparseHamiltonianBase(); } } \ No newline at end of file diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/StateVectorKokkos.hpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/StateVectorKokkos.hpp index a3f0951b24..02064e811f 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/StateVectorKokkos.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/StateVectorKokkos.hpp @@ -17,6 +17,7 @@ */ #pragma once +#include #include #include #include @@ -170,7 +171,7 @@ class StateVectorKokkos final * * @param num_qubits Number of qubits */ - StateVectorKokkos(ComplexT *hostdata_, size_t length, + StateVectorKokkos(ComplexT *hostdata_, std::size_t length, const Kokkos::InitializationSettings &kokkos_args = {}) : StateVectorKokkos(log2(length), kokkos_args) { PL_ABORT_IF_NOT(isPerfectPowerOf2(length), @@ -178,12 +179,20 @@ class StateVectorKokkos final HostToDevice(hostdata_, length); } + StateVectorKokkos(std::complex *hostdata_, std::size_t length, + const Kokkos::InitializationSettings &kokkos_args = {}) + : StateVectorKokkos(log2(length), kokkos_args) { + PL_ABORT_IF_NOT(isPerfectPowerOf2(length), + "The size of provided data must be a power of 2."); + HostToDevice(reinterpret_cast(hostdata_), length); + } + /** * @brief Create a new state vector from data on the host. * * @param num_qubits Number of qubits */ - StateVectorKokkos(const ComplexT *hostdata_, size_t length, + StateVectorKokkos(const ComplexT *hostdata_, std::size_t length, const Kokkos::InitializationSettings &kokkos_args = {}) : StateVectorKokkos(log2(length), kokkos_args) { PL_ABORT_IF_NOT(isPerfectPowerOf2(length), @@ -692,7 +701,7 @@ class StateVectorKokkos final * @param new_data data pointer to new data. * @param new_size size of underlying data storage. */ - void updateData(ComplexT *new_data, size_t new_size) { + void updateData(ComplexT *new_data, std::size_t new_size) { updateData(KokkosVector(new_data, new_size)); } @@ -744,7 +753,7 @@ class StateVectorKokkos final * @brief Copy data from the host space to the device space. * */ - inline void HostToDevice(ComplexT *sv, size_t length) { + inline void HostToDevice(ComplexT *sv, std::size_t length) { Kokkos::deep_copy(*data_, UnmanagedComplexHostView(sv, length)); } @@ -752,7 +761,7 @@ class StateVectorKokkos final * @brief Copy data from the device space to the host space. * */ - inline void DeviceToHost(ComplexT *sv, size_t length) const { + inline void DeviceToHost(ComplexT *sv, std::size_t length) const { Kokkos::deep_copy(UnmanagedComplexHostView(sv, length), *data_); } From dccfe3e1ac64216da93254a3f087471cf38fef73 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 16:45:42 +0000 Subject: [PATCH 17/43] Send backend-specific ops in respective modules. --- .../core/src/bindings/Bindings.hpp | 88 ++----------------- .../lightning_gpu/bindings/LGPUBindings.hpp | 67 ++++++++++++++ .../bindings/LKokkosBindings.hpp | 54 ++++++++++++ .../bindings/LQubitBindings.hpp | 9 ++ 4 files changed, 135 insertions(+), 83 deletions(-) diff --git a/pennylane_lightning/core/src/bindings/Bindings.hpp b/pennylane_lightning/core/src/bindings/Bindings.hpp index 511c4ef798..ac2fdd8b7f 100644 --- a/pennylane_lightning/core/src/bindings/Bindings.hpp +++ b/pennylane_lightning/core/src/bindings/Bindings.hpp @@ -287,7 +287,8 @@ void registerInfo(py::module_ &m) { * @tparam StateVectorT * @param m Pybind module */ -template void registerObservables(py::module_ &m) { +template +void registerBackendAgnosticObservables(py::module_ &m) { using PrecisionT = typename StateVectorT::PrecisionT; // Statevector's precision. using ComplexT = @@ -406,87 +407,6 @@ template void registerObservables(py::module_ &m) { return self == other_cast; }, "Compare two observables"); - -#ifdef _ENABLE_PLGPU - class_name = "SparseHamiltonianC" + bitsize; - using np_arr_sparse_ind = typename std::conditional< - std::is_same::value, - py::array_t, - py::array_t>::type; - using IdxT = typename SparseHamiltonian::IdxT; - py::class_, - std::shared_ptr>, - Observable>(m, class_name.c_str(), - py::module_local()) - .def(py::init([](const np_arr_c &data, const np_arr_sparse_ind &indices, - const np_arr_sparse_ind &offsets, - const std::vector &wires) { - const py::buffer_info buffer_data = data.request(); - const auto *data_ptr = static_cast(buffer_data.ptr); - - const py::buffer_info buffer_indices = indices.request(); - const auto *indices_ptr = - static_cast(buffer_indices.ptr); - - const py::buffer_info buffer_offsets = offsets.request(); - const auto *offsets_ptr = - static_cast(buffer_offsets.ptr); - - return SparseHamiltonian{ - std::vector({data_ptr, data_ptr + data.size()}), - std::vector({indices_ptr, indices_ptr + indices.size()}), - std::vector({offsets_ptr, offsets_ptr + offsets.size()}), - wires}; - })) - .def("__repr__", &SparseHamiltonian::getObsName) - .def("get_wires", &SparseHamiltonian::getWires, - "Get wires of observables") - .def( - "__eq__", - [](const SparseHamiltonian &self, - py::handle other) -> bool { - if (!py::isinstance>(other)) { - return false; - } - auto other_cast = other.cast>(); - return self == other_cast; - }, - "Compare two observables"); -#endif - -#if _ENABLE_PLKOKKOS == 1 - class_name = "SparseHamiltonianC" + bitsize; - py::class_, - std::shared_ptr>, - Observable>(m, class_name.c_str(), - py::module_local()) - .def(py::init([](const np_arr_c &data, - const std::vector &indices, - const std::vector &indptr, - const std::vector &wires) { - using ComplexT = typename StateVectorT::ComplexT; - const py::buffer_info buffer_data = data.request(); - const auto *data_ptr = static_cast(buffer_data.ptr); - - return SparseHamiltonian{ - std::vector({data_ptr, data_ptr + data.size()}), - indices, indptr, wires}; - })) - .def("__repr__", &SparseHamiltonian::getObsName) - .def("get_wires", &SparseHamiltonian::getWires, - "Get wires of observables") - .def( - "__eq__", - [](const SparseHamiltonian &self, - py::handle other) -> bool { - if (!py::isinstance>(other)) { - return false; - } - auto other_cast = other.cast>(); - return self == other_cast; - }, - "Compare two observables"); -#endif } /** @@ -708,7 +628,9 @@ template void lightningClassBindings(py::module_ &m) { /* Observables submodule */ py::module_ obs_submodule = m.def_submodule("observables", "Submodule for observables classes."); - registerObservables(obs_submodule); + // registerBackendAgnosticObservables(obs_submodule); + registerBackendAgnosticObservables(obs_submodule); + registerBackendSpecificObservables(obs_submodule); //***********************************************************************// // Measurements 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 3f4e81b015..2ebf7d3f95 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp @@ -264,6 +264,73 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { "Variance of a sparse Hamiltonian."); } +/** + * @brief Register backend specific observables. + * + * @tparam StateVectorT + * @param m Pybind module + */ +template +void registerBackendSpecificObservables(py::module_ &m) { + using PrecisionT = + typename StateVectorT::PrecisionT; // Statevector's precision. + using ComplexT = + typename StateVectorT::ComplexT; // Statevector's complex type. + using ParamT = PrecisionT; // Parameter's data precision + + const std::string bitsize = + std::to_string(sizeof(std::complex) * 8); + + using np_arr_c = py::array_t, py::array::c_style>; + + std::string class_name; + + class_name = "SparseHamiltonianC" + bitsize; + using np_arr_sparse_ind = typename std::conditional< + std::is_same::value, + py::array_t, + py::array_t>::type; + using IdxT = typename SparseHamiltonian::IdxT; + py::class_, + std::shared_ptr>, + Observable>(m, class_name.c_str(), + py::module_local()) + .def(py::init([](const np_arr_c &data, const np_arr_sparse_ind &indices, + const np_arr_sparse_ind &offsets, + const std::vector &wires) { + const py::buffer_info buffer_data = data.request(); + const auto *data_ptr = static_cast(buffer_data.ptr); + + const py::buffer_info buffer_indices = indices.request(); + const auto *indices_ptr = + static_cast(buffer_indices.ptr); + + const py::buffer_info buffer_offsets = offsets.request(); + const auto *offsets_ptr = + static_cast(buffer_offsets.ptr); + + return SparseHamiltonian{ + std::vector({data_ptr, data_ptr + data.size()}), + std::vector({indices_ptr, indices_ptr + indices.size()}), + std::vector({offsets_ptr, offsets_ptr + offsets.size()}), + wires}; + })) + .def("__repr__", &SparseHamiltonian::getObsName) + .def("get_wires", &SparseHamiltonian::getWires, + "Get wires of observables") + .def( + "__eq__", + [](const SparseHamiltonian &self, + py::handle other) -> bool { + if (!py::isinstance>(other)) { + return false; + } + auto other_cast = other.cast>(); + return self == other_cast; + }, + "Compare two observables"); +} + /** * @brief Register backend specific adjoint Jacobian methods. * diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/bindings/LKokkosBindings.hpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/bindings/LKokkosBindings.hpp index bd9d89d72f..6432864c4e 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/bindings/LKokkosBindings.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/bindings/LKokkosBindings.hpp @@ -24,6 +24,7 @@ #include "ConstantUtil.hpp" // lookup #include "GateOperation.hpp" #include "MeasurementsKokkos.hpp" +#include "ObservablesKokkos.hpp" #include "StateVectorKokkos.hpp" #include "TypeList.hpp" #include "Util.hpp" // exp2 @@ -33,6 +34,7 @@ namespace { using namespace Pennylane::Bindings; using namespace Pennylane::LightningKokkos::Algorithms; using namespace Pennylane::LightningKokkos::Measures; +using namespace Pennylane::LightningKokkos::Observables; using Kokkos::InitializationSettings; using Pennylane::LightningKokkos::StateVectorKokkos; using Pennylane::Util::exp2; @@ -214,6 +216,58 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { "Variance of a sparse Hamiltonian."); } +/** + * @brief Register observable classes. + * + * @tparam StateVectorT + * @param m Pybind module + */ +template +void registerBackendSpecificObservables(py::module_ &m) { + using PrecisionT = + typename StateVectorT::PrecisionT; // Statevector's precision. + using ParamT = PrecisionT; // Parameter's data precision + + const std::string bitsize = + std::to_string(sizeof(std::complex) * 8); + + using np_arr_c = py::array_t, py::array::c_style>; + + std::string class_name; + + class_name = "SparseHamiltonianC" + bitsize; + py::class_, + std::shared_ptr>, + Observable>(m, class_name.c_str(), + py::module_local()) + .def(py::init([](const np_arr_c &data, + const std::vector &indices, + const std::vector &indptr, + const std::vector &wires) { + using ComplexT = typename StateVectorT::ComplexT; + const py::buffer_info buffer_data = data.request(); + const auto *data_ptr = static_cast(buffer_data.ptr); + + return SparseHamiltonian{ + std::vector({data_ptr, data_ptr + data.size()}), + indices, indptr, wires}; + })) + .def("__repr__", &SparseHamiltonian::getObsName) + .def("get_wires", &SparseHamiltonian::getWires, + "Get wires of observables") + .def( + "__eq__", + [](const SparseHamiltonian &self, + py::handle other) -> bool { + if (!py::isinstance>(other)) { + return false; + } + auto other_cast = other.cast>(); + return self == other_cast; + }, + "Compare two observables"); +} + /** * @brief Register backend specific adjoint Jacobian methods. * diff --git a/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp b/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp index 190a9c6525..7cd5a1ee2e 100644 --- a/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp @@ -180,6 +180,15 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { }); } +/** + * @brief Register backend specific observables. + * + * @tparam StateVectorT + * @param m Pybind module + */ +template +void registerBackendSpecificObservables([[maybe_unused]] py::module_ &m) {} + /** * @brief Register Vector Jacobian Product. */ From a0141775e4c9bd31e4402877f6a6fdf33975d778 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 17:28:10 +0000 Subject: [PATCH 18/43] Fix sparse_hamiltonianmpi_c and add getWires test. --- pennylane_lightning/core/_serialize.py | 53 +++++++------------ .../observables/tests/Test_Observables.cpp | 5 ++ 2 files changed, 25 insertions(+), 33 deletions(-) diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index 49d46768a7..e1e93b3b21 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -95,26 +95,17 @@ def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False if use_mpi: self.use_mpi = use_mpi - self.statevectormpi_c128 = lightning_ops.StateVectorMPIC128 - self.named_obsmpi_c64 = lightning_ops.observablesMPI.NamedObsMPIC64 - self.named_obsmpi_c128 = lightning_ops.observablesMPI.NamedObsMPIC128 - self.hermitian_obsmpi_c64 = lightning_ops.observablesMPI.HermitianObsMPIC64 - self.hermitian_obsmpi_c128 = lightning_ops.observablesMPI.HermitianObsMPIC128 - self.tensor_prod_obsmpi_c64 = lightning_ops.observablesMPI.TensorProdObsMPIC64 - self.tensor_prod_obsmpi_c128 = lightning_ops.observablesMPI.TensorProdObsMPIC128 - self.hamiltonianmpi_c64 = lightning_ops.observablesMPI.HamiltonianMPIC64 - self.hamiltonianmpi_c128 = lightning_ops.observablesMPI.HamiltonianMPIC128 - - if self.device_name == "lightning.gpu": - self.sparse_hamiltonianmpi_c64 = ( - lightning_ops.observablesMPI.SparseHamiltonianMPIC64 - ) - self.sparse_hamiltonianmpi_c128 = ( - lightning_ops.observablesMPI.SparseHamiltoniaMPInC128 - ) - else: - ValueError(f"{self.device_name} does not support MPI.") - self.mpi_manager = lightning_ops.MPIManager + self.statevector_mpi_c128 = lightning_ops.StateVectorMPIC128 + self.named_obs_mpi_c64 = lightning_ops.observablesMPI.NamedObsMPIC64 + self.named_obs_mpi_c128 = lightning_ops.observablesMPI.NamedObsMPIC128 + self.hermitian_obs_mpi_c64 = lightning_ops.observablesMPI.HermitianObsMPIC64 + self.hermitian_obs_mpi_c128 = lightning_ops.observablesMPI.HermitianObsMPIC128 + self.tensor_prod_obs_mpi_c64 = lightning_ops.observablesMPI.TensorProdObsMPIC64 + self.tensor_prod_obs_mpi_c128 = lightning_ops.observablesMPI.TensorProdObsMPIC128 + self.hamiltonian_mpi_c64 = lightning_ops.observablesMPI.HamiltonianMPIC64 + self.hamiltonian_mpi_c128 = lightning_ops.observablesMPI.HamiltonianMPIC128 + + self._mpi_manager = lightning_ops.MPIManager @property def ctype(self): @@ -129,46 +120,42 @@ def rtype(self): @property def sv_type(self): if self.use_mpi: - return self.statevectormpi_c128 + return self.statevector_mpi_c128 return self.statevector_c128 @property def named_obs(self): """Named observable matching ``use_csingle`` precision.""" if self.use_mpi: - return self.named_obsmpi_c64 if self.use_csingle else self.named_obsmpi_c128 + return self.named_obs_mpi_c64 if self.use_csingle else self.named_obs_mpi_c128 return self.named_obs_c64 if self.use_csingle else self.named_obs_c128 @property def hermitian_obs(self): """Hermitian observable matching ``use_csingle`` precision.""" if self.use_mpi: - return self.hermitian_obsmpi_c64 if self.use_csingle else self.hermitian_obsmpi_c128 + return self.hermitian_obs_mpi_c64 if self.use_csingle else self.hermitian_obs_mpi_c128 return self.hermitian_obs_c64 if self.use_csingle else self.hermitian_obs_c128 @property def tensor_obs(self): """Tensor product observable matching ``use_csingle`` precision.""" if self.use_mpi: - return self.tensor_prod_obsmpi_c64 if self.use_csingle else self.tensor_prod_obsmpi_c128 + return ( + self.tensor_prod_obs_mpi_c64 if self.use_csingle else self.tensor_prod_obs_mpi_c128 + ) return self.tensor_prod_obs_c64 if self.use_csingle else self.tensor_prod_obs_c128 @property def hamiltonian_obs(self): """Hamiltonian observable matching ``use_csingle`` precision.""" if self.use_mpi: - return self.hamiltonianmpi_c64 if self.use_csingle else self.hamiltonianmpi_c128 + return self.hamiltonian_mpi_c64 if self.use_csingle else self.hamiltonian_mpi_c128 return self.hamiltonian_c64 if self.use_csingle else self.hamiltonian_c128 @property def sparse_hamiltonian_obs(self): """SparseHamiltonian observable matching ``use_csingle`` precision.""" - if self.use_mpi: - return ( - self.sparse_hamiltonianmpi_c64 - if self.use_csingle - else self.sparse_hamiltonianmpi_c128 - ) return self.sparse_hamiltonian_c64 if self.use_csingle else self.sparse_hamiltonian_c128 def _named_obs(self, observable, wires_map: dict): @@ -211,9 +198,9 @@ def _sparse_hamiltonian(self, observable, wires_map: dict): H_sparse = SparseHamiltonian(Hmat, wires=range(1)) spm = H_sparse.sparse_matrix() # Only root 0 needs the overall sparsematrix data - if self.mpi_manager().getRank() == 0: + if self._mpi_manager().getRank() == 0: spm = observable.sparse_matrix() - self.mpi_manager().Barrier() + self._mpi_manager().Barrier() else: spm = observable.sparse_matrix() spm = observable.sparse_matrix() diff --git a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp index 12f25ef667..8890f8c121 100644 --- a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp +++ b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp @@ -484,6 +484,11 @@ template void testSparseHamiltonianBase() { ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, {0, 1, 2}); + DYNAMIC_SECTION("getWires - " + << StateVectorToName::name) { + REQUIRE(sparseH->getWires() == std::vector{0, 1, 2}); + } + DYNAMIC_SECTION("getObsName - " << StateVectorToName::name) { REQUIRE(sparseH->getObsName() == From 2f9d14bf3c60abe8bbf3d8f1e3e82180b1923b15 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 19:08:49 +0000 Subject: [PATCH 19/43] Add sparseH diff capability in LQ. --- pennylane_lightning/core/_serialize.py | 10 +-- .../observables/tests/Test_Observables.cpp | 14 ++-- .../bindings/LQubitBindings.hpp | 51 +++++++++++++- .../observables/ObservablesLQubit.cpp | 6 ++ .../observables/ObservablesLQubit.hpp | 69 +++++++++++++++++++ tests/test_adjoint_jacobian.py | 8 --- tests/test_serialize.py | 57 ++++++++------- 7 files changed, 159 insertions(+), 56 deletions(-) diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index e1e93b3b21..4534b79aea 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -87,9 +87,8 @@ def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False self.tensor_prod_obs_c128 = lightning_ops.observables.TensorProdObsC128 self.hamiltonian_c64 = lightning_ops.observables.HamiltonianC64 self.hamiltonian_c128 = lightning_ops.observables.HamiltonianC128 - if device_name in ["lightning.gpu", "lightning.kokkos"]: - self.sparse_hamiltonian_c64 = lightning_ops.observables.SparseHamiltonianC64 - self.sparse_hamiltonian_c128 = lightning_ops.observables.SparseHamiltonianC128 + self.sparse_hamiltonian_c64 = lightning_ops.observables.SparseHamiltonianC64 + self.sparse_hamiltonian_c128 = lightning_ops.observables.SparseHamiltonianC128 self.use_mpi = False @@ -241,10 +240,7 @@ def _ob(self, observable, wires_map): return self._tensor_ob(observable, wires_map) if observable.name == "Hamiltonian": return self._hamiltonian(observable, wires_map) - if ( - self.device_name in ["lightning.gpu", "lightning.kokkos"] - and observable.name == "SparseHamiltonian" - ): + if observable.name == "SparseHamiltonian": return self._sparse_hamiltonian(observable, wires_map) if isinstance(observable, (PauliX, PauliY, PauliZ, Identity, Hadamard)): return self._named_obs(observable, wires_map) diff --git a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp index 8890f8c121..f6877e69a9 100644 --- a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp +++ b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp @@ -30,8 +30,8 @@ /// @cond DEV namespace { using namespace Pennylane::Observables; - using Pennylane::Util::createProductState; +using Pennylane::Util::createRandomStateVectorData; using Pennylane::Util::createZeroState; using Pennylane::Util::isApproxEqual; using Pennylane::Util::LightningException; @@ -41,7 +41,6 @@ using Pennylane::Util::TestVector; #ifdef _ENABLE_PLQUBIT constexpr bool BACKEND_FOUND = true; -constexpr bool SPARSE_HAM_SUPPORTED = false; #include "TestHelpersStateVectors.hpp" // TestStateVectorBackends, StateVectorToName @@ -53,7 +52,6 @@ using namespace Pennylane::LightningQubit::Util; #elif _ENABLE_PLKOKKOS == 1 constexpr bool BACKEND_FOUND = true; -constexpr bool SPARSE_HAM_SUPPORTED = true; #include "TestHelpersStateVectors.hpp" // TestStateVectorBackends, StateVectorToName @@ -65,7 +63,6 @@ using namespace Pennylane::LightningKokkos::Util; #elif _ENABLE_PLGPU == 1 constexpr bool BACKEND_FOUND = true; -constexpr bool SPARSE_HAM_SUPPORTED = true; #include "TestHelpersStateVectors.hpp" @@ -77,7 +74,6 @@ using namespace Pennylane::LightningGPU::Util; #else constexpr bool BACKEND_FOUND = false; -constexpr bool SPARSE_HAM_SUPPORTED = false; using TestStateVectorBackends = Pennylane::Util::TypeList; template struct StateVectorToName {}; @@ -484,12 +480,12 @@ template void testSparseHamiltonianBase() { ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, {0, 1, 2}); - DYNAMIC_SECTION("getWires - " + DYNAMIC_SECTION("SparseHamiltonianBase - getWires - " << StateVectorToName::name) { REQUIRE(sparseH->getWires() == std::vector{0, 1, 2}); } - DYNAMIC_SECTION("getObsName - " + DYNAMIC_SECTION("SparseHamiltonianBase - getObsName - " << StateVectorToName::name) { REQUIRE(sparseH->getObsName() == "SparseHamiltonian: {\n" @@ -503,7 +499,7 @@ template void testSparseHamiltonianBase() { "}"); } - DYNAMIC_SECTION("applyInPlace must fail - " + DYNAMIC_SECTION("SparseHamiltonianBase - applyInPlace must fail - " << StateVectorToName::name) { auto init_state = @@ -521,7 +517,7 @@ template void testSparseHamiltonianBase() { TEST_CASE("Methods implemented in the SparseHamiltonianBase class", "[SparseHamiltonianBase]") { - if constexpr (SPARSE_HAM_SUPPORTED) { + if constexpr (BACKEND_FOUND) { testSparseHamiltonianBase(); } } \ No newline at end of file diff --git a/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp b/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp index 7cd5a1ee2e..2f83f2f39d 100644 --- a/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp @@ -25,15 +25,17 @@ #include "DynamicDispatcher.hpp" #include "GateOperation.hpp" #include "MeasurementsLQubit.hpp" +#include "ObservablesLQubit.hpp" #include "StateVectorLQubitRaw.hpp" #include "TypeList.hpp" #include "VectorJacobianProduct.hpp" /// @cond DEV namespace { -using namespace Pennylane::LightningQubit::Measures; -using namespace Pennylane::LightningQubit::Algorithms; using namespace Pennylane::Bindings; +using namespace Pennylane::LightningQubit::Algorithms; +using namespace Pennylane::LightningQubit::Measures; +using namespace Pennylane::LightningQubit::Observables; using Pennylane::LightningQubit::StateVectorLQubitRaw; } // namespace /// @endcond @@ -187,7 +189,50 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { * @param m Pybind module */ template -void registerBackendSpecificObservables([[maybe_unused]] py::module_ &m) {} +void registerBackendSpecificObservables([[maybe_unused]] py::module_ &m) { + using PrecisionT = + typename StateVectorT::PrecisionT; // Statevector's precision. + using ParamT = PrecisionT; // Parameter's data precision + + const std::string bitsize = + std::to_string(sizeof(std::complex) * 8); + + using np_arr_c = py::array_t, py::array::c_style>; + + std::string class_name; + + class_name = "SparseHamiltonianC" + bitsize; + py::class_, + std::shared_ptr>, + Observable>(m, class_name.c_str(), + py::module_local()) + .def(py::init([](const np_arr_c &data, + const std::vector &indices, + const std::vector &indptr, + const std::vector &wires) { + using ComplexT = typename StateVectorT::ComplexT; + const py::buffer_info buffer_data = data.request(); + const auto *data_ptr = static_cast(buffer_data.ptr); + + return SparseHamiltonian{ + std::vector({data_ptr, data_ptr + data.size()}), + indices, indptr, wires}; + })) + .def("__repr__", &SparseHamiltonian::getObsName) + .def("get_wires", &SparseHamiltonian::getWires, + "Get wires of observables") + .def( + "__eq__", + [](const SparseHamiltonian &self, + py::handle other) -> bool { + if (!py::isinstance>(other)) { + return false; + } + auto other_cast = other.cast>(); + return self == other_cast; + }, + "Compare two observables"); +} /** * @brief Register Vector Jacobian Product. diff --git a/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.cpp b/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.cpp index 0a1fb54c48..e45e2c1572 100644 --- a/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.cpp @@ -41,3 +41,9 @@ template class Observables::Hamiltonian>; template class Observables::Hamiltonian>; template class Observables::Hamiltonian>; + +template class Observables::SparseHamiltonian>; +template class Observables::SparseHamiltonian>; + +template class Observables::SparseHamiltonian>; +template class Observables::SparseHamiltonian>; diff --git a/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.hpp b/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.hpp index 3433a3fcc3..c68d119526 100644 --- a/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.hpp @@ -27,6 +27,7 @@ #include "LinearAlgebra.hpp" // scaleAndAdd #include "Macros.hpp" // use_openmp #include "Observables.hpp" +#include "SparseLinAlg.hpp" #include "StateVectorLQubitManaged.hpp" #include "StateVectorLQubitRaw.hpp" #include "Util.hpp" @@ -359,4 +360,72 @@ class Hamiltonian final : public HamiltonianBase { } }; +/** + * @brief Sparse representation of Hamiltonian + * + */ +template +class SparseHamiltonian final : public SparseHamiltonianBase { + private: + using BaseType = SparseHamiltonianBase; + + public: + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + using IdxT = typename BaseType::IdxT; + + /** + * @brief Create a SparseHamiltonian from data, indices and offsets in CSR + * format. + * + * @param data Arguments to construct data + * @param indices Arguments to construct indices + * @param offsets Arguments to construct offsets + * @param wires Arguments to construct wires + */ + template + explicit SparseHamiltonian(T1 &&data, T2 &&indices, T3 &&offsets, + T4 &&wires) + : BaseType{data, indices, offsets, wires} {} + + /** + * @brief Convenient wrapper for the constructor as the constructor does not + * convert the std::shared_ptr with a derived class correctly. + * + * This function is useful as std::make_shared does not handle + * brace-enclosed initializer list correctly. + * + * @param data Argument to construct data + * @param indices Argument to construct indices + * @param offsets Argument to construct ofsets + * @param wires Argument to construct wires + */ + static auto create(std::initializer_list data, + std::initializer_list indices, + std::initializer_list offsets, + std::initializer_list wires) + -> std::shared_ptr> { + return std::shared_ptr>( + new SparseHamiltonian{ + std::move(data), std::move(indices), std::move(offsets), + std::move(wires)}); + } + + /** + * @brief Updates the statevector SV:->SV', where SV' = a*H*SV, and where H + * is a sparse Hamiltonian. + * + */ + void applyInPlace(StateVectorT &sv) const override { + PL_ABORT_IF_NOT(this->wires_.size() == sv.getNumQubits(), + "SparseH wire count does not match state-vector size"); + auto operator_vector = Util::apply_Sparse_Matrix( + sv.getData(), sv.getLength(), this->offsets_.data(), + this->offsets_.size(), this->indices_.data(), this->data_.data(), + this->data_.size()); + + sv.updateData(operator_vector); + } +}; + } // namespace Pennylane::LightningQubit::Observables \ No newline at end of file diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py index c9b338cdc3..c851e1b489 100644 --- a/tests/test_adjoint_jacobian.py +++ b/tests/test_adjoint_jacobian.py @@ -920,10 +920,6 @@ def circuit(params): @pytest.mark.skipif(not ld._CPP_BINARY_AVAILABLE, reason="Lightning binary required") -@pytest.mark.skipif( - device_name not in ["lightning.gpu", "lightning.kokkos"], - reason="SparseHamiltonian only supported by Lightning-Kokkos", -) def test_tape_qchem_sparse(tol): """Tests the circuit Ansatz with a QChem Hamiltonian produces correct results""" @@ -956,10 +952,6 @@ def circuit(params): @pytest.mark.skipif(not ld._CPP_BINARY_AVAILABLE, reason="Lightning binary required") -@pytest.mark.skipif( - device_name not in ["lightning.gpu", "lightning.kokkos"], - reason="SparseHamiltonian only supported by Lightning-Kokkos", -) @pytest.mark.parametrize( "returns", [ diff --git a/tests/test_serialize.py b/tests/test_serialize.py index 92938b3406..360ac0e71b 100644 --- a/tests/test_serialize.py +++ b/tests/test_serialize.py @@ -60,6 +60,8 @@ TensorProdObsC128, HamiltonianC64, HamiltonianC128, + SparseHamiltonianC64, + SparseHamiltonianC128, ) @@ -70,41 +72,38 @@ def test_wrong_device_name(): QuantumScriptSerializer("thunder.qubit") -obs_list = [ - (qml.PauliZ(0), NamedObsC128), - (qml.PauliZ(0) @ qml.PauliZ(1), TensorProdObsC128), - (qml.Hadamard(0), NamedObsC128), - (qml.Hermitian(np.eye(2), wires=0), HermitianObsC128), - ( - qml.PauliZ(0) @ qml.Hadamard(1) @ (0.1 * (qml.PauliZ(2) + qml.PauliX(3))), - HamiltonianC128, - ), - ( +@pytest.mark.parametrize( + "obs,obs_type", + [ + (qml.PauliZ(0), NamedObsC128), + (qml.PauliZ(0) @ qml.PauliZ(1), TensorProdObsC128), + (qml.Hadamard(0), NamedObsC128), + (qml.Hermitian(np.eye(2), wires=0), HermitianObsC128), ( - qml.Hermitian(np.eye(2), wires=0) - @ qml.Hermitian(np.eye(2), wires=1) - @ qml.Projector([0], wires=2) + qml.PauliZ(0) @ qml.Hadamard(1) @ (0.1 * (qml.PauliZ(2) + qml.PauliX(3))), + HamiltonianC128, ), - TensorProdObsC128, - ), - ( - qml.PauliZ(0) @ qml.Hermitian(np.eye(2), wires=1) @ qml.Projector([0], wires=2), - TensorProdObsC128, - ), - (qml.Projector([0], wires=0), HermitianObsC128), - (qml.Hamiltonian([1], [qml.PauliZ(0)]), HamiltonianC128), - (qml.sum(qml.Hadamard(0), qml.PauliX(1)), HermitianObsC128), -] -if device_name in ["lightning.gpu", "lightning.kokkos"]: - obs_list += [ + ( + ( + qml.Hermitian(np.eye(2), wires=0) + @ qml.Hermitian(np.eye(2), wires=1) + @ qml.Projector([0], wires=2) + ), + TensorProdObsC128, + ), + ( + qml.PauliZ(0) @ qml.Hermitian(np.eye(2), wires=1) @ qml.Projector([0], wires=2), + TensorProdObsC128, + ), + (qml.Projector([0], wires=0), HermitianObsC128), + (qml.Hamiltonian([1], [qml.PauliZ(0)]), HamiltonianC128), + (qml.sum(qml.Hadamard(0), qml.PauliX(1)), HermitianObsC128), ( qml.SparseHamiltonian(qml.Hamiltonian([1], [qml.PauliZ(0)]).sparse_matrix(), wires=[0]), SparseHamiltonianC128, ), - ] - - -@pytest.mark.parametrize("obs,obs_type", obs_list) + ], +) def test_obs_returns_expected_type(obs, obs_type): """Tests that observables get serialized to the expected type.""" assert isinstance( From 812a1c6f2fb120c3cf8ff6f8aa46b6ca2267183c Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 15:28:42 -0400 Subject: [PATCH 20/43] Add sparse Hamiltonian support for Lightning-Kokkos (#527) * Use more generic indices in base std::size_t. * Quick and dirty spham bindings. * Add sparse_ham serialization. * Add sparse_ham tests in tests/test_adjoint_jacobian.py' * Bug fix sparse product. * Fix python bindings LGPU idxT * Fix serial tests and update changelog. * Fix tidy & sparse adjoint test device name. * Fix tidy warning for sparse_ham. * Send backend-specific ops in respective modules. * Fix sparse_hamiltonianmpi_c and add getWires test. --- .github/CHANGELOG.md | 3 + pennylane_lightning/core/_serialize.py | 101 ++++++++++-------- .../core/src/bindings/Bindings.hpp | 51 +-------- .../core/src/observables/Observables.hpp | 24 +++-- .../observables/tests/Test_Observables.cpp | 38 +++++-- .../lightning_gpu/bindings/LGPUBindings.hpp | 67 ++++++++++++ .../observables/ObservablesGPU.hpp | 12 +-- .../lightning_kokkos/StateVectorKokkos.hpp | 19 +++- .../bindings/LKokkosBindings.hpp | 54 ++++++++++ .../observables/ObservablesKokkos.cpp | 3 + .../observables/ObservablesKokkos.hpp | 71 ++++++++++++ .../bindings/LQubitBindings.hpp | 9 ++ tests/test_adjoint_jacobian.py | 67 +++++++++++- tests/test_serialize.py | 59 +++++----- 14 files changed, 425 insertions(+), 153 deletions(-) diff --git a/.github/CHANGELOG.md b/.github/CHANGELOG.md index 2288658d22..9ef71c78df 100644 --- a/.github/CHANGELOG.md +++ b/.github/CHANGELOG.md @@ -2,6 +2,9 @@ ### New features since last release +* Add `SparseHamiltonian` support for Lightning-Kokkos. + [(#527)] (https://github.com/PennyLaneAI/pennylane-lightning/pull/527) + * Integrate python/pybind layer of distributed Lightning-GPU into the Lightning monorepo with python unit tests. [(#518)] (https://github.com/PennyLaneAI/pennylane-lightning/pull/518) diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index eb05f41709..e1e93b3b21 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -25,6 +25,8 @@ Identity, StatePrep, Rot, + Hamiltonian, + SparseHamiltonian, ) from pennylane.operation import Tensor from pennylane.tape import QuantumTape @@ -85,31 +87,25 @@ def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False self.tensor_prod_obs_c128 = lightning_ops.observables.TensorProdObsC128 self.hamiltonian_c64 = lightning_ops.observables.HamiltonianC64 self.hamiltonian_c128 = lightning_ops.observables.HamiltonianC128 - - if self.device_name == "lightning.gpu": - self.sparsehamiltonian_c64 = lightning_ops.observables.SparseHamiltonianC64 - self.sparsehamiltonian_c128 = lightning_ops.observables.SparseHamiltonianC128 + if device_name in ["lightning.gpu", "lightning.kokkos"]: + self.sparse_hamiltonian_c64 = lightning_ops.observables.SparseHamiltonianC64 + self.sparse_hamiltonian_c128 = lightning_ops.observables.SparseHamiltonianC128 self.use_mpi = False if use_mpi: self.use_mpi = use_mpi - self.statevectormpi_c128 = lightning_ops.StateVectorMPIC128 - self.named_obsmpi_c64 = lightning_ops.observablesMPI.NamedObsMPIC64 - self.named_obsmpi_c128 = lightning_ops.observablesMPI.NamedObsMPIC128 - self.hermitian_obsmpi_c64 = lightning_ops.observablesMPI.HermitianObsMPIC64 - self.hermitian_obsmpi_c128 = lightning_ops.observablesMPI.HermitianObsMPIC128 - self.tensor_prod_obsmpi_c64 = lightning_ops.observablesMPI.TensorProdObsMPIC64 - self.tensor_prod_obsmpi_c128 = lightning_ops.observablesMPI.TensorProdObsMPIC128 - self.hamiltonianmpi_c64 = lightning_ops.observablesMPI.HamiltonianMPIC64 - self.hamiltonianmpi_c128 = lightning_ops.observablesMPI.HamiltonianMPIC128 - - if self.device_name == "lightning.gpu": - self.sparsehamiltonianmpi_c64 = lightning_ops.observablesMPI.SparseHamiltonianMPIC64 - self.sparsehamiltonianmpi_c128 = ( - lightning_ops.observablesMPI.SparseHamiltonianMPIC128 - ) - self.mpi_manager = lightning_ops.MPIManager + self.statevector_mpi_c128 = lightning_ops.StateVectorMPIC128 + self.named_obs_mpi_c64 = lightning_ops.observablesMPI.NamedObsMPIC64 + self.named_obs_mpi_c128 = lightning_ops.observablesMPI.NamedObsMPIC128 + self.hermitian_obs_mpi_c64 = lightning_ops.observablesMPI.HermitianObsMPIC64 + self.hermitian_obs_mpi_c128 = lightning_ops.observablesMPI.HermitianObsMPIC128 + self.tensor_prod_obs_mpi_c64 = lightning_ops.observablesMPI.TensorProdObsMPIC64 + self.tensor_prod_obs_mpi_c128 = lightning_ops.observablesMPI.TensorProdObsMPIC128 + self.hamiltonian_mpi_c64 = lightning_ops.observablesMPI.HamiltonianMPIC64 + self.hamiltonian_mpi_c128 = lightning_ops.observablesMPI.HamiltonianMPIC128 + + self._mpi_manager = lightning_ops.MPIManager @property def ctype(self): @@ -124,47 +120,43 @@ def rtype(self): @property def sv_type(self): if self.use_mpi: - return self.statevectormpi_c128 + return self.statevector_mpi_c128 return self.statevector_c128 @property def named_obs(self): """Named observable matching ``use_csingle`` precision.""" if self.use_mpi: - return self.named_obsmpi_c64 if self.use_csingle else self.named_obsmpi_c128 + return self.named_obs_mpi_c64 if self.use_csingle else self.named_obs_mpi_c128 return self.named_obs_c64 if self.use_csingle else self.named_obs_c128 @property def hermitian_obs(self): """Hermitian observable matching ``use_csingle`` precision.""" if self.use_mpi: - return self.hermitian_obsmpi_c64 if self.use_csingle else self.hermitian_obsmpi_c128 + return self.hermitian_obs_mpi_c64 if self.use_csingle else self.hermitian_obs_mpi_c128 return self.hermitian_obs_c64 if self.use_csingle else self.hermitian_obs_c128 @property def tensor_obs(self): """Tensor product observable matching ``use_csingle`` precision.""" if self.use_mpi: - return self.tensor_prod_obsmpi_c64 if self.use_csingle else self.tensor_prod_obsmpi_c128 + return ( + self.tensor_prod_obs_mpi_c64 if self.use_csingle else self.tensor_prod_obs_mpi_c128 + ) return self.tensor_prod_obs_c64 if self.use_csingle else self.tensor_prod_obs_c128 @property def hamiltonian_obs(self): """Hamiltonian observable matching ``use_csingle`` precision.""" if self.use_mpi: - return self.hamiltonianmpi_c64 if self.use_csingle else self.hamiltonianmpi_c128 + return self.hamiltonian_mpi_c64 if self.use_csingle else self.hamiltonian_mpi_c128 return self.hamiltonian_c64 if self.use_csingle else self.hamiltonian_c128 @property - def sparsehamiltonian_obs(self): - """Sparse Hamiltonian observable matching ``use_csingle`` precision.""" - if self.use_mpi: - return ( - self.sparsehamiltonianmpi_c64 - if self.use_csingle - else self.sparsehamiltonianmpi_c128 - ) - return self.sparsehamiltonian_c64 if self.use_csingle else self.sparsehamiltonian_c128 + def sparse_hamiltonian_obs(self): + """SparseHamiltonian observable matching ``use_csingle`` precision.""" + return self.sparse_hamiltonian_c64 if self.use_csingle else self.sparse_hamiltonian_c128 def _named_obs(self, observable, wires_map: dict): """Serializes a Named observable""" @@ -190,25 +182,37 @@ def _hamiltonian(self, observable, wires_map: dict): terms = [self._ob(t, wires_map) for t in observable.ops] return self.hamiltonian_obs(coeffs, terms) - def _sparsehamiltonian(self, observable, wires_map: dict): - wires = [] - wires_list = observable.wires.tolist() - wires.extend([wires_map[w] for w in wires_list]) + def _sparse_hamiltonian(self, observable, wires_map: dict): + """Serialize an observable (Sparse Hamiltonian) + + Args: + observable (Observable): the input observable (Sparse Hamiltonian) + wire_map (dict): a dictionary mapping input wires to the device's backend wires + + Returns: + sparse_hamiltonian_obs (SparseHamiltonianC64 or SparseHamiltonianC128): A Sparse Hamiltonian observable object compatible with the C++ backend + """ + if self.use_mpi: - Hmat = qml.Hamiltonian([1.0], [qml.Identity(0)]).sparse_matrix() - H_sparse = qml.SparseHamiltonian(Hmat, wires=range(1)) + Hmat = Hamiltonian([1.0], [Identity(0)]).sparse_matrix() + H_sparse = SparseHamiltonian(Hmat, wires=range(1)) spm = H_sparse.sparse_matrix() # Only root 0 needs the overall sparsematrix data - if self.mpi_manager().getRank() == 0: + if self._mpi_manager().getRank() == 0: spm = observable.sparse_matrix() - self.mpi_manager().Barrier() + self._mpi_manager().Barrier() else: spm = observable.sparse_matrix() + spm = observable.sparse_matrix() data = np.array(spm.data).astype(self.ctype) - indices = np.array(spm.indices).astype(self.rtype) - offsets = np.array(spm.indptr).astype(self.rtype) + indices = np.array(spm.indices).astype(np.int64) + offsets = np.array(spm.indptr).astype(np.int64) + + wires = [] + wires_list = observable.wires.tolist() + wires.extend([wires_map[w] for w in wires_list]) - return self.sparsehamiltonian_obs(data, indices, offsets, wires) + return self.sparse_hamiltonian_obs(data, indices, offsets, wires) def _pauli_word(self, observable, wires_map: dict): """Serialize a :class:`pennylane.pauli.PauliWord` into a Named or Tensor observable.""" @@ -237,8 +241,11 @@ def _ob(self, observable, wires_map): return self._tensor_ob(observable, wires_map) if observable.name == "Hamiltonian": return self._hamiltonian(observable, wires_map) - if observable.name == "SparseHamiltonian" and self.device_name != "lightning.qubit": - return self._sparsehamiltonian(observable, wires_map) + if ( + self.device_name in ["lightning.gpu", "lightning.kokkos"] + and observable.name == "SparseHamiltonian" + ): + return self._sparse_hamiltonian(observable, wires_map) if isinstance(observable, (PauliX, PauliY, PauliZ, Identity, Hadamard)): return self._named_obs(observable, wires_map) if observable._pauli_rep is not None: diff --git a/pennylane_lightning/core/src/bindings/Bindings.hpp b/pennylane_lightning/core/src/bindings/Bindings.hpp index 1c8fa5ba6c..ac2fdd8b7f 100644 --- a/pennylane_lightning/core/src/bindings/Bindings.hpp +++ b/pennylane_lightning/core/src/bindings/Bindings.hpp @@ -287,7 +287,8 @@ void registerInfo(py::module_ &m) { * @tparam StateVectorT * @param m Pybind module */ -template void registerObservables(py::module_ &m) { +template +void registerBackendAgnosticObservables(py::module_ &m) { using PrecisionT = typename StateVectorT::PrecisionT; // Statevector's precision. using ComplexT = @@ -299,10 +300,6 @@ template void registerObservables(py::module_ &m) { using np_arr_c = py::array_t, py::array::c_style>; using np_arr_r = py::array_t; - using np_arr_sparse_ind = typename std::conditional< - std::is_same::value, - py::array_t, - py::array_t>::type; std::string class_name; @@ -410,46 +407,6 @@ template void registerObservables(py::module_ &m) { return self == other_cast; }, "Compare two observables"); -#ifdef _ENABLE_PLGPU - class_name = "SparseHamiltonianC" + bitsize; - using SpIDX = typename SparseHamiltonian::IdxT; - py::class_, - std::shared_ptr>, - Observable>(m, class_name.c_str(), - py::module_local()) - .def(py::init([](const np_arr_c &data, const np_arr_sparse_ind &indices, - const np_arr_sparse_ind &offsets, - const std::vector &wires) { - const py::buffer_info buffer_data = data.request(); - const auto *data_ptr = static_cast(buffer_data.ptr); - - const py::buffer_info buffer_indices = indices.request(); - const auto *indices_ptr = static_cast(buffer_indices.ptr); - - const py::buffer_info buffer_offsets = offsets.request(); - const auto *offsets_ptr = static_cast(buffer_offsets.ptr); - - return SparseHamiltonian{ - std::vector({data_ptr, data_ptr + data.size()}), - std::vector({indices_ptr, indices_ptr + indices.size()}), - std::vector({offsets_ptr, offsets_ptr + offsets.size()}), - wires}; - })) - .def("__repr__", &SparseHamiltonian::getObsName) - .def("get_wires", &SparseHamiltonian::getWires, - "Get wires of observables") - .def( - "__eq__", - [](const SparseHamiltonian &self, - py::handle other) -> bool { - if (!py::isinstance>(other)) { - return false; - } - auto other_cast = other.cast>(); - return self == other_cast; - }, - "Compare two observables"); -#endif } /** @@ -671,7 +628,9 @@ template void lightningClassBindings(py::module_ &m) { /* Observables submodule */ py::module_ obs_submodule = m.def_submodule("observables", "Submodule for observables classes."); - registerObservables(obs_submodule); + // registerBackendAgnosticObservables(obs_submodule); + registerBackendAgnosticObservables(obs_submodule); + registerBackendSpecificObservables(obs_submodule); //***********************************************************************// // Measurements diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index ae062cc531..d7cd3a1f8f 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -414,6 +414,7 @@ class HamiltonianBase : public Observable { } }; +// NOLINTBEGIN /** * @brief Sparse representation of SparseHamiltonian * @@ -424,8 +425,8 @@ class SparseHamiltonianBase : public Observable { public: using PrecisionT = typename StateVectorT::PrecisionT; using ComplexT = typename StateVectorT::ComplexT; - // cuSparse required index type #ifdef _ENABLE_PLGPU + // cuSparse required index type using IdxT = typename std::conditional::value, int32_t, int64_t>::type; @@ -480,7 +481,7 @@ class SparseHamiltonianBase : public Observable { * * @param arg1 Argument to construct data * @param arg2 Argument to construct indices - * @param arg3 Argument to construct ofsets + * @param arg3 Argument to construct offsets * @param arg4 Argument to construct wires */ static auto create(std::initializer_list arg1, @@ -504,18 +505,20 @@ class SparseHamiltonianBase : public Observable { using Pennylane::Util::operator<<; std::ostringstream ss; ss << "SparseHamiltonian: {\n'data' : \n"; - for (const auto &d : data_) - // Note that for LGPU backend, ComplexT is std::complex as of 0.33 + for (const auto &d : data_) { // Note that for LGPU backend, ComplexT is + // std::complex as of 0.33 // release Need to revisit it once we set ComplexT as cuComplex // later - ss << "{" << d.real() << ", " << d.imag() << "}," - << "\n"; + ss << "{" << d.real() << ", " << d.imag() << "}, "; + } ss << ",\n'indices' : \n"; - for (const auto &i : indices_) - ss << i << std::endl; + for (const auto &i : indices_) { + ss << i << ", "; + } ss << ",\n'offsets' : \n"; - for (const auto &o : offsets_) - ss << o; + for (const auto &o : offsets_) { + ss << o << ", "; + } ss << "\n}"; return ss.str(); } @@ -526,5 +529,6 @@ class SparseHamiltonianBase : public Observable { return wires_; }; }; +// NOLINTEND } // namespace Pennylane::Observables \ No newline at end of file diff --git a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp index b6ff4758d6..8890f8c121 100644 --- a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp +++ b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp @@ -41,6 +41,7 @@ using Pennylane::Util::TestVector; #ifdef _ENABLE_PLQUBIT constexpr bool BACKEND_FOUND = true; +constexpr bool SPARSE_HAM_SUPPORTED = false; #include "TestHelpersStateVectors.hpp" // TestStateVectorBackends, StateVectorToName @@ -52,6 +53,7 @@ using namespace Pennylane::LightningQubit::Util; #elif _ENABLE_PLKOKKOS == 1 constexpr bool BACKEND_FOUND = true; +constexpr bool SPARSE_HAM_SUPPORTED = true; #include "TestHelpersStateVectors.hpp" // TestStateVectorBackends, StateVectorToName @@ -63,6 +65,7 @@ using namespace Pennylane::LightningKokkos::Util; #elif _ENABLE_PLGPU == 1 constexpr bool BACKEND_FOUND = true; +constexpr bool SPARSE_HAM_SUPPORTED = true; #include "TestHelpersStateVectors.hpp" @@ -74,6 +77,7 @@ using namespace Pennylane::LightningGPU::Util; #else constexpr bool BACKEND_FOUND = false; +constexpr bool SPARSE_HAM_SUPPORTED = false; using TestStateVectorBackends = Pennylane::Util::TypeList; template struct StateVectorToName {}; @@ -474,15 +478,33 @@ template void testSparseHamiltonianBase() { const std::size_t num_qubits = 3; std::mt19937 re{1337}; - DYNAMIC_SECTION("applyInPlace must fail - " + auto sparseH = SparseHamiltonianBase::create( + {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, + {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, {0, 1, 2}); + + DYNAMIC_SECTION("getWires - " + << StateVectorToName::name) { + REQUIRE(sparseH->getWires() == std::vector{0, 1, 2}); + } + + DYNAMIC_SECTION("getObsName - " << StateVectorToName::name) { + REQUIRE(sparseH->getObsName() == + "SparseHamiltonian: {\n" + "'data' : \n" + "{1, 0}, {1, 0}, {1, 0}, {1, 0}, {1, 0}, {1, 0}, {1, 0}, " + "{1, 0}, ,\n" + "'indices' : \n" + "7, 6, 5, 4, 3, 2, 1, 0, ,\n" + "'offsets' : \n" + "0, 1, 2, 3, 4, 5, 6, 7, 8, \n" + "}"); + } - auto sparseH = SparseHamiltonianBase::create( - {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, - ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, - ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, - {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, - {0, 1, 2}); + DYNAMIC_SECTION("applyInPlace must fail - " + << StateVectorToName::name) { auto init_state = createRandomStateVectorData(re, num_qubits); @@ -499,7 +521,7 @@ template void testSparseHamiltonianBase() { TEST_CASE("Methods implemented in the SparseHamiltonianBase class", "[SparseHamiltonianBase]") { - if constexpr (BACKEND_FOUND) { + if constexpr (SPARSE_HAM_SUPPORTED) { testSparseHamiltonianBase(); } } \ No newline at end of file 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 3f4e81b015..2ebf7d3f95 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindings.hpp @@ -264,6 +264,73 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { "Variance of a sparse Hamiltonian."); } +/** + * @brief Register backend specific observables. + * + * @tparam StateVectorT + * @param m Pybind module + */ +template +void registerBackendSpecificObservables(py::module_ &m) { + using PrecisionT = + typename StateVectorT::PrecisionT; // Statevector's precision. + using ComplexT = + typename StateVectorT::ComplexT; // Statevector's complex type. + using ParamT = PrecisionT; // Parameter's data precision + + const std::string bitsize = + std::to_string(sizeof(std::complex) * 8); + + using np_arr_c = py::array_t, py::array::c_style>; + + std::string class_name; + + class_name = "SparseHamiltonianC" + bitsize; + using np_arr_sparse_ind = typename std::conditional< + std::is_same::value, + py::array_t, + py::array_t>::type; + using IdxT = typename SparseHamiltonian::IdxT; + py::class_, + std::shared_ptr>, + Observable>(m, class_name.c_str(), + py::module_local()) + .def(py::init([](const np_arr_c &data, const np_arr_sparse_ind &indices, + const np_arr_sparse_ind &offsets, + const std::vector &wires) { + const py::buffer_info buffer_data = data.request(); + const auto *data_ptr = static_cast(buffer_data.ptr); + + const py::buffer_info buffer_indices = indices.request(); + const auto *indices_ptr = + static_cast(buffer_indices.ptr); + + const py::buffer_info buffer_offsets = offsets.request(); + const auto *offsets_ptr = + static_cast(buffer_offsets.ptr); + + return SparseHamiltonian{ + std::vector({data_ptr, data_ptr + data.size()}), + std::vector({indices_ptr, indices_ptr + indices.size()}), + std::vector({offsets_ptr, offsets_ptr + offsets.size()}), + wires}; + })) + .def("__repr__", &SparseHamiltonian::getObsName) + .def("get_wires", &SparseHamiltonian::getWires, + "Get wires of observables") + .def( + "__eq__", + [](const SparseHamiltonian &self, + py::handle other) -> bool { + if (!py::isinstance>(other)) { + return false; + } + auto other_cast = other.cast>(); + return self == other_cast; + }, + "Compare two observables"); +} + /** * @brief Register backend specific adjoint Jacobian methods. * diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp index cf7898340e..6d0d0a94e7 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPU.hpp @@ -215,18 +215,15 @@ class Hamiltonian final : public HamiltonianBase { */ template class SparseHamiltonian final : public SparseHamiltonianBase { + private: + using BaseType = SparseHamiltonianBase; + public: using PrecisionT = typename StateVectorT::PrecisionT; using ComplexT = typename StateVectorT::ComplexT; // cuSparse required index type - using IdxT = - typename std::conditional::value, - int32_t, int64_t>::type; - - private: - using BaseType = SparseHamiltonianBase; + using IdxT = typename BaseType::IdxT; - public: /** * @brief Create a SparseHamiltonian from data, indices and offsets in CSR * format. @@ -291,7 +288,6 @@ class SparseHamiltonian final : public SparseHamiltonianBase { this->indices_.data(), this->data_.data(), static_cast(this->data_.size()), sv.getData(), d_sv_prime->getData(), device_id, stream_id, handle); - sv.updateData(std::move(d_sv_prime)); } }; diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/StateVectorKokkos.hpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/StateVectorKokkos.hpp index a3f0951b24..02064e811f 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/StateVectorKokkos.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/StateVectorKokkos.hpp @@ -17,6 +17,7 @@ */ #pragma once +#include #include #include #include @@ -170,7 +171,7 @@ class StateVectorKokkos final * * @param num_qubits Number of qubits */ - StateVectorKokkos(ComplexT *hostdata_, size_t length, + StateVectorKokkos(ComplexT *hostdata_, std::size_t length, const Kokkos::InitializationSettings &kokkos_args = {}) : StateVectorKokkos(log2(length), kokkos_args) { PL_ABORT_IF_NOT(isPerfectPowerOf2(length), @@ -178,12 +179,20 @@ class StateVectorKokkos final HostToDevice(hostdata_, length); } + StateVectorKokkos(std::complex *hostdata_, std::size_t length, + const Kokkos::InitializationSettings &kokkos_args = {}) + : StateVectorKokkos(log2(length), kokkos_args) { + PL_ABORT_IF_NOT(isPerfectPowerOf2(length), + "The size of provided data must be a power of 2."); + HostToDevice(reinterpret_cast(hostdata_), length); + } + /** * @brief Create a new state vector from data on the host. * * @param num_qubits Number of qubits */ - StateVectorKokkos(const ComplexT *hostdata_, size_t length, + StateVectorKokkos(const ComplexT *hostdata_, std::size_t length, const Kokkos::InitializationSettings &kokkos_args = {}) : StateVectorKokkos(log2(length), kokkos_args) { PL_ABORT_IF_NOT(isPerfectPowerOf2(length), @@ -692,7 +701,7 @@ class StateVectorKokkos final * @param new_data data pointer to new data. * @param new_size size of underlying data storage. */ - void updateData(ComplexT *new_data, size_t new_size) { + void updateData(ComplexT *new_data, std::size_t new_size) { updateData(KokkosVector(new_data, new_size)); } @@ -744,7 +753,7 @@ class StateVectorKokkos final * @brief Copy data from the host space to the device space. * */ - inline void HostToDevice(ComplexT *sv, size_t length) { + inline void HostToDevice(ComplexT *sv, std::size_t length) { Kokkos::deep_copy(*data_, UnmanagedComplexHostView(sv, length)); } @@ -752,7 +761,7 @@ class StateVectorKokkos final * @brief Copy data from the device space to the host space. * */ - inline void DeviceToHost(ComplexT *sv, size_t length) const { + inline void DeviceToHost(ComplexT *sv, std::size_t length) const { Kokkos::deep_copy(UnmanagedComplexHostView(sv, length), *data_); } diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/bindings/LKokkosBindings.hpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/bindings/LKokkosBindings.hpp index bd9d89d72f..6432864c4e 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/bindings/LKokkosBindings.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/bindings/LKokkosBindings.hpp @@ -24,6 +24,7 @@ #include "ConstantUtil.hpp" // lookup #include "GateOperation.hpp" #include "MeasurementsKokkos.hpp" +#include "ObservablesKokkos.hpp" #include "StateVectorKokkos.hpp" #include "TypeList.hpp" #include "Util.hpp" // exp2 @@ -33,6 +34,7 @@ namespace { using namespace Pennylane::Bindings; using namespace Pennylane::LightningKokkos::Algorithms; using namespace Pennylane::LightningKokkos::Measures; +using namespace Pennylane::LightningKokkos::Observables; using Kokkos::InitializationSettings; using Pennylane::LightningKokkos::StateVectorKokkos; using Pennylane::Util::exp2; @@ -214,6 +216,58 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { "Variance of a sparse Hamiltonian."); } +/** + * @brief Register observable classes. + * + * @tparam StateVectorT + * @param m Pybind module + */ +template +void registerBackendSpecificObservables(py::module_ &m) { + using PrecisionT = + typename StateVectorT::PrecisionT; // Statevector's precision. + using ParamT = PrecisionT; // Parameter's data precision + + const std::string bitsize = + std::to_string(sizeof(std::complex) * 8); + + using np_arr_c = py::array_t, py::array::c_style>; + + std::string class_name; + + class_name = "SparseHamiltonianC" + bitsize; + py::class_, + std::shared_ptr>, + Observable>(m, class_name.c_str(), + py::module_local()) + .def(py::init([](const np_arr_c &data, + const std::vector &indices, + const std::vector &indptr, + const std::vector &wires) { + using ComplexT = typename StateVectorT::ComplexT; + const py::buffer_info buffer_data = data.request(); + const auto *data_ptr = static_cast(buffer_data.ptr); + + return SparseHamiltonian{ + std::vector({data_ptr, data_ptr + data.size()}), + indices, indptr, wires}; + })) + .def("__repr__", &SparseHamiltonian::getObsName) + .def("get_wires", &SparseHamiltonian::getWires, + "Get wires of observables") + .def( + "__eq__", + [](const SparseHamiltonian &self, + py::handle other) -> bool { + if (!py::isinstance>(other)) { + return false; + } + auto other_cast = other.cast>(); + return self == other_cast; + }, + "Compare two observables"); +} + /** * @brief Register backend specific adjoint Jacobian methods. * diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.cpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.cpp index d90f3e6019..66192b934a 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.cpp @@ -28,3 +28,6 @@ template class Observables::TensorProdObs>; template class Observables::Hamiltonian>; template class Observables::Hamiltonian>; + +template class Observables::SparseHamiltonian>; +template class Observables::SparseHamiltonian>; diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp index a0371df7f2..c3fae6b3ea 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/ObservablesKokkos.hpp @@ -31,6 +31,7 @@ namespace { using namespace Pennylane::Util; using namespace Pennylane::Observables; using Pennylane::LightningKokkos::StateVectorKokkos; +using Pennylane::LightningKokkos::Util::SparseMV_Kokkos; } // namespace /// @endcond @@ -199,6 +200,76 @@ class Hamiltonian final : public HamiltonianBase { } }; +/** + * @brief Sparse representation of Hamiltonian + * + */ +template +class SparseHamiltonian final : public SparseHamiltonianBase { + private: + using BaseType = SparseHamiltonianBase; + + public: + using PrecisionT = typename StateVectorT::PrecisionT; + using ComplexT = typename StateVectorT::ComplexT; + using IdxT = typename BaseType::IdxT; + + /** + * @brief Create a SparseHamiltonian from data, indices and offsets in CSR + * format. + * + * @param data Arguments to construct data + * @param indices Arguments to construct indices + * @param offsets Arguments to construct offsets + * @param wires Arguments to construct wires + */ + template + explicit SparseHamiltonian(T1 &&data, T2 &&indices, T3 &&offsets, + T4 &&wires) + : BaseType{data, indices, offsets, wires} {} + + /** + * @brief Convenient wrapper for the constructor as the constructor does not + * convert the std::shared_ptr with a derived class correctly. + * + * This function is useful as std::make_shared does not handle + * brace-enclosed initializer list correctly. + * + * @param data Argument to construct data + * @param indices Argument to construct indices + * @param offsets Argument to construct ofsets + * @param wires Argument to construct wires + */ + static auto create(std::initializer_list data, + std::initializer_list indices, + std::initializer_list offsets, + std::initializer_list wires) + -> std::shared_ptr> { + return std::shared_ptr>( + new SparseHamiltonian{ + std::move(data), std::move(indices), std::move(offsets), + std::move(wires)}); + } + + /** + * @brief Updates the statevector SV:->SV', where SV' = a*H*SV, and where H + * is a sparse Hamiltonian. + * + */ + void applyInPlace(StateVectorT &sv) const override { + PL_ABORT_IF_NOT(this->wires_.size() == sv.getNumQubits(), + "SparseH wire count does not match state-vector size"); + StateVectorT d_sv_prime(sv.getNumQubits()); + + SparseMV_Kokkos( + sv.getView(), d_sv_prime.getView(), this->offsets_.data(), + this->offsets_.size(), this->indices_.data(), this->data_.data(), + this->data_.size()); + + sv.updateData(d_sv_prime); + } +}; + /// @cond DEV namespace detail { using Pennylane::LightningKokkos::Util::axpy_Kokkos; diff --git a/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp b/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp index 190a9c6525..7cd5a1ee2e 100644 --- a/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_qubit/bindings/LQubitBindings.hpp @@ -180,6 +180,15 @@ void registerBackendSpecificMeasurements(PyClass &pyclass) { }); } +/** + * @brief Register backend specific observables. + * + * @tparam StateVectorT + * @param m Pybind module + */ +template +void registerBackendSpecificObservables([[maybe_unused]] py::module_ &m) {} + /** * @brief Register Vector Jacobian Product. */ diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py index 69fef4ed11..c9b338cdc3 100644 --- a/tests/test_adjoint_jacobian.py +++ b/tests/test_adjoint_jacobian.py @@ -25,6 +25,7 @@ from pennylane import QNode, qnode from pennylane import qchem + I, X, Y, Z = ( np.eye(2), qml.PauliX.compute_matrix(), @@ -918,9 +919,10 @@ def circuit(params): assert np.allclose(qml.grad(circuit_ld)(params), qml.grad(circuit_dq)(params), tol) +@pytest.mark.skipif(not ld._CPP_BINARY_AVAILABLE, reason="Lightning binary required") @pytest.mark.skipif( - device_name != "lightning.gpu" or not ld._CPP_BINARY_AVAILABLE, - reason="Lightning binary required", + device_name not in ["lightning.gpu", "lightning.kokkos"], + reason="SparseHamiltonian only supported by Lightning-Kokkos", ) def test_tape_qchem_sparse(tol): """Tests the circuit Ansatz with a QChem Hamiltonian produces correct results""" @@ -950,6 +952,64 @@ def circuit(params): assert np.allclose(qml.grad(circuit_ld)(params), qml.grad(circuit_dq)(params), tol) +custom_wires = ["alice", 3.14, -1, 0] + + +@pytest.mark.skipif(not ld._CPP_BINARY_AVAILABLE, reason="Lightning binary required") +@pytest.mark.skipif( + device_name not in ["lightning.gpu", "lightning.kokkos"], + reason="SparseHamiltonian only supported by Lightning-Kokkos", +) +@pytest.mark.parametrize( + "returns", + [ + qml.SparseHamiltonian( + qml.Hamiltonian( + [0.1], + [qml.PauliX(wires=custom_wires[0]) @ qml.PauliZ(wires=custom_wires[1])], + ).sparse_matrix(custom_wires), + wires=custom_wires, + ), + qml.SparseHamiltonian( + qml.Hamiltonian( + [2.0], + [qml.PauliX(wires=custom_wires[2]) @ qml.PauliZ(wires=custom_wires[0])], + ).sparse_matrix(custom_wires), + wires=custom_wires, + ), + qml.SparseHamiltonian( + qml.Hamiltonian( + [1.1], + [qml.PauliX(wires=custom_wires[0]) @ qml.PauliZ(wires=custom_wires[2])], + ).sparse_matrix(custom_wires), + wires=custom_wires, + ), + ], +) +def test_adjoint_SparseHamiltonian(returns): + """Integration tests that compare to default.qubit for a large circuit containing parametrized + operations and when using custom wire labels""" + + dev_kokkos = qml.device(device_name, wires=custom_wires) + dev_default = qml.device("default.qubit", wires=custom_wires) + + def circuit(params): + circuit_ansatz(params, wires=custom_wires) + return qml.expval(returns) + + n_params = 30 + np.random.seed(1337) + params = np.random.rand(n_params) + + qnode_kokkos = qml.QNode(circuit, dev_kokkos, diff_method="adjoint") + qnode_default = qml.QNode(circuit, dev_default, diff_method="parameter-shift") + + j_kokkos = qml.jacobian(qnode_kokkos)(params) + j_default = qml.jacobian(qnode_default)(params) + + assert np.allclose(j_kokkos, j_default) + + @pytest.mark.parametrize( "returns", [ @@ -1030,9 +1090,6 @@ def casted_to_array_batched(params): assert np.allclose(j_def, j_lightning_batched) -custom_wires = ["alice", 3.14, -1, 0] - - @pytest.mark.parametrize( "returns", [ diff --git a/tests/test_serialize.py b/tests/test_serialize.py index ab6df3c26e..92938b3406 100644 --- a/tests/test_serialize.py +++ b/tests/test_serialize.py @@ -34,6 +34,8 @@ TensorProdObsC128, HamiltonianC64, HamiltonianC128, + SparseHamiltonianC64, + SparseHamiltonianC128, ) elif device_name == "lightning.gpu": from pennylane_lightning.lightning_gpu_ops.observables import ( @@ -45,6 +47,8 @@ TensorProdObsC128, HamiltonianC64, HamiltonianC128, + SparseHamiltonianC64, + SparseHamiltonianC128, ) else: from pennylane_lightning.lightning_qubit_ops.observables import ( @@ -66,34 +70,41 @@ def test_wrong_device_name(): QuantumScriptSerializer("thunder.qubit") -@pytest.mark.parametrize( - "obs,obs_type", - [ - (qml.PauliZ(0), NamedObsC128), - (qml.PauliZ(0) @ qml.PauliZ(1), TensorProdObsC128), - (qml.Hadamard(0), NamedObsC128), - (qml.Hermitian(np.eye(2), wires=0), HermitianObsC128), - ( - qml.PauliZ(0) @ qml.Hadamard(1) @ (0.1 * (qml.PauliZ(2) + qml.PauliX(3))), - HamiltonianC128, - ), +obs_list = [ + (qml.PauliZ(0), NamedObsC128), + (qml.PauliZ(0) @ qml.PauliZ(1), TensorProdObsC128), + (qml.Hadamard(0), NamedObsC128), + (qml.Hermitian(np.eye(2), wires=0), HermitianObsC128), + ( + qml.PauliZ(0) @ qml.Hadamard(1) @ (0.1 * (qml.PauliZ(2) + qml.PauliX(3))), + HamiltonianC128, + ), + ( ( - ( - qml.Hermitian(np.eye(2), wires=0) - @ qml.Hermitian(np.eye(2), wires=1) - @ qml.Projector([0], wires=2) - ), - TensorProdObsC128, + qml.Hermitian(np.eye(2), wires=0) + @ qml.Hermitian(np.eye(2), wires=1) + @ qml.Projector([0], wires=2) ), + TensorProdObsC128, + ), + ( + qml.PauliZ(0) @ qml.Hermitian(np.eye(2), wires=1) @ qml.Projector([0], wires=2), + TensorProdObsC128, + ), + (qml.Projector([0], wires=0), HermitianObsC128), + (qml.Hamiltonian([1], [qml.PauliZ(0)]), HamiltonianC128), + (qml.sum(qml.Hadamard(0), qml.PauliX(1)), HermitianObsC128), +] +if device_name in ["lightning.gpu", "lightning.kokkos"]: + obs_list += [ ( - qml.PauliZ(0) @ qml.Hermitian(np.eye(2), wires=1) @ qml.Projector([0], wires=2), - TensorProdObsC128, + qml.SparseHamiltonian(qml.Hamiltonian([1], [qml.PauliZ(0)]).sparse_matrix(), wires=[0]), + SparseHamiltonianC128, ), - (qml.Projector([0], wires=0), HermitianObsC128), - (qml.Hamiltonian([1], [qml.PauliZ(0)]), HamiltonianC128), - (qml.sum(qml.Hadamard(0), qml.PauliX(1)), HermitianObsC128), - ], -) + ] + + +@pytest.mark.parametrize("obs,obs_type", obs_list) def test_obs_returns_expected_type(obs, obs_type): """Tests that observables get serialized to the expected type.""" assert isinstance( From 8b2f752809efdf9f28f8b7f0e2d622c15b713c40 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 20:19:53 +0000 Subject: [PATCH 21/43] Fix clang tidy --- pennylane_lightning/core/src/observables/Observables.hpp | 4 ++-- .../lightning_qubit/observables/ObservablesLQubit.hpp | 2 ++ 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index 37b2ad4e1b..0ea9cbedd4 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -414,7 +414,6 @@ class HamiltonianBase : public Observable { } }; -// NOLINTBEGIN /** * @brief Sparse representation of SparseHamiltonian * @@ -488,10 +487,12 @@ class SparseHamiltonianBase : public Observable { std::initializer_list arg3, std::initializer_list arg4) -> std::shared_ptr> { + // NOLINTBEGIN(*-move-const-arg) return std::shared_ptr>( new SparseHamiltonianBase{ std::move(arg1), std::move(arg2), std::move(arg3), std::move(arg4)}); + // NOLINTEND(*-move-const-arg) } void applyInPlace([[maybe_unused]] StateVectorT &sv) const override { @@ -525,6 +526,5 @@ class SparseHamiltonianBase : public Observable { return wires_; }; }; -// NOLINTEND } // namespace Pennylane::Observables \ No newline at end of file diff --git a/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.hpp b/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.hpp index c68d119526..b659969037 100644 --- a/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_qubit/observables/ObservablesLQubit.hpp @@ -405,10 +405,12 @@ class SparseHamiltonian final : public SparseHamiltonianBase { std::initializer_list offsets, std::initializer_list wires) -> std::shared_ptr> { + // NOLINTBEGIN(*-move-const-arg) return std::shared_ptr>( new SparseHamiltonian{ std::move(data), std::move(indices), std::move(offsets), std::move(wires)}); + // NOLINTEND(*-move-const-arg) } /** From cf2866bdf16e9a6f49c1a1386e09ec663be18c71 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 20:23:28 +0000 Subject: [PATCH 22/43] Comment workflows but tidy. --- .github/workflows/changelog_reminder.yml | 2 +- .github/workflows/tests_gpu.yml | 2 +- .github/workflows/tests_gpu_cu11.yml | 2 +- .github/workflows/tests_linux.yml | 2 +- .github/workflows/tests_linux_x86_mpi.yml | 2 +- .github/workflows/tests_windows.yml | 2 +- .github/workflows/tests_without_binary.yml | 2 +- .github/workflows/update_dev_version.yml | 2 +- .github/workflows/wheel_linux_x86_64.yml | 2 +- .github/workflows/wheel_linux_x86_64_cu11.yml | 2 +- .github/workflows/wheel_macos_arm64.yml | 2 +- .github/workflows/wheel_macos_x86_64.yml | 2 +- .github/workflows/wheel_noarch.yml | 2 +- .github/workflows/wheel_win_x86_64.yml | 2 +- 14 files changed, 14 insertions(+), 14 deletions(-) diff --git a/.github/workflows/changelog_reminder.yml b/.github/workflows/changelog_reminder.yml index 2eb934c165..4d3b120541 100644 --- a/.github/workflows/changelog_reminder.yml +++ b/.github/workflows/changelog_reminder.yml @@ -1,5 +1,5 @@ on: - pull_request: + #pull_request: types: [opened, ready_for_review] name: Changelog Reminder diff --git a/.github/workflows/tests_gpu.yml b/.github/workflows/tests_gpu.yml index bbf240d91c..f5747f3763 100644 --- a/.github/workflows/tests_gpu.yml +++ b/.github/workflows/tests_gpu.yml @@ -1,6 +1,6 @@ name: Testing (GPU) on: - pull_request: + #pull_request: push: branches: - master diff --git a/.github/workflows/tests_gpu_cu11.yml b/.github/workflows/tests_gpu_cu11.yml index cb15aca14e..16715ed7eb 100644 --- a/.github/workflows/tests_gpu_cu11.yml +++ b/.github/workflows/tests_gpu_cu11.yml @@ -1,6 +1,6 @@ name: Testing::Linux::x86_64::LGPU on: - pull_request: + #pull_request: push: branches: - master diff --git a/.github/workflows/tests_linux.yml b/.github/workflows/tests_linux.yml index ac8a05c593..7fcb15d701 100644 --- a/.github/workflows/tests_linux.yml +++ b/.github/workflows/tests_linux.yml @@ -10,7 +10,7 @@ on: type: string required: true description: The version of PennyLane to use. Valid values are either 'stable' (most recent git-tag) or 'latest' (most recent commit from master) - pull_request: + #pull_request: push: branches: - master diff --git a/.github/workflows/tests_linux_x86_mpi.yml b/.github/workflows/tests_linux_x86_mpi.yml index ff6e54aa7c..8c34bd80a3 100644 --- a/.github/workflows/tests_linux_x86_mpi.yml +++ b/.github/workflows/tests_linux_x86_mpi.yml @@ -14,7 +14,7 @@ on: push: branches: - main - pull_request: + #pull_request: env: COVERAGE_FLAGS: "--cov=pennylane_lightning --cov-report=term-missing --cov-report=xml:./coverage.xml --no-flaky-report -p no:warnings --tb=native" diff --git a/.github/workflows/tests_windows.yml b/.github/workflows/tests_windows.yml index 36f34bf23b..98cd80b5ee 100644 --- a/.github/workflows/tests_windows.yml +++ b/.github/workflows/tests_windows.yml @@ -3,7 +3,7 @@ on: push: branches: - master - pull_request: + #pull_request: concurrency: group: tests_windows-${{ github.ref }} diff --git a/.github/workflows/tests_without_binary.yml b/.github/workflows/tests_without_binary.yml index bb21da4cfb..845b215e96 100644 --- a/.github/workflows/tests_without_binary.yml +++ b/.github/workflows/tests_without_binary.yml @@ -10,7 +10,7 @@ on: type: string required: true description: The version of PennyLane to use. Valid values are either 'stable' (most recent git-tag) or 'latest' (most recent commit from master) - pull_request: + #pull_request: push: branches: - master diff --git a/.github/workflows/update_dev_version.yml b/.github/workflows/update_dev_version.yml index bfa66e1e8b..02062bdb5c 100644 --- a/.github/workflows/update_dev_version.yml +++ b/.github/workflows/update_dev_version.yml @@ -1,6 +1,6 @@ name: Update dev version automatically on: - pull_request: + #pull_request: jobs: update-dev-version: diff --git a/.github/workflows/wheel_linux_x86_64.yml b/.github/workflows/wheel_linux_x86_64.yml index 50a3f81c72..abb74ca216 100644 --- a/.github/workflows/wheel_linux_x86_64.yml +++ b/.github/workflows/wheel_linux_x86_64.yml @@ -9,7 +9,7 @@ env: GCC_VERSION: 11 on: - pull_request: + #pull_request: push: branches: - master diff --git a/.github/workflows/wheel_linux_x86_64_cu11.yml b/.github/workflows/wheel_linux_x86_64_cu11.yml index 2b270c4b0e..c4dd47bc68 100644 --- a/.github/workflows/wheel_linux_x86_64_cu11.yml +++ b/.github/workflows/wheel_linux_x86_64_cu11.yml @@ -11,7 +11,7 @@ env: CUDA_VERSION_MINOR: 5 on: - pull_request: + #pull_request: push: branches: - master diff --git a/.github/workflows/wheel_macos_arm64.yml b/.github/workflows/wheel_macos_arm64.yml index 4e0449038c..dde609f04a 100644 --- a/.github/workflows/wheel_macos_arm64.yml +++ b/.github/workflows/wheel_macos_arm64.yml @@ -6,7 +6,7 @@ name: Wheel::MacOS::ARM # **Who does it impact**: Wheels to be uploaded to PyPI. on: - pull_request: + #pull_request: push: branches: - master diff --git a/.github/workflows/wheel_macos_x86_64.yml b/.github/workflows/wheel_macos_x86_64.yml index 82e6220524..d5e1d53441 100644 --- a/.github/workflows/wheel_macos_x86_64.yml +++ b/.github/workflows/wheel_macos_x86_64.yml @@ -6,7 +6,7 @@ name: Wheel::MacOS::Intel # **Who does it impact**: Wheels to be uploaded to PyPI. on: - pull_request: + #pull_request: push: branches: - master diff --git a/.github/workflows/wheel_noarch.yml b/.github/workflows/wheel_noarch.yml index d3e6622730..000a800cf2 100644 --- a/.github/workflows/wheel_noarch.yml +++ b/.github/workflows/wheel_noarch.yml @@ -6,7 +6,7 @@ name: Wheel::Any::None # **Who does it impact**: Wheels to be uploaded to PyPI. on: - pull_request: + #pull_request: push: branches: - master diff --git a/.github/workflows/wheel_win_x86_64.yml b/.github/workflows/wheel_win_x86_64.yml index 0d0b0989ea..9654103161 100644 --- a/.github/workflows/wheel_win_x86_64.yml +++ b/.github/workflows/wheel_win_x86_64.yml @@ -6,7 +6,7 @@ name: Wheel::Windows::x86_64 # **Who does it impact**: Wheels to be uploaded to PyPI. on: - pull_request: + #pull_request: push: branches: - master From ccace74d0767d693b7c93e28e0a60a858d8b96aa Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 20:25:39 +0000 Subject: [PATCH 23/43] Fix tidy warn --- pennylane_lightning/core/src/observables/Observables.hpp | 9 ++------- 1 file changed, 2 insertions(+), 7 deletions(-) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index 0ea9cbedd4..67d77f9d32 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -443,13 +443,8 @@ class SparseHamiltonianBase : public Observable { isEqual(const Observable &other) const override { const auto &other_cast = static_cast &>(other); - - if (data_ != other_cast.data_ || indices_ != other_cast.indices_ || - offsets_ != other_cast.offsets_) { - return false; - } - - return true; + return data_ != other_cast.data_ || indices_ != other_cast.indices_ || + offsets_ != other_cast.offsets_; } public: From ae615e47b5a8f7683b395c0d4b072846acb3e167 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 20:42:13 +0000 Subject: [PATCH 24/43] Add override to sp::getWires --- pennylane_lightning/core/src/observables/Observables.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index 67d77f9d32..f2a709c84f 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -517,7 +517,7 @@ class SparseHamiltonianBase : public Observable { /** * @brief Get the wires the observable applies to. */ - [[nodiscard]] auto getWires() const -> std::vector { + [[nodiscard]] auto getWires() const -> std::vector override { return wires_; }; }; From 4d108f3a755ebbe505915d8a771eb22434e812f0 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 21:03:45 +0000 Subject: [PATCH 25/43] Restore triggers --- .github/workflows/changelog_reminder.yml | 2 +- .github/workflows/tests_gpu.yml | 2 +- .github/workflows/tests_gpu_cu11.yml | 2 +- .github/workflows/tests_linux.yml | 2 +- .github/workflows/tests_linux_x86_mpi.yml | 2 +- .github/workflows/tests_windows.yml | 2 +- .github/workflows/tests_without_binary.yml | 2 +- .github/workflows/update_dev_version.yml | 2 +- .github/workflows/wheel_linux_x86_64.yml | 2 +- .github/workflows/wheel_linux_x86_64_cu11.yml | 2 +- .github/workflows/wheel_macos_arm64.yml | 2 +- .github/workflows/wheel_macos_x86_64.yml | 2 +- .github/workflows/wheel_noarch.yml | 2 +- .github/workflows/wheel_win_x86_64.yml | 2 +- 14 files changed, 14 insertions(+), 14 deletions(-) diff --git a/.github/workflows/changelog_reminder.yml b/.github/workflows/changelog_reminder.yml index 4d3b120541..2eb934c165 100644 --- a/.github/workflows/changelog_reminder.yml +++ b/.github/workflows/changelog_reminder.yml @@ -1,5 +1,5 @@ on: - #pull_request: + pull_request: types: [opened, ready_for_review] name: Changelog Reminder diff --git a/.github/workflows/tests_gpu.yml b/.github/workflows/tests_gpu.yml index f5747f3763..bbf240d91c 100644 --- a/.github/workflows/tests_gpu.yml +++ b/.github/workflows/tests_gpu.yml @@ -1,6 +1,6 @@ name: Testing (GPU) on: - #pull_request: + pull_request: push: branches: - master diff --git a/.github/workflows/tests_gpu_cu11.yml b/.github/workflows/tests_gpu_cu11.yml index 16715ed7eb..cb15aca14e 100644 --- a/.github/workflows/tests_gpu_cu11.yml +++ b/.github/workflows/tests_gpu_cu11.yml @@ -1,6 +1,6 @@ name: Testing::Linux::x86_64::LGPU on: - #pull_request: + pull_request: push: branches: - master diff --git a/.github/workflows/tests_linux.yml b/.github/workflows/tests_linux.yml index 7fcb15d701..ac8a05c593 100644 --- a/.github/workflows/tests_linux.yml +++ b/.github/workflows/tests_linux.yml @@ -10,7 +10,7 @@ on: type: string required: true description: The version of PennyLane to use. Valid values are either 'stable' (most recent git-tag) or 'latest' (most recent commit from master) - #pull_request: + pull_request: push: branches: - master diff --git a/.github/workflows/tests_linux_x86_mpi.yml b/.github/workflows/tests_linux_x86_mpi.yml index 8c34bd80a3..ff6e54aa7c 100644 --- a/.github/workflows/tests_linux_x86_mpi.yml +++ b/.github/workflows/tests_linux_x86_mpi.yml @@ -14,7 +14,7 @@ on: push: branches: - main - #pull_request: + pull_request: env: COVERAGE_FLAGS: "--cov=pennylane_lightning --cov-report=term-missing --cov-report=xml:./coverage.xml --no-flaky-report -p no:warnings --tb=native" diff --git a/.github/workflows/tests_windows.yml b/.github/workflows/tests_windows.yml index 98cd80b5ee..36f34bf23b 100644 --- a/.github/workflows/tests_windows.yml +++ b/.github/workflows/tests_windows.yml @@ -3,7 +3,7 @@ on: push: branches: - master - #pull_request: + pull_request: concurrency: group: tests_windows-${{ github.ref }} diff --git a/.github/workflows/tests_without_binary.yml b/.github/workflows/tests_without_binary.yml index 845b215e96..bb21da4cfb 100644 --- a/.github/workflows/tests_without_binary.yml +++ b/.github/workflows/tests_without_binary.yml @@ -10,7 +10,7 @@ on: type: string required: true description: The version of PennyLane to use. Valid values are either 'stable' (most recent git-tag) or 'latest' (most recent commit from master) - #pull_request: + pull_request: push: branches: - master diff --git a/.github/workflows/update_dev_version.yml b/.github/workflows/update_dev_version.yml index 02062bdb5c..bfa66e1e8b 100644 --- a/.github/workflows/update_dev_version.yml +++ b/.github/workflows/update_dev_version.yml @@ -1,6 +1,6 @@ name: Update dev version automatically on: - #pull_request: + pull_request: jobs: update-dev-version: diff --git a/.github/workflows/wheel_linux_x86_64.yml b/.github/workflows/wheel_linux_x86_64.yml index abb74ca216..50a3f81c72 100644 --- a/.github/workflows/wheel_linux_x86_64.yml +++ b/.github/workflows/wheel_linux_x86_64.yml @@ -9,7 +9,7 @@ env: GCC_VERSION: 11 on: - #pull_request: + pull_request: push: branches: - master diff --git a/.github/workflows/wheel_linux_x86_64_cu11.yml b/.github/workflows/wheel_linux_x86_64_cu11.yml index c4dd47bc68..2b270c4b0e 100644 --- a/.github/workflows/wheel_linux_x86_64_cu11.yml +++ b/.github/workflows/wheel_linux_x86_64_cu11.yml @@ -11,7 +11,7 @@ env: CUDA_VERSION_MINOR: 5 on: - #pull_request: + pull_request: push: branches: - master diff --git a/.github/workflows/wheel_macos_arm64.yml b/.github/workflows/wheel_macos_arm64.yml index dde609f04a..4e0449038c 100644 --- a/.github/workflows/wheel_macos_arm64.yml +++ b/.github/workflows/wheel_macos_arm64.yml @@ -6,7 +6,7 @@ name: Wheel::MacOS::ARM # **Who does it impact**: Wheels to be uploaded to PyPI. on: - #pull_request: + pull_request: push: branches: - master diff --git a/.github/workflows/wheel_macos_x86_64.yml b/.github/workflows/wheel_macos_x86_64.yml index d5e1d53441..82e6220524 100644 --- a/.github/workflows/wheel_macos_x86_64.yml +++ b/.github/workflows/wheel_macos_x86_64.yml @@ -6,7 +6,7 @@ name: Wheel::MacOS::Intel # **Who does it impact**: Wheels to be uploaded to PyPI. on: - #pull_request: + pull_request: push: branches: - master diff --git a/.github/workflows/wheel_noarch.yml b/.github/workflows/wheel_noarch.yml index 000a800cf2..d3e6622730 100644 --- a/.github/workflows/wheel_noarch.yml +++ b/.github/workflows/wheel_noarch.yml @@ -6,7 +6,7 @@ name: Wheel::Any::None # **Who does it impact**: Wheels to be uploaded to PyPI. on: - #pull_request: + pull_request: push: branches: - master diff --git a/.github/workflows/wheel_win_x86_64.yml b/.github/workflows/wheel_win_x86_64.yml index 9654103161..0d0b0989ea 100644 --- a/.github/workflows/wheel_win_x86_64.yml +++ b/.github/workflows/wheel_win_x86_64.yml @@ -6,7 +6,7 @@ name: Wheel::Windows::x86_64 # **Who does it impact**: Wheels to be uploaded to PyPI. on: - #pull_request: + pull_request: push: branches: - master From f015e2f3ff5b33327e97edc849d81a9d7a1635d6 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Mon, 23 Oct 2023 23:48:29 +0000 Subject: [PATCH 26/43] Update tests_linux_x86_mpi.yml --- .github/workflows/tests_linux_x86_mpi.yml | 82 +++++++++++++++++++---- 1 file changed, 69 insertions(+), 13 deletions(-) diff --git a/.github/workflows/tests_linux_x86_mpi.yml b/.github/workflows/tests_linux_x86_mpi.yml index ff6e54aa7c..e43e9ca241 100644 --- a/.github/workflows/tests_linux_x86_mpi.yml +++ b/.github/workflows/tests_linux_x86_mpi.yml @@ -133,7 +133,7 @@ jobs: for file in *runner_mpi ; do /opt/mpi/${{ matrix.mpilib }}/bin/mpirun -np 2 ./$file --order lex --reporter junit --out ./tests/results/report_$file.xml; done; lcov --directory . -b ../pennylane_lightning/src --capture --output-file coverage.info lcov --remove coverage.info '/usr/*' --output-file coverage.info - mv coverage.info coverage-${{ github.job }}-lightning_gpu.info + mv coverage.info coverage-${{ github.job }}-lightning_gpu_${{ matrix.mpilib }}.info - name: Upload test results uses: actions/upload-artifact@v3 @@ -148,16 +148,13 @@ jobs: with: if-no-files-found: error name: ubuntu-codecov-results-cpp - path: ./Build/coverage-${{ github.job }}-lightning_gpu.info + path: ./Build/coverage-${{ github.job }}-lightning_gpu_${{ matrix.mpilib }}.info - name: Cleanup if: always() run: | rm -rf ${{ steps.setup_venv.outputs.venv_name }} - rm -rf * - rm -rf .git - rm -rf .gitignore - rm -rf .github + rm -rf * .git .gitignore .github pip cache purge @@ -169,7 +166,6 @@ jobs: - x64 - ubuntu-22.04 - multi-gpu - needs: ["cpp_tests"] strategy: max-parallel: 1 matrix: @@ -249,15 +245,75 @@ jobs: - name: Run PennyLane-Lightning-GPU unit tests run: | source /etc/profile.d/modules.sh && module use /opt/modules/ && module load ${{ matrix.mpilib }} - PL_DEVICE=lightning.gpu python -m pytest ./tests/ - PL_DEVICE=lightning.gpu /opt/mpi/${{ matrix.mpilib }}/bin/mpirun -np 2 python -m pytest ./mpitests/test_adjoint_jacobian.py + PL_DEVICE=lightning.gpu python -m pytest ./tests/ $COVERAGE_FLAGS + PL_DEVICE=lightning.gpu /opt/mpi/${{ matrix.mpilib }}/bin/mpirun -np 2 python -m pytest ./mpitests $COVERAGE_FLAGS + # PL_DEVICE=lightning.gpu /opt/mpi/${{ matrix.mpilib }}/bin/mpirun --oversubscribe -n 4 pytest -s -x mpitests/test_device.py -k test_create_device $COVERAGE_FLAGS + mv coverage.xml coverage-${{ github.job }}-lightning_gpu_${{ matrix.mpilib }}.xml + + - name: Upload code coverage results + uses: actions/upload-artifact@v3 + with: + name: ubuntu-codecov-results-python + path: coverage-${{ github.job }}-lightning_gpu_${{ matrix.mpilib }}.xml + if-no-files-found: error + + - name: Cleanup + if: always() + run: | + rm -rf ${{ steps.setup_venv.outputs.venv_name }} + rm -rf * .git .gitignore .github + pip cache purge + + upload-to-codecov-linux-cpp: + needs: ["cpp_tests"] + name: Upload coverage data to codecov + runs-on: ubuntu-latest + steps: + - name: Checkout + uses: actions/checkout@v3 + + - name: Download coverage reports + uses: actions/download-artifact@v3 + with: + name: ubuntu-codecov-results-cpp + + - name: Upload to Codecov + uses: codecov/codecov-action@v3 + with: + fail_ci_if_error: true + verbose: true + token: ${{ secrets.CODECOV_TOKEN }} + + - name: Cleanup + if: always() + run: | + rm -rf ${{ steps.setup_venv.outputs.venv_name }} + rm -rf * .git .gitignore .github + pip cache purge + + upload-to-codecov-linux-python: + needs: ["python_tests"] + name: Upload coverage data to codecov + runs-on: ubuntu-latest + steps: + - name: Checkout + uses: actions/checkout@v3 + + - name: Download coverage reports + uses: actions/download-artifact@v3 + with: + name: ubuntu-codecov-results-python + + - name: Upload to Codecov + uses: codecov/codecov-action@v3 + with: + fail_ci_if_error: true + verbose: true + token: ${{ secrets.CODECOV_TOKEN }} - name: Cleanup if: always() run: | rm -rf ${{ steps.setup_venv.outputs.venv_name }} - rm -rf * - rm -rf .git - rm -rf .gitignore - rm -rf .github + rm -rf * .git .gitignore .github pip cache purge From ca93578e33eb1293f009b5dd2f58c9014eb943f1 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 12:45:58 +0000 Subject: [PATCH 27/43] Add constructibility tests. --- .../core/src/observables/Observables.hpp | 2 ++ .../observables/tests/Test_ObservablesGPU.cpp | 14 ++++++++++++++ .../observables/tests/Test_ObservablesKokkos.cpp | 14 ++++++++++++++ .../observables/tests/Test_ObservablesLQubit.cpp | 15 +++++++++++++++ 4 files changed, 45 insertions(+) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index f2a709c84f..b1c5665f1c 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -439,6 +439,7 @@ class SparseHamiltonianBase : public Observable { std::vector wires_; private: + // LCOV_EXCL_START [[nodiscard]] bool isEqual(const Observable &other) const override { const auto &other_cast = @@ -446,6 +447,7 @@ class SparseHamiltonianBase : public Observable { return data_ != other_cast.data_ || indices_ != other_cast.indices_ || offsets_ != other_cast.offsets_; } + // LCOV_EXCL_STOP public: /** diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/Test_ObservablesGPU.cpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/Test_ObservablesGPU.cpp index 816b285695..398f664ffc 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/Test_ObservablesGPU.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/tests/Test_ObservablesGPU.cpp @@ -154,6 +154,20 @@ TEMPLATE_PRODUCT_TEST_CASE("Hamiltonian", "[Observables]", } } +TEMPLATE_PRODUCT_TEST_CASE("SparseHamiltonian", "[Observables]", + (StateVectorCudaManaged), (float, double)) { + using StateVectorT = TestType; + using SparseHamiltonianT = SparseHamiltonian; + + SECTION("Copy constructibility") { + REQUIRE(std::is_copy_constructible_v); + } + + SECTION("Move constructibility") { + REQUIRE(std::is_move_constructible_v); + } +} + TEMPLATE_PRODUCT_TEST_CASE("Observables::HermitianHasher", "[Observables]", (StateVectorCudaManaged), (float, double)) { using StateVectorT = TestType; diff --git a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/tests/Test_ObservablesKokkos.cpp b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/tests/Test_ObservablesKokkos.cpp index aa616d2e33..5d402b1e2d 100644 --- a/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/tests/Test_ObservablesKokkos.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_kokkos/observables/tests/Test_ObservablesKokkos.cpp @@ -153,6 +153,20 @@ TEMPLATE_PRODUCT_TEST_CASE("Hamiltonian", "[Observables]", (StateVectorKokkos), } } +TEMPLATE_PRODUCT_TEST_CASE("SparseHamiltonian", "[Observables]", + (StateVectorKokkos), (float, double)) { + using StateVectorT = TestType; + using SparseHamiltonianT = SparseHamiltonian; + + SECTION("Copy constructibility") { + REQUIRE(std::is_copy_constructible_v); + } + + SECTION("Move constructibility") { + REQUIRE(std::is_move_constructible_v); + } +} + TEMPLATE_PRODUCT_TEST_CASE("Hamiltonian::ApplyInPlace", "[Observables]", (StateVectorKokkos), (float, double)) { using StateVectorT = TestType; diff --git a/pennylane_lightning/core/src/simulators/lightning_qubit/observables/tests/Test_ObservablesLQubit.cpp b/pennylane_lightning/core/src/simulators/lightning_qubit/observables/tests/Test_ObservablesLQubit.cpp index 624248da20..4ef59b04c5 100644 --- a/pennylane_lightning/core/src/simulators/lightning_qubit/observables/tests/Test_ObservablesLQubit.cpp +++ b/pennylane_lightning/core/src/simulators/lightning_qubit/observables/tests/Test_ObservablesLQubit.cpp @@ -159,6 +159,21 @@ TEMPLATE_PRODUCT_TEST_CASE("Hamiltonian", "[Observables]", } } +TEMPLATE_PRODUCT_TEST_CASE("SparseHamiltonian", "[Observables]", + (StateVectorLQubitManaged, StateVectorLQubitRaw), + (float, double)) { + using StateVectorT = TestType; + using SparseHamiltonianT = SparseHamiltonian; + + SECTION("Copy constructibility") { + REQUIRE(std::is_copy_constructible_v); + } + + SECTION("Move constructibility") { + REQUIRE(std::is_move_constructible_v); + } +} + TEMPLATE_PRODUCT_TEST_CASE("Hamiltonian::ApplyInPlace", "[Observables]", (StateVectorLQubitManaged, StateVectorLQubitRaw), (float, double)) { From 62850d7f23a22b2a27523a1886aa2f8ae0ff7227 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 13:16:02 +0000 Subject: [PATCH 28/43] Move L-Kokkos-CUDA tests to workflow call, called from tests_gpu_cu11.yml. --- .github/workflows/tests_gpu_cu11.yml | 7 +++++++ .github/workflows/{tests_gpu.yml => tests_gpu_kokkos.yml} | 5 +---- 2 files changed, 8 insertions(+), 4 deletions(-) rename .github/workflows/{tests_gpu.yml => tests_gpu_kokkos.yml} (99%) diff --git a/.github/workflows/tests_gpu_cu11.yml b/.github/workflows/tests_gpu_cu11.yml index cb15aca14e..5fb82ca491 100644 --- a/.github/workflows/tests_gpu_cu11.yml +++ b/.github/workflows/tests_gpu_cu11.yml @@ -17,7 +17,14 @@ concurrency: cancel-in-progress: true jobs: + + # Workflow call to avoid toe stepping + test_lightning_kokkos: + name: "Test Lightning-Kokkos" + uses: ./.github/workflows/tests_gpu_kokkos.yml + builddeps: + needs: ["test_lightning_kokkos"] runs-on: - self-hosted - ubuntu-22.04 diff --git a/.github/workflows/tests_gpu.yml b/.github/workflows/tests_gpu_kokkos.yml similarity index 99% rename from .github/workflows/tests_gpu.yml rename to .github/workflows/tests_gpu_kokkos.yml index bbf240d91c..15f097fb73 100644 --- a/.github/workflows/tests_gpu.yml +++ b/.github/workflows/tests_gpu_kokkos.yml @@ -1,9 +1,6 @@ name: Testing (GPU) on: - pull_request: - push: - branches: - - master + workflow_call: env: CI_CUDA_ARCH: 86 From a72e2f785d24421b07ea58da579676d66a9e1a11 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 14:00:57 +0000 Subject: [PATCH 29/43] Remove GPU deadlock. --- .github/workflows/tests_gpu_cu11.yml | 12 ++++-------- .github/workflows/tests_gpu_kokkos.yml | 7 +++++-- 2 files changed, 9 insertions(+), 10 deletions(-) diff --git a/.github/workflows/tests_gpu_cu11.yml b/.github/workflows/tests_gpu_cu11.yml index 5fb82ca491..afc27d2810 100644 --- a/.github/workflows/tests_gpu_cu11.yml +++ b/.github/workflows/tests_gpu_cu11.yml @@ -1,9 +1,9 @@ name: Testing::Linux::x86_64::LGPU on: - pull_request: - push: - branches: - - master + workflow_run: + workflows: ["Testing::LKokkos::CUDA"] + types: + - completed env: CI_CUDA_ARCH: 86 @@ -18,10 +18,6 @@ concurrency: jobs: - # Workflow call to avoid toe stepping - test_lightning_kokkos: - name: "Test Lightning-Kokkos" - uses: ./.github/workflows/tests_gpu_kokkos.yml builddeps: needs: ["test_lightning_kokkos"] diff --git a/.github/workflows/tests_gpu_kokkos.yml b/.github/workflows/tests_gpu_kokkos.yml index 15f097fb73..bc4591b29a 100644 --- a/.github/workflows/tests_gpu_kokkos.yml +++ b/.github/workflows/tests_gpu_kokkos.yml @@ -1,6 +1,9 @@ -name: Testing (GPU) +name: Testing::LKokkos::CUDA on: - workflow_call: + pull_request: + push: + branches: + - master env: CI_CUDA_ARCH: 86 From 58f742b60df0a35e341db890c8d9e6c3b2a16698 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 15:24:00 +0000 Subject: [PATCH 30/43] Bug fix Python MPI. --- .../core/src/simulators/lightning_gpu/CMakeLists.txt | 4 ++++ .../src/simulators/lightning_gpu/bindings/LGPUBindingsMPI.hpp | 2 +- 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/CMakeLists.txt b/pennylane_lightning/core/src/simulators/lightning_gpu/CMakeLists.txt index adec26a5af..5d5a336f98 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/CMakeLists.txt +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/CMakeLists.txt @@ -31,6 +31,10 @@ add_library(${PL_BACKEND} STATIC ${LGPU_FILES}) target_compile_options(lightning_compile_options INTERFACE "-D_ENABLE_PLGPU=1") +if(ENABLE_MPI) + target_compile_options(lightning_compile_options INTERFACE "-D_ENABLE_PLGPU_MPI=1") +endif() + ########################## ## Enforce C++ Standard ## ########################## diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindingsMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindingsMPI.hpp index 25bee731a4..1ca4670fe7 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindingsMPI.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/bindings/LGPUBindingsMPI.hpp @@ -154,7 +154,7 @@ void registerBackendClassSpecificBindingsMPI(PyClass &pyclass) { "Get the GPU index for the statevector data.") .def("numQubits", &StateVectorT::getNumQubits) .def("dataLength", &StateVectorT::getLength) - .def("resetGPU", &StateVectorT::initSV_MPI) + .def("resetGPU", &StateVectorT::initSV) .def( "apply", [](StateVectorT &sv, const std::string &str, From 9535f48592e76e987ec48740d43178409297b57d Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 15:45:09 +0000 Subject: [PATCH 31/43] Upload both outputs. --- .github/workflows/tests_linux_x86_mpi.yml | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/.github/workflows/tests_linux_x86_mpi.yml b/.github/workflows/tests_linux_x86_mpi.yml index e43e9ca241..03d88cba06 100644 --- a/.github/workflows/tests_linux_x86_mpi.yml +++ b/.github/workflows/tests_linux_x86_mpi.yml @@ -246,15 +246,16 @@ jobs: run: | source /etc/profile.d/modules.sh && module use /opt/modules/ && module load ${{ matrix.mpilib }} PL_DEVICE=lightning.gpu python -m pytest ./tests/ $COVERAGE_FLAGS + mv coverage.xml coverage-${{ github.job }}-lightning_gpu_${{ matrix.mpilib }}-0.xml PL_DEVICE=lightning.gpu /opt/mpi/${{ matrix.mpilib }}/bin/mpirun -np 2 python -m pytest ./mpitests $COVERAGE_FLAGS + mv coverage.xml coverage-${{ github.job }}-lightning_gpu_${{ matrix.mpilib }}-1.xml # PL_DEVICE=lightning.gpu /opt/mpi/${{ matrix.mpilib }}/bin/mpirun --oversubscribe -n 4 pytest -s -x mpitests/test_device.py -k test_create_device $COVERAGE_FLAGS - mv coverage.xml coverage-${{ github.job }}-lightning_gpu_${{ matrix.mpilib }}.xml - name: Upload code coverage results uses: actions/upload-artifact@v3 with: name: ubuntu-codecov-results-python - path: coverage-${{ github.job }}-lightning_gpu_${{ matrix.mpilib }}.xml + path: coverage-${{ github.job }}-lightning_gpu_${{ matrix.mpilib }}-*.xml if-no-files-found: error - name: Cleanup From 6648230e086e52b76457dde8be87fd13e82ab3e8 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 16:01:23 +0000 Subject: [PATCH 32/43] Update gcc version in format.yml. --- .github/workflows/format.yml | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/.github/workflows/format.yml b/.github/workflows/format.yml index 6cbde05023..2cb5d689bd 100644 --- a/.github/workflows/format.yml +++ b/.github/workflows/format.yml @@ -84,7 +84,7 @@ jobs: cp -rf ${{ github.workspace}}/Kokkos_install/${{ matrix.exec_model }}/* Kokkos/ - name: Install dependencies - run: sudo apt update && sudo apt -y install clang-tidy-14 cmake g++-10 ninja-build libomp-14-dev + run: sudo apt update && sudo apt -y install clang-tidy-14 cmake gcc-11 g++-11 ninja-build libomp-14-dev env: DEBIAN_FRONTEND: noninteractive @@ -96,5 +96,6 @@ jobs: -DBUILD_TESTS=ON \ -DENABLE_WARNINGS=ON \ -DPL_BACKEND=${{ matrix.pl_backend }} \ - -DCMAKE_CXX_COMPILER="$(which g++-10)" + -DCMAKE_CXX_COMPILER="$(which g++-11)" \ + -DCMAKE_CXX_COMPILER="$(which gcc-11)" cmake --build ./Build \ No newline at end of file From 200d81de388bdea06029eef81ba019eaa7711b17 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 13:23:28 -0400 Subject: [PATCH 33/43] Update .github/CHANGELOG.md [skip ci] Co-authored-by: Amintor Dusko <87949283+AmintorDusko@users.noreply.github.com> --- .github/CHANGELOG.md | 3 +++ 1 file changed, 3 insertions(+) diff --git a/.github/CHANGELOG.md b/.github/CHANGELOG.md index 2438bebf27..2dd314c424 100644 --- a/.github/CHANGELOG.md +++ b/.github/CHANGELOG.md @@ -2,6 +2,9 @@ ### New features since last release +* Add `SparseHamiltonian` support for Lightning-GPU. + [(#526)] (https://github.com/PennyLaneAI/pennylane-lightning/pull/526) + * Add `SparseHamiltonian` support for Lightning-Kokkos. [(#527)] (https://github.com/PennyLaneAI/pennylane-lightning/pull/527) From 1a91db590a5727ec3ad761ca6b40123286509f85 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 13:23:43 -0400 Subject: [PATCH 34/43] Update .github/workflows/tests_gpu_kokkos.yml [skip ci] Co-authored-by: Amintor Dusko <87949283+AmintorDusko@users.noreply.github.com> --- .github/workflows/tests_gpu_kokkos.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/tests_gpu_kokkos.yml b/.github/workflows/tests_gpu_kokkos.yml index bc4591b29a..40b9462f03 100644 --- a/.github/workflows/tests_gpu_kokkos.yml +++ b/.github/workflows/tests_gpu_kokkos.yml @@ -1,4 +1,4 @@ -name: Testing::LKokkos::CUDA +name: Testing::LKokkos::GPU on: pull_request: push: From 32236d11e36af9ddb12046ae2e1bb08b9d74ee76 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 18:00:32 +0000 Subject: [PATCH 35/43] rename argn [skip ci] --- .github/workflows/tests_gpu_cu11.yml | 3 -- .../core/src/bindings/Bindings.hpp | 1 - .../core/src/observables/Observables.hpp | 34 +++++++++---------- 3 files changed, 17 insertions(+), 21 deletions(-) diff --git a/.github/workflows/tests_gpu_cu11.yml b/.github/workflows/tests_gpu_cu11.yml index afc27d2810..ff42368bf6 100644 --- a/.github/workflows/tests_gpu_cu11.yml +++ b/.github/workflows/tests_gpu_cu11.yml @@ -17,10 +17,7 @@ concurrency: cancel-in-progress: true jobs: - - builddeps: - needs: ["test_lightning_kokkos"] runs-on: - self-hosted - ubuntu-22.04 diff --git a/pennylane_lightning/core/src/bindings/Bindings.hpp b/pennylane_lightning/core/src/bindings/Bindings.hpp index ac2fdd8b7f..1368ee3a94 100644 --- a/pennylane_lightning/core/src/bindings/Bindings.hpp +++ b/pennylane_lightning/core/src/bindings/Bindings.hpp @@ -628,7 +628,6 @@ template void lightningClassBindings(py::module_ &m) { /* Observables submodule */ py::module_ obs_submodule = m.def_submodule("observables", "Submodule for observables classes."); - // registerBackendAgnosticObservables(obs_submodule); registerBackendAgnosticObservables(obs_submodule); registerBackendSpecificObservables(obs_submodule); diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index b1c5665f1c..be0eb63273 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -454,16 +454,16 @@ class SparseHamiltonianBase : public Observable { * @brief Create a SparseHamiltonianBase from data, indices and offsets in * CSR format. * - * @param arg1 Arguments to construct data - * @param arg2 Arguments to construct indices - * @param arg3 Arguments to construct offsets - * @param arg4 Arguments to construct wires + * @param data Arguments to construct data + * @param indices Arguments to construct indices + * @param offsets Arguments to construct offsets + * @param wires Arguments to construct wires */ template > - SparseHamiltonianBase(T1 &&arg1, T2 &&arg2, T3 &&arg3, T4 &&arg4) - : data_{std::forward(arg1)}, indices_{std::forward(arg2)}, - offsets_{std::forward(arg3)}, wires_{std::forward(arg4)} { + SparseHamiltonianBase(T1 &&data, T2 &&indices, T3 &&offsets, T4 &&wires) + : data_{std::forward(data)}, indices_{std::forward(indices)}, + offsets_{std::forward(offsets)}, wires_{std::forward(wires)} { PL_ASSERT(data_.size() == indices_.size()); } @@ -474,21 +474,21 @@ class SparseHamiltonianBase : public Observable { * This function is useful as std::make_shared does not handle * brace-enclosed initializer list correctly. * - * @param arg1 Argument to construct data - * @param arg2 Argument to construct indices - * @param arg3 Argument to construct offsets - * @param arg4 Argument to construct wires + * @param data Argument to construct data + * @param indices Argument to construct indices + * @param offsets Argument to construct offsets + * @param wires Argument to construct wires */ - static auto create(std::initializer_list arg1, - std::initializer_list arg2, - std::initializer_list arg3, - std::initializer_list arg4) + static auto create(std::initializer_list data, + std::initializer_list indices, + std::initializer_list offsets, + std::initializer_list wires) -> std::shared_ptr> { // NOLINTBEGIN(*-move-const-arg) return std::shared_ptr>( new SparseHamiltonianBase{ - std::move(arg1), std::move(arg2), std::move(arg3), - std::move(arg4)}); + std::move(data), std::move(indices), std::move(offsets), + std::move(wires)}); // NOLINTEND(*-move-const-arg) } From 999b2c2dc55dd0583e36d52dad0e540933524708 Mon Sep 17 00:00:00 2001 From: Shuli <08cnbj@gmail.com> Date: Tue, 24 Oct 2023 18:18:06 +0000 Subject: [PATCH 36/43] Remove unused lines [skip ci] --- .../observables/ObservablesGPUMPI.hpp | 76 ------------------- 1 file changed, 76 deletions(-) diff --git a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp index b6433eeaae..94f5e45739 100644 --- a/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp +++ b/pennylane_lightning/core/src/simulators/lightning_gpu/observables/ObservablesGPUMPI.hpp @@ -283,33 +283,8 @@ class SparseHamiltonianMPI final : public SparseHamiltonianBase { } using CFP_t = typename StateVectorT::CFP_t; - /* - // Distribute sparse matrix across multi-nodes/multi-gpus - size_t num_rows = size_t{1} << sv.getTotalNumQubits(); - size_t local_num_rows = size_t{1} << sv.getNumLocalQubits(); - - std::vector>> csrmatrix_blocks; - - if (mpi_manager.getRank() == 0) { - csrmatrix_blocks = splitCSRMatrix( - mpi_manager, num_rows, this->offsets_.data(), - this->indices_.data(), this->data_.data()); - } - mpi_manager.Barrier(); - - std::vector> localCSRMatVector; - for (size_t i = 0; i < mpi_manager.getSize(); i++) { - auto localCSRMat = scatterCSRMatrix( - mpi_manager, csrmatrix_blocks[i], local_num_rows, 0); - localCSRMatVector.push_back(localCSRMat); - } - - mpi_manager.Barrier(); - */ - auto device_id = sv.getDataBuffer().getDevTag().getDeviceID(); auto stream_id = sv.getDataBuffer().getDevTag().getStreamID(); - // cusparseHandle_t handle = sv.getCusparseHandle(); const size_t length_local = size_t{1} << sv.getNumLocalQubits(); @@ -327,57 +302,6 @@ class SparseHamiltonianMPI final : public SparseHamiltonianBase { d_sv_prime->getData(), device_id, stream_id, sv.getCusparseHandle()); - /* - std::unique_ptr> d_tmp = - std::make_unique>(length_local, device_id, - stream_id, true); - - for (size_t i = 0; i < mpi_manager.getSize(); i++) { - size_t color = 0; - auto &localCSRMatrix = localCSRMatVector[i]; - - if (localCSRMatrix.getValues().size() != 0) { - color = 1; - SparseMV_cuSparse( - localCSRMatrix.getCsrOffsets().data(), - static_cast(localCSRMatrix.getCsrOffsets().size()), - localCSRMatrix.getColumns().data(), - localCSRMatrix.getValues().data(), - static_cast(localCSRMatrix.getValues().size()), - sv.getData(), d_sv_prime->getData(), device_id, stream_id, handle); - } - PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); - mpi_manager.Barrier(); - - if (mpi_manager.getRank() == i) { - color = 1; - if (localCSRMatrix.getValues().size() == 0) { - d_tmp->zeroInit(); - } - } - PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); - mpi_manager.Barrier(); - - auto new_mpi_manager = - mpi_manager.split(color, mpi_manager.getRank()); - int reduce_root_rank = -1; - - if (mpi_manager.getRank() == i) { - reduce_root_rank = new_mpi_manager.getRank(); - } - - mpi_manager.template Bcast(reduce_root_rank, i); - - if (new_mpi_manager.getComm() != MPI_COMM_NULL) { - new_mpi_manager.template Reduce( - d_tmp->getData(), d_sv_prime->getData(), length_local, - reduce_root_rank, "sum"); - } - PL_CUDA_IS_SUCCESS(cudaDeviceSynchronize()); - mpi_manager.Barrier(); - } - */ - sv.CopyGpuDataToGpuIn(d_sv_prime->getData(), d_sv_prime->getLength()); mpi_manager.Barrier(); } From e3f4854524ee6699e328912d4cb3fa95d78c4544 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 19:00:04 +0000 Subject: [PATCH 37/43] Fix SparseHamiltonianBase::isEqual. [skip ci] --- .../core/src/observables/Observables.hpp | 6 ++--- .../observables/tests/Test_Observables.cpp | 25 +++++++++++++++++++ 2 files changed, 27 insertions(+), 4 deletions(-) diff --git a/pennylane_lightning/core/src/observables/Observables.hpp b/pennylane_lightning/core/src/observables/Observables.hpp index be0eb63273..183a9fcd82 100644 --- a/pennylane_lightning/core/src/observables/Observables.hpp +++ b/pennylane_lightning/core/src/observables/Observables.hpp @@ -439,15 +439,13 @@ class SparseHamiltonianBase : public Observable { std::vector wires_; private: - // LCOV_EXCL_START [[nodiscard]] bool isEqual(const Observable &other) const override { const auto &other_cast = static_cast &>(other); - return data_ != other_cast.data_ || indices_ != other_cast.indices_ || - offsets_ != other_cast.offsets_; + return data_ == other_cast.data_ && indices_ == other_cast.indices_ && + offsets_ == other_cast.offsets_ && (wires_ == other_cast.wires_); } - // LCOV_EXCL_STOP public: /** diff --git a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp index f6877e69a9..b0541bd3e1 100644 --- a/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp +++ b/pennylane_lightning/core/src/observables/tests/Test_Observables.cpp @@ -480,6 +480,31 @@ template void testSparseHamiltonianBase() { ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, {0, 1, 2}); + DYNAMIC_SECTION("SparseHamiltonianBase - isEqual - " + << StateVectorToName::name) { + auto sparseH0 = SparseHamiltonianBase::create( + {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, + {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, + {0, 1, 2}); + auto sparseH1 = SparseHamiltonianBase::create( + {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, + {7, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, + {0, 1, 2}); + auto sparseH2 = SparseHamiltonianBase::create( + {ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}, + ComplexT{1.0, 0.0}, ComplexT{1.0, 0.0}}, + {8, 6, 5, 4, 3, 2, 1, 0}, {0, 1, 2, 3, 4, 5, 6, 7, 8}, + {0, 1, 2}); + + REQUIRE(*sparseH0 == *sparseH1); + REQUIRE(*sparseH0 != *sparseH2); + } + DYNAMIC_SECTION("SparseHamiltonianBase - getWires - " << StateVectorToName::name) { REQUIRE(sparseH->getWires() == std::vector{0, 1, 2}); From c26c99b9ac393ee318b713741beacc65474003a6 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 19:06:40 +0000 Subject: [PATCH 38/43] Trigger CI From 031d3fb68c3a2412db2dc184ffd8b543da79a40b Mon Sep 17 00:00:00 2001 From: Dev version update bot Date: Tue, 24 Oct 2023 19:07:23 +0000 Subject: [PATCH 39/43] Auto update version --- 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 b31d7d35ba..d276e92d29 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.33.0-dev23" +__version__ = "0.33.0-dev24" From 4c319a9b04756db926dbb7833f9320d72d612437 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 19:08:31 +0000 Subject: [PATCH 40/43] Trigger CI From 3a53340228f86d6d27836543a81c25afe1b4235f Mon Sep 17 00:00:00 2001 From: Shuli <08cnbj@gmail.com> Date: Tue, 24 Oct 2023 19:52:26 +0000 Subject: [PATCH 41/43] resolve comments --- .github/workflows/tests_gpu_cu11.yml | 8 ++++---- pennylane_lightning/core/_serialize.py | 7 ++++--- 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/.github/workflows/tests_gpu_cu11.yml b/.github/workflows/tests_gpu_cu11.yml index ff42368bf6..cb15aca14e 100644 --- a/.github/workflows/tests_gpu_cu11.yml +++ b/.github/workflows/tests_gpu_cu11.yml @@ -1,9 +1,9 @@ name: Testing::Linux::x86_64::LGPU on: - workflow_run: - workflows: ["Testing::LKokkos::CUDA"] - types: - - completed + pull_request: + push: + branches: + - master env: CI_CUDA_ARCH: 86 diff --git a/pennylane_lightning/core/_serialize.py b/pennylane_lightning/core/_serialize.py index 3e43fca34a..12a17a0cc5 100644 --- a/pennylane_lightning/core/_serialize.py +++ b/pennylane_lightning/core/_serialize.py @@ -78,6 +78,7 @@ def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False ) from exception else: raise DeviceError(f'The device name "{device_name}" is not a valid option.') + self.statevector_c64 = lightning_ops.StateVectorC64 self.statevector_c128 = lightning_ops.StateVectorC128 self.named_obs_c64 = lightning_ops.observables.NamedObsC64 self.named_obs_c128 = lightning_ops.observables.NamedObsC128 @@ -94,6 +95,7 @@ def __init__(self, device_name, use_csingle: bool = False, use_mpi: bool = False if use_mpi: self._use_mpi = use_mpi + self.statevector_mpi_c64 = lightning_ops.StateVectorMPIC64 self.statevector_mpi_c128 = lightning_ops.StateVectorMPIC128 self.named_obs_mpi_c64 = lightning_ops.observablesMPI.NamedObsMPIC64 self.named_obs_mpi_c128 = lightning_ops.observablesMPI.NamedObsMPIC128 @@ -120,8 +122,8 @@ def rtype(self): def sv_type(self): """State vector matching ``use_csingle`` precision (and MPI if it is supported).""" if self._use_mpi: - return self.statevector_mpi_c128 - return self.statevector_c128 + return self.statevector_mpi_c64 if self.use_csingle else self.statevector_mpi_c128 + return self.statevector_c64 if self.use_csingle else self.statevector_c128 @property def named_obs(self): @@ -203,7 +205,6 @@ def _sparse_hamiltonian(self, observable, wires_map: dict): self._mpi_manager().Barrier() else: spm = observable.sparse_matrix() - spm = observable.sparse_matrix() data = np.array(spm.data).astype(self.ctype) indices = np.array(spm.indices).astype(np.int64) offsets = np.array(spm.indptr).astype(np.int64) From 3df482b0e25ce36e732fa2d6a3c98895c9ecdf22 Mon Sep 17 00:00:00 2001 From: Shuli <08cnbj@gmail.com> Date: Tue, 24 Oct 2023 19:57:09 +0000 Subject: [PATCH 42/43] rename dev_kokkos to dev --- tests/test_adjoint_jacobian.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tests/test_adjoint_jacobian.py b/tests/test_adjoint_jacobian.py index c851e1b489..5e685d8d74 100644 --- a/tests/test_adjoint_jacobian.py +++ b/tests/test_adjoint_jacobian.py @@ -982,7 +982,7 @@ def test_adjoint_SparseHamiltonian(returns): """Integration tests that compare to default.qubit for a large circuit containing parametrized operations and when using custom wire labels""" - dev_kokkos = qml.device(device_name, wires=custom_wires) + dev = qml.device(device_name, wires=custom_wires) dev_default = qml.device("default.qubit", wires=custom_wires) def circuit(params): @@ -993,13 +993,13 @@ def circuit(params): np.random.seed(1337) params = np.random.rand(n_params) - qnode_kokkos = qml.QNode(circuit, dev_kokkos, diff_method="adjoint") + qnode = qml.QNode(circuit, dev, diff_method="adjoint") qnode_default = qml.QNode(circuit, dev_default, diff_method="parameter-shift") - j_kokkos = qml.jacobian(qnode_kokkos)(params) + j_device = qml.jacobian(qnode)(params) j_default = qml.jacobian(qnode_default)(params) - assert np.allclose(j_kokkos, j_default) + assert np.allclose(j_device, j_default) @pytest.mark.parametrize( From b7eb831fdf8f0c22dee0457e99b9d231bcfc5438 Mon Sep 17 00:00:00 2001 From: Vincent Michaud-Rioux Date: Tue, 24 Oct 2023 20:16:09 +0000 Subject: [PATCH 43/43] Fix tidy. --- .github/workflows/format.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/format.yml b/.github/workflows/format.yml index 2cb5d689bd..350312d253 100644 --- a/.github/workflows/format.yml +++ b/.github/workflows/format.yml @@ -97,5 +97,5 @@ jobs: -DENABLE_WARNINGS=ON \ -DPL_BACKEND=${{ matrix.pl_backend }} \ -DCMAKE_CXX_COMPILER="$(which g++-11)" \ - -DCMAKE_CXX_COMPILER="$(which gcc-11)" + -DCMAKE_C_COMPILER="$(which gcc-11)" cmake --build ./Build \ No newline at end of file