Merge branch 'main' into backendv2

hitomitak · Aug 25, 2023 · 2f875e7 · 2f875e7
2 parents 9366882 + d8beeab
commit 2f875e7
Show file tree

Hide file tree

Showing 28 changed files with 320 additions and 136 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -11,8 +11,15 @@
 cmake_minimum_required(VERSION 3.8 FATAL_ERROR)
 file(STRINGS "qiskit_aer/VERSION.txt" VERSION_NUM)
 
+# Add CUDA to the project if needed.
+set(EXTRA_LANGUAGES "")
+if(AER_THRUST_BACKEND STREQUAL "CUDA")
+  list(APPEND EXTRA_LANGUAGES CUDA)
+endif()
+
 include(CheckLanguage)
-project(qasm_simulator VERSION ${VERSION_NUM} LANGUAGES CXX C)
+project(qasm_simulator VERSION ${VERSION_NUM} LANGUAGES CXX C ${EXTRA_LANGUAGES} )
+
 
 list(APPEND CMAKE_MODULE_PATH ${CMAKE_SOURCE_DIR}/cmake)
 list(APPEND CMAKE_PREFIX_PATH ${CMAKE_SOURCE_DIR}/cmake)
@@ -284,7 +291,7 @@ if(AER_THRUST_SUPPORTED)
 			set(CMAKE_CUDA_RUNTIME_LIBRARY None)
 
 
-			set(CUDA_NVCC_FLAGS "${AER_CUDA_ARCH_FLAGS_EXPAND} -DAER_THRUST_CUDA -I${AER_SIMULATOR_CPP_SRC_DIR} -isystem ${AER_SIMULATOR_CPP_SRC_DIR}/third-party/headers -use_fast_math --expt-extended-lambda")
+			set(CUDA_NVCC_FLAGS "${AER_CUDA_ARCH_FLAGS_EXPAND} -DAER_THRUST_GPU -DAER_THRUST_CUDA -I${AER_SIMULATOR_CPP_SRC_DIR} -isystem ${AER_SIMULATOR_CPP_SRC_DIR}/third-party/headers -use_fast_math --expt-extended-lambda")
 			set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -I${PYTHON_SITE_PATH}/cuquantum/include")
 			set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -I${PYTHON_SITE_PATH}/cutensor/include")
 			set(THRUST_DEPENDANT_LIBS "${THRUST_DEPENDANT_LIBS} -Wl,--disable-new-dtags")
@@ -311,7 +318,7 @@ if(AER_THRUST_SUPPORTED)
 			string(STRIP ${CUDA_NVCC_FLAGS} CUDA_NVCC_FLAGS)
 			string(STRIP ${THRUST_DEPENDANT_LIBS} THRUST_DEPENDANT_LIBS)
 		else()
-			set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} ${AER_CUDA_ARCH_FLAGS_EXPAND} -DAER_THRUST_CUDA -I${AER_SIMULATOR_CPP_SRC_DIR} -isystem ${AER_SIMULATOR_CPP_SRC_DIR}/third-party/headers -use_fast_math --expt-extended-lambda")
+			set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} ${AER_CUDA_ARCH_FLAGS_EXPAND} -DAER_THRUST_GPU -DAER_THRUST_CUDA -I${AER_SIMULATOR_CPP_SRC_DIR} -isystem ${AER_SIMULATOR_CPP_SRC_DIR}/third-party/headers -use_fast_math --expt-extended-lambda")
 
 			set(AER_COMPILER_DEFINITIONS ${AER_COMPILER_DEFINITIONS} THRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA)
 			set(THRUST_DEPENDANT_LIBS "-L${CUDA_TOOLKIT_ROOT_DIR}/lib64")
@@ -455,7 +462,6 @@ else() # Standalone build
 			LINK_FLAGS ${AER_LINKER_FLAGS}
 			RUNTIME_OUTPUT_DIRECTORY_DEBUG Debug
 			RUNTIME_OUTPUT_DIRECTORY_RELEASE Release)
-		enable_language(CUDA)
 	endfunction()
 
 	function(build_cpu target src_file is_exec)
@@ -519,7 +525,9 @@ else() # Standalone build
 		if(BUILD_TESTS AND NOT AER_MPI)
 			add_executable(test_libaer "${PROJECT_SOURCE_DIR}/test/runtime/runtime_sample.c")
 			target_include_directories(test_libaer PUBLIC "${PROJECT_SOURCE_DIR}/contrib/runtime/")
-			set_target_properties(test_libaer PROPERTIES RUNTIME_OUTPUT_DIRECTORY_RELEASE bin)
+      # AER_LINKER_FLAGS carry eventual OpenMP linking flags.
+			set_target_properties(test_libaer PROPERTIES RUNTIME_OUTPUT_DIRECTORY_RELEASE bin
+                                              		 LINK_FLAGS ${AER_LINKER_FLAGS})
 			target_link_libraries(test_libaer PRIVATE ${AER_LIBRARIES})
 			target_link_libraries(test_libaer PRIVATE aer)
 			add_test(NAME aer_runtime_test COMMAND bin/test_libaer)

diff --git a/qiskit_aer/backends/backend_utils.py b/qiskit_aer/backends/backend_utils.py
@@ -214,6 +214,7 @@
             "delay",
             "pauli",
             "reset",
+            "ecr",
         ]
     ),
     "extended_stabilizer": sorted(

diff --git a/qiskit_aer/backends/qasm_simulator.py b/qiskit_aer/backends/qasm_simulator.py
@@ -786,6 +786,7 @@ def _method_basis_gates(self):
                     "swap",
                     "delay",
                     "pauli",
+                    "ecr",
                 ]
             )
         if method == "extended_stabilizer":

diff --git a/qiskit_aer/quantum_info/states/aer_densitymatrix.py b/qiskit_aer/quantum_info/states/aer_densitymatrix.py
@@ -274,7 +274,7 @@ def from_label(cls, label):
 
     @staticmethod
     def from_int(i, dims):
-        size = np.product(dims)
+        size = np.prod(dims)
         state = np.zeros((size, size), dtype=complex)
         state[i, i] = 1.0
         return AerDensityMatrix(state, dims=dims)

diff --git a/qiskit_aer/quantum_info/states/aer_statevector.py b/qiskit_aer/quantum_info/states/aer_statevector.py
@@ -281,7 +281,7 @@ def from_label(cls, label):
 
     @staticmethod
     def from_int(i, dims):
-        size = np.product(dims)
+        size = np.prod(dims)
         state = np.zeros(size, dtype=complex)
         state[i] = 1.0
         return AerStatevector(state, dims=dims)
diff --git a/releasenotes/notes/correct_measure_in_state_controller-a92692fd7083c476.yaml b/releasenotes/notes/correct_measure_in_state_controller-a92692fd7083c476.yaml
@@ -0,0 +1,6 @@
+---
+fixes:
+  - |
+    ``measure`` in libaer.so wrongly read classical memory by assuming
+    opposite ordering of its indices. This fix corrects the assumed ordering.
+    This change affects only libaer.so and not python applications.
diff --git a/releasenotes/notes/enhancement_ecr_for_stabilizer_simulator-00110a1b39d35054.yaml b/releasenotes/notes/enhancement_ecr_for_stabilizer_simulator-00110a1b39d35054.yaml
@@ -0,0 +1,6 @@
+---
+features:
+  - |
+    Enable 2-qubit gate ECR for aer_stabilizer_simulator. Refer to
+    `#1883 <https://github.com/Qiskit/qiskit-aer/issues/1883` for more
+    details.
diff --git a/src/controllers/aer_controller.hpp b/src/controllers/aer_controller.hpp
@@ -285,7 +285,7 @@ void Controller::set_config(const Config &config) {
     sim_device_ = Device::ThrustCPU;
 #endif
   } else if (sim_device_name_ == "GPU") {
-#ifndef AER_THRUST_CUDA
+#ifndef AER_THRUST_GPU
     throw std::runtime_error(
         "Simulation device \"GPU\" is not supported on this system");
 #else
@@ -421,7 +421,7 @@ size_t Controller::get_system_memory_mb() {
 
 size_t Controller::get_gpu_memory_mb() {
   size_t total_physical_memory = 0;
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   int iDev, nDev, j;
   if (cudaGetDeviceCount(&nDev) != cudaSuccess) {
     cudaGetLastError();

diff --git a/src/controllers/state_controller.hpp b/src/controllers/state_controller.hpp
@@ -439,7 +439,7 @@ class AerState {
 };
 
 bool AerState::is_gpu(bool raise_error) const {
-#ifndef AER_THRUST_CUDA
+#ifndef AER_THRUST_GPU
   if (raise_error)
     throw std::runtime_error(
         "Simulation device \"GPU\" is not supported on this system");
@@ -1313,8 +1313,9 @@ uint_t AerState::apply_measure(const reg_t &qubits) {
 
   uint_t bitstring = 0;
   uint_t bit = 1;
+  uint_t mem_size = state_->creg().memory_size();
   for (const auto &qubit : qubits) {
-    if (state_->creg().creg_memory()[qubit] == '1')
+    if (state_->creg().creg_memory()[mem_size - qubit - 1] == '1')
       bitstring |= bit;
     bit <<= 1;
   }

diff --git a/src/misc/gpu_static_properties.hpp b/src/misc/gpu_static_properties.hpp
@@ -0,0 +1,24 @@
+/**
+ * This code is part of Qiskit.
+ *
+ * (C) Copyright AMD 2023.
+ *
+ * This code is licensed under the Apache License, Version 2.0. You may
+ * obtain a copy of this license in the LICENSE.txt file in the root directory
+ * of this source tree or at http://www.apache.org/licenses/LICENSE-2.0.
+ *
+ * Any modifications or derivative works of this code must retain this
+ * copyright notice, and modified files need to carry a notice indicating
+ * that they have been altered from the originals.
+ */
+#ifndef __GPU_STATIC_PRIORITIES_H__
+#define __GPU_STATIC_PRIORITIES_H__
+
+#ifdef AER_THRUST_CUDA
+// In CUDA warpSize could not be a compile-time constant so we use 32 directly.
+#define _WS 32
+// Maximum number of threads in a block.
+#define _MAX_THD 1024
+#endif // AER_THRUST_CUDA
+
+#endif //__GPU_STATIC_PRIORITIES_H__
diff --git a/src/misc/wrap_thrust.hpp b/src/misc/wrap_thrust.hpp
@@ -38,7 +38,7 @@ DISABLE_WARNING_PUSH
 #include <thrust/iterator/permutation_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
 #include <thrust/device_vector.h>
 #endif
 #include <thrust/host_vector.h>

diff --git a/src/simulators/batch_shots_executor.hpp b/src/simulators/batch_shots_executor.hpp
@@ -268,7 +268,7 @@ void BatchShotsExecutor<state_t>::run_circuit_shots(
   }
 #endif
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   if (Base::sim_device_ == Device::GPU) {
     int nDev;
     if (cudaGetDeviceCount(&nDev) != cudaSuccess) {

diff --git a/src/simulators/circuit_executor.hpp b/src/simulators/circuit_executor.hpp
@@ -294,7 +294,7 @@ void Executor<state_t>::set_config(const Config &config) {
   }
 
   // set target GPUs
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   int nDev = 0;
   if (cudaGetDeviceCount(&nDev) != cudaSuccess) {
     cudaGetLastError();
@@ -332,7 +332,7 @@ size_t Executor<state_t>::get_system_memory_mb() {
 template <class state_t>
 size_t Executor<state_t>::get_gpu_memory_mb() {
   size_t total_physical_memory = 0;
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   for (int_t iDev = 0; iDev < target_gpus_.size(); iDev++) {
     size_t freeMem, totalMem;
     cudaSetDevice(target_gpus_[iDev]);

diff --git a/src/simulators/density_matrix/densitymatrix_thrust.hpp b/src/simulators/density_matrix/densitymatrix_thrust.hpp
@@ -51,7 +51,7 @@ class DensityMatrixThrust : public UnitaryMatrixThrust<data_t> {
   //-----------------------------------------------------------------------
 
   // Return the string name of the class
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   static std::string name() { return "density_matrix_gpu"; }
 #else
   static std::string name() { return "density_matrix_thrust"; }

diff --git a/src/simulators/parallel_state_executor.hpp b/src/simulators/parallel_state_executor.hpp
@@ -796,7 +796,7 @@ void ParallelStateExecutor<state_t>::apply_ops_chunks(InputIterator first,
   }
 
   if (Base::sim_device_ == Device::GPU) {
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
     int nDev;
     if (cudaGetDeviceCount(&nDev) != cudaSuccess) {
       cudaGetLastError();

diff --git a/src/simulators/stabilizer/stabilizer_state.hpp b/src/simulators/stabilizer/stabilizer_state.hpp
@@ -39,9 +39,25 @@ const Operations::OpSet StateOpSet(
      OpType::save_state, OpType::set_stabilizer, OpType::jump, OpType::mark},
     // Gates
     {"CX", "cx", "cy", "cz", "swap", "id", "x", "y", "z", "h", "s", "sdg", "sx",
-     "sxdg", "delay", "pauli"});
-
-enum class Gates { id, x, y, z, h, s, sdg, sx, sxdg, cx, cy, cz, swap, pauli };
+     "sxdg", "delay", "pauli", "ecr"});
+
+enum class Gates {
+  id,
+  x,
+  y,
+  z,
+  h,
+  s,
+  sdg,
+  sx,
+  sxdg,
+  cx,
+  cy,
+  cz,
+  swap,
+  pauli,
+  ecr
+};
 
 //============================================================================
 // Stabilizer Table state class
@@ -181,12 +197,13 @@ const stringmap_t<Gates> State::gateset_({
     {"sx", Gates::sx},     // Sqrt X gate.
     {"sxdg", Gates::sxdg}, // Inverse Sqrt X gate.
     // Two-qubit gates
-    {"CX", Gates::cx},      // Controlled-X gate (CNOT)
-    {"cx", Gates::cx},      // Controlled-X gate (CNOT),
-    {"cy", Gates::cy},      // Controlled-Y gate
-    {"cz", Gates::cz},      // Controlled-Z gate
-    {"swap", Gates::swap},  // SWAP gate
-    {"pauli", Gates::pauli} // Pauli gate
+    {"CX", Gates::cx},       // Controlled-X gate (CNOT)
+    {"cx", Gates::cx},       // Controlled-X gate (CNOT),
+    {"cy", Gates::cy},       // Controlled-Y gate
+    {"cz", Gates::cz},       // Controlled-Z gate
+    {"swap", Gates::swap},   // SWAP gate
+    {"pauli", Gates::pauli}, // Pauli gate
+    {"ecr", Gates::ecr}      // ECR gate
 });
 
 //============================================================================
@@ -342,6 +359,16 @@ void State::apply_gate(const Operations::Op &op) {
   case Gates::pauli:
     apply_pauli(op.qubits, op.string_params[0]);
     break;
+  case Gates::ecr:
+    BaseState::qreg_.append_h(op.qubits[1]);
+    BaseState::qreg_.append_s(op.qubits[0]);
+    BaseState::qreg_.append_z(op.qubits[1]); // sdg(1)
+    BaseState::qreg_.append_s(op.qubits[1]); // sdg(1)
+    BaseState::qreg_.append_h(op.qubits[1]);
+    BaseState::qreg_.append_cx(op.qubits[0], op.qubits[1]);
+    BaseState::qreg_.append_x(op.qubits[0]);
+    BaseState::qreg_.append_x(op.qubits[1]);
+    break;
   default:
     // We shouldn't reach here unless there is a bug in gateset
     throw std::invalid_argument(

diff --git a/src/simulators/statevector/chunk/chunk.hpp b/src/simulators/statevector/chunk/chunk.hpp
@@ -230,7 +230,7 @@ class Chunk {
     return chunk_container_.lock()->trace(chunk_pos_, row, count);
   }
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   cudaStream_t stream(void) {
     return std::static_pointer_cast<DeviceChunkContainer<data_t>>(
                chunk_container_.lock())

diff --git a/src/simulators/statevector/chunk/chunk_container.hpp b/src/simulators/statevector/chunk/chunk_container.hpp
@@ -49,7 +49,7 @@ DISABLE_WARNING_POP
 #define QV_PROBABILITY_BUFFER_SIZE 4
 #define QV_NUM_INTERNAL_REGS 4
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
 #define AERDeviceVector thrust::device_vector
 #else
 #define AERDeviceVector thrust::host_vector
@@ -58,7 +58,7 @@ DISABLE_WARNING_POP
 
 #include "framework/utils.hpp"
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
 #include "simulators/statevector/chunk/cuda_kernels.hpp"
 #endif
 
@@ -144,7 +144,7 @@ class ChunkContainer
 
   virtual void set_device(void) const {}
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   virtual cudaStream_t stream(uint_t iChunk) const { return nullptr; }
 #endif
 
@@ -395,7 +395,7 @@ void ChunkContainer<data_t>::Execute(Function func, uint_t iChunk,
       conditional_bit_ = -1; // reset conditional
   }
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   cudaStream_t strm = stream(iChunk);
   if (strm) {
     uint_t nt, nb;
@@ -457,7 +457,7 @@ template <typename Function>
 void ChunkContainer<data_t>::ExecuteSum(double *pSum, Function func,
                                         uint_t iChunk, uint_t count) const {
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   uint_t size = count * func.size(chunk_bits_);
 
   set_device();
@@ -637,7 +637,7 @@ template <typename Function>
 void ChunkContainer<data_t>::ExecuteSum2(double *pSum, Function func,
                                          uint_t iChunk, uint_t count) const {
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   uint_t size = count * func.size(chunk_bits_);
 
   set_device();
@@ -816,7 +816,7 @@ void ChunkContainer<data_t>::apply_matrix(
   } else {
     auto qubits_sorted = qubits;
     std::sort(qubits_sorted.begin(), qubits_sorted.end());
-#ifndef AER_THRUST_CUDA
+#ifndef AER_THRUST_GPU
     if (N == 3) {
       StoreMatrix(mat, iChunk);
       Execute(MatrixMult8x8<data_t>(qubits, qubits_sorted), iChunk, gid, count);

diff --git a/src/simulators/statevector/chunk/chunk_manager.hpp b/src/simulators/statevector/chunk/chunk_manager.hpp
@@ -122,7 +122,7 @@ ChunkManager<data_t>::ChunkManager() {
   num_places_ = 1;
 #else
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
   if (cudaGetDeviceCount(&num_devices_) == cudaSuccess) {
     num_places_ = num_devices_;
   } else {
@@ -248,7 +248,7 @@ uint_t ChunkManager<data_t>::Allocate(int chunk_bits, int nqubits,
 
       num_buffers = AER_MAX_BUFFERS;
 
-#ifdef AER_THRUST_CUDA
+#ifdef AER_THRUST_GPU
       num_places_ = num_devices_;
       if (num_threads_per_group_ > 1)
         num_places_ *= num_threads_per_group_;