Skip to content

Commit 03c0219

Browse files
sun-9545sunojsun-9545sunoj
authored
Fix #993: Support >32 Qubit simulations on AMD GPUs via 3D grid folding (#1016)
This PR provides a comprehensive solution to Issue #993, addressing the **hipErrorInvalidConfiguration** encountered when dispatching circuits beyond 31 qubits on AMD hardware. By refactoring the dispatch logic and indexing, qsim can now successfully run simulations of **32+ qubits on high-memory devices like the AMD MI300X.** The Solution Resolved Dispatch Limits: Implemented 3D grid folding in CreateGrid to bypass the hardware-specific 1D $x$-dimension limit (65,535 blocks). Large workloads are now distributed across $(x, y, z)$ dimensions, supporting the massive thread counts required for high-qubit states. 64-bit Indexing: Replaced 32-bit signed integers with uint64_t for state-vector addressing. This prevents index overflow when the state space exceeds $2^{31}$ amplitudes, which occurs at the 32-qubit boundary. Unlocked >32 Qubit Support: **Full State-Vector:** Successfully verified 34 qubits (~128GB VRAM) on a single MI300X. **Hybrid Simulation:** Introduced a **GPU-accelerated hybrid simulator (qsimh_base_cuda.cu)** to enable 32+ runs by partitioning the state space into manageable segments. Verification & Benchmarks **Environment:** AMD MI300X (192GB), ROCm 7.1.0. **Regression**: Small-scale circuits (< 30 qubits) run with 100% accuracy. **Stress Test**: Verified a 50-qubit hybrid simulation with 100% GPU utilization and sustained 750W power draw. **Correctness:** Confirmed that 64-bit block IDs are correctly calculated across multi-dimensional grids. **Modified Files** **apps/qsimh_base_cuda.cu** (New Hybrid Simulator) **lib/cuda2hip.h** (ROCm compatibility) **lib/simulator_cuda.h** (3D Dispatch) **lib/simulator_cuda_kernels.h** (64-bit Kernels) **lib/statespace_cuda.h** (Grid Folding) **lib/statespace_cuda_kernels.h** (Block ID Helpers) **lib/vectorspace_cuda.h** (Namespace Isolation)
1 parent 7c89d16 commit 03c0219

6 files changed

Lines changed: 124 additions & 90 deletions

File tree

‎lib/simulator_cuda.h‎

Lines changed: 26 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -344,13 +344,13 @@ class SimulatorCUDA final {
344344

345345
unsigned k = 5 + G;
346346
unsigned n = num_qubits > k ? num_qubits - k : 0;
347-
unsigned size = unsigned{1} << n;
347+
uint64_t size = uint64_t{1} << n;
348348
unsigned threads = 64U;
349-
unsigned blocks = std::max(1U, size / 2);
349+
uint64_t blocks = std::max(uint64_t{1}, size / 2);
350350

351351
IndicesH<G> d_i(d_ws);
352352

353-
ApplyGateH_Kernel<G><<<blocks, threads>>>(
353+
ApplyGateH_Kernel<G><<<CreateGrid(blocks), threads>>>(
354354
(fp_type*) d_ws, d_i.xss, d_i.ms, state.get());
355355
}
356356

@@ -368,13 +368,13 @@ class SimulatorCUDA final {
368368

369369
unsigned k = 5 + num_effective_qs;
370370
unsigned n = num_qubits > k ? num_qubits - k : 0;
371-
unsigned size = unsigned{1} << n;
371+
uint64_t size = uint64_t{1} << n;
372372
unsigned threads = 32;
373-
unsigned blocks = size;
373+
uint64_t blocks = size;
374374

375375
IndicesL<G> d_i(d_ws);
376376

377-
ApplyGateL_Kernel<G><<<blocks, threads>>>(
377+
ApplyGateL_Kernel<G><<<CreateGrid(blocks), threads>>>(
378378
(fp_type*) d_ws, d_i.xss, d_i.ms, d_i.qis, d_i.tis,
379379
1 << num_effective_qs, state.get());
380380
}
@@ -401,13 +401,13 @@ class SimulatorCUDA final {
401401

402402
unsigned k = 5 + G + cqs.size();
403403
unsigned n = num_qubits > k ? num_qubits - k : 0;
404-
unsigned size = unsigned{1} << n;
404+
uint64_t size = uint64_t{1} << n;
405405
unsigned threads = 64U;
406-
unsigned blocks = std::max(1U, size / 2);
406+
uint64_t blocks = std::max(uint64_t{1}, size / 2);
407407

408408
IndicesH<G> d_i(d_ws);
409409

410-
ApplyControlledGateH_Kernel<G><<<blocks, threads>>>(
410+
ApplyControlledGateH_Kernel<G><<<CreateGrid(blocks), threads>>>(
411411
(fp_type*) d_ws, d_i.xss, d_i.ms, num_aqs + 1, cvalsh, state.get());
412412
}
413413

@@ -426,13 +426,13 @@ class SimulatorCUDA final {
426426

427427
unsigned k = 5 + G + cqs.size();
428428
unsigned n = num_qubits > k ? num_qubits - k : 0;
429-
unsigned size = unsigned{1} << n;
429+
uint64_t size = uint64_t{1} << n;
430430
unsigned threads = 32;
431-
unsigned blocks = size;
431+
uint64_t blocks = size;
432432

433433
IndicesL<G> d_i(d_ws);
434434

435-
ApplyControlledGateLH_Kernel<G><<<blocks, threads>>>(
435+
ApplyControlledGateLH_Kernel<G><<<CreateGrid(blocks), threads>>>(
436436
(fp_type*) d_ws, d_i.xss, d_i.ms, d_i.qis, d_i.tis,
437437
d.num_aqs + 1, d.cvalsh, 1 << d.num_effective_qs, state.get());
438438
}
@@ -452,13 +452,13 @@ class SimulatorCUDA final {
452452

453453
unsigned k = 5 + G + cqs.size();
454454
unsigned n = num_qubits > k ? num_qubits - k : 0;
455-
unsigned size = unsigned{1} << n;
455+
uint64_t size = uint64_t{1} << n;
456456
unsigned threads = 32;
457-
unsigned blocks = size;
457+
uint64_t blocks = size;
458458

459459
IndicesLC<G> d_i(d_ws);
460460

461-
ApplyControlledGateL_Kernel<G><<<blocks, threads>>>(
461+
ApplyControlledGateL_Kernel<G><<<CreateGrid(blocks), threads>>>(
462462
(fp_type*) d_ws, d_i.xss, d_i.ms, d_i.qis, d_i.tis, d_i.cis,
463463
d.num_aqs + 1, d.cvalsh, 1 << d.num_effective_qs,
464464
1 << (5 - d.remaining_low_cqs), state.get());
@@ -479,11 +479,11 @@ class SimulatorCUDA final {
479479

480480
unsigned k = 5 + G;
481481
unsigned n = num_qubits > k ? num_qubits - k : 0;
482-
unsigned size = unsigned{1} << n;
482+
uint64_t size = uint64_t{1} << n;
483483

484484
unsigned s = std::min(n >= 14 ? n - 14 : 0, 4U);
485485
unsigned threads = 64U;
486-
unsigned blocks = std::max(1U, (size / 2) >> s);
486+
uint64_t blocks = std::max(uint64_t{1}, (size / 2) >> s);
487487
unsigned num_iterations_per_block = 1 << s;
488488

489489
constexpr unsigned m = 16;
@@ -493,7 +493,7 @@ class SimulatorCUDA final {
493493

494494
IndicesH<G> d_i(d_ws);
495495

496-
ExpectationValueH_Kernel<G><<<blocks, threads>>>(
496+
ExpectationValueH_Kernel<G><<<CreateGrid(blocks), threads>>>(
497497
(fp_type*) d_ws, d_i.xss, d_i.ms, num_iterations_per_block,
498498
state.get(), Plus<double>(), d_res1);
499499

@@ -517,11 +517,11 @@ class SimulatorCUDA final {
517517

518518
unsigned k = 5 + num_effective_qs;
519519
unsigned n = num_qubits > k ? num_qubits - k : 0;
520-
unsigned size = unsigned{1} << n;
520+
uint64_t size = uint64_t{1} << n;
521521

522522
unsigned s = std::min(n >= 13 ? n - 13 : 0, 5U);
523523
unsigned threads = 32;
524-
unsigned blocks = size >> s;
524+
uint64_t blocks = size >> s;
525525
unsigned num_iterations_per_block = 1 << s;
526526

527527
constexpr unsigned m = 16;
@@ -531,7 +531,7 @@ class SimulatorCUDA final {
531531

532532
IndicesL<G> d_i(d_ws);
533533

534-
ExpectationValueL_Kernel<G><<<blocks, threads>>>(
534+
ExpectationValueL_Kernel<G><<<CreateGrid(blocks), threads>>>(
535535
(fp_type*) d_ws, d_i.xss, d_i.ms, d_i.qis, d_i.tis,
536536
num_iterations_per_block, state.get(), Plus<double>(), d_res1);
537537

@@ -542,18 +542,18 @@ class SimulatorCUDA final {
542542

543543
template <unsigned m>
544544
std::complex<double> ExpectationValueReduceFinal(
545-
unsigned blocks, double mul,
545+
uint64_t blocks, double mul,
546546
const Complex* d_res1, Complex* d_res2) const {
547547
Complex res2[m];
548548

549549
if (blocks <= 16) {
550550
ErrorCheck(cudaMemcpy(res2, d_res1, blocks * sizeof(Complex),
551551
cudaMemcpyDeviceToHost));
552552
} else {
553-
unsigned threads2 = std::min(1024U, blocks);
554-
unsigned blocks2 = std::min(m, blocks / threads2);
553+
unsigned threads2 = std::min(uint64_t{1024}, blocks);
554+
uint64_t blocks2 = std::min(uint64_t{m}, blocks / threads2);
555555

556-
unsigned dblocks = std::max(1U, blocks / (blocks2 * threads2));
556+
unsigned dblocks = std::max(uint64_t{1}, blocks / (blocks2 * threads2));
557557
unsigned bytes = threads2 * sizeof(Complex);
558558

559559
Reduce2Kernel<Complex><<<blocks2, threads2, bytes>>>(
@@ -568,7 +568,7 @@ class SimulatorCUDA final {
568568
double re = 0;
569569
double im = 0;
570570

571-
for (unsigned i = 0; i < blocks; ++i) {
571+
for (uint64_t i = 0; i < blocks; ++i) {
572572
re += res2[i].re;
573573
im += res2[i].im;
574574
}

‎lib/simulator_cuda_kernels.h‎

Lines changed: 27 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -18,13 +18,13 @@
1818
#ifdef __NVCC__
1919
#include <cuda.h>
2020
#include <cuda_runtime.h>
21-
22-
#include "util_cuda.h"
2321
#elif __HIP__
2422
#include <hip/hip_runtime.h>
2523
#include "cuda2hip.h"
2624
#endif
2725

26+
#include "util_cuda.h"
27+
2828
namespace qsim {
2929

3030
template <unsigned G, typename fp_type, typename idx_type>
@@ -33,6 +33,8 @@ __global__ void ApplyGateH_Kernel(
3333
const idx_type* __restrict__ mss, fp_type* __restrict__ rstate) {
3434
// blockDim.x must be equal to 64.
3535

36+
uint64_t blockId = GetBlockId();
37+
3638
static_assert(G < 7, "gates acting on more than 6 qubits are not supported.");
3739

3840
constexpr unsigned gsize = 1 << G;
@@ -61,7 +63,7 @@ __global__ void ApplyGateH_Kernel(
6163

6264
__syncthreads();
6365

64-
idx_type i = (64 * idx_type{blockIdx.x} + threadIdx.x) & 0xffffffffffe0;
66+
idx_type i = (64 * idx_type{blockId} + threadIdx.x) & 0xffffffffffe0;
6567
idx_type ii = i & mss[0];
6668
for (unsigned j = 1; j <= G; ++j) {
6769
i *= 2;
@@ -115,6 +117,8 @@ __global__ void ApplyGateL_Kernel(
115117
fp_type* __restrict__ rstate) {
116118
// blockDim.x must be equal to 32.
117119

120+
uint64_t blockId = GetBlockId();
121+
118122
static_assert(G < 7, "gates acting on more than 6 qubits are not supported.");
119123

120124
constexpr unsigned gsize = 1 << G;
@@ -137,7 +141,7 @@ __global__ void ApplyGateL_Kernel(
137141
}
138142
}
139143

140-
idx_type i = 32 * idx_type{blockIdx.x};
144+
idx_type i = 32 * idx_type{blockId};
141145
idx_type ii = i & mss[0];
142146
for (unsigned j = 1; j <= G; ++j) {
143147
i *= 2;
@@ -204,6 +208,8 @@ __global__ void ApplyControlledGateH_Kernel(
204208
fp_type* __restrict__ rstate) {
205209
// blockDim.x must be equal to 64.
206210

211+
uint64_t blockId = GetBlockId();
212+
207213
static_assert(G < 7, "gates acting on more than 6 qubits are not supported.");
208214

209215
constexpr unsigned gsize = 1 << G;
@@ -232,7 +238,7 @@ __global__ void ApplyControlledGateH_Kernel(
232238

233239
__syncthreads();
234240

235-
idx_type i = (64 * idx_type{blockIdx.x} + threadIdx.x) & 0xffffffffffe0;
241+
idx_type i = (64 * idx_type{blockId} + threadIdx.x) & 0xffffffffffe0;
236242
idx_type ii = i & mss[0];
237243
for (unsigned j = 1; j < num_mss; ++j) {
238244
i *= 2;
@@ -288,6 +294,8 @@ __global__ void ApplyControlledGateLH_Kernel(
288294
unsigned esize, fp_type* __restrict__ rstate) {
289295
// blockDim.x must be equal to 32.
290296

297+
uint64_t blockId = GetBlockId();
298+
291299
static_assert(G < 7, "gates acting on more than 6 qubits are not supported.");
292300

293301
constexpr unsigned gsize = 1 << G;
@@ -300,7 +308,7 @@ __global__ void ApplyControlledGateLH_Kernel(
300308
__shared__ fp_type rs0[32][gsize + 1], is0[32][gsize + 1];
301309
__shared__ fp_type v[2 * gsize * rows];
302310

303-
idx_type i = 32 * idx_type{blockIdx.x};
311+
idx_type i = 32 * idx_type{blockId};
304312
idx_type ii = i & mss[0];
305313
for (unsigned j = 1; j < num_mss; ++j) {
306314
i *= 2;
@@ -381,6 +389,8 @@ __global__ void ApplyControlledGateL_Kernel(
381389
fp_type* __restrict__ rstate) {
382390
// blockDim.x must be equal to 32.
383391

392+
uint64_t blockId = GetBlockId();
393+
384394
static_assert(G < 7, "gates acting on more than 6 qubits are not supported.");
385395

386396
constexpr unsigned gsize = 1 << G;
@@ -393,7 +403,7 @@ __global__ void ApplyControlledGateL_Kernel(
393403
__shared__ fp_type rs0[32][gsize + 1], is0[32][gsize + 1];
394404
__shared__ fp_type v[2 * gsize * rows];
395405

396-
idx_type i = 32 * idx_type{blockIdx.x};
406+
idx_type i = 32 * idx_type{blockId};
397407
idx_type ii = i & mss[0];
398408
for (unsigned j = 1; j < num_mss; ++j) {
399409
i *= 2;
@@ -477,6 +487,8 @@ __global__ void ExpectationValueH_Kernel(
477487
const fp_type* __restrict__ rstate, Op op, cfp_type* __restrict__ result) {
478488
// blockDim.x must be equal to 64.
479489

490+
uint64_t blockId = GetBlockId();
491+
480492
static_assert(G < 7, "gates acting on more than 6 qubits are not supported.");
481493

482494
constexpr unsigned gsize = 1 << G;
@@ -508,7 +520,7 @@ __global__ void ExpectationValueH_Kernel(
508520
double im = 0;
509521

510522
for (unsigned iter = 0; iter < num_iterations_per_block; ++iter) {
511-
idx_type b = num_iterations_per_block * idx_type{blockIdx.x} + iter;
523+
idx_type b = num_iterations_per_block * idx_type{blockId} + iter;
512524

513525
idx_type i = (64 * b + threadIdx.x) & 0xffffffffffe0;
514526
idx_type ii = i & mss[0];
@@ -573,8 +585,8 @@ __global__ void ExpectationValueH_Kernel(
573585
__syncthreads();
574586

575587
if (threadIdx.x == 0) {
576-
result[blockIdx.x].re = partial2[0].re + partial2[1].re;
577-
result[blockIdx.x].im = partial2[0].im + partial2[1].im;
588+
result[blockId].re = partial2[0].re + partial2[1].re;
589+
result[blockId].im = partial2[0].im + partial2[1].im;
578590
}
579591
}
580592

@@ -587,6 +599,8 @@ __global__ void ExpectationValueL_Kernel(
587599
const fp_type* __restrict__ rstate, Op op, cfp_type* __restrict__ result) {
588600
// blockDim.x must be equal to 32.
589601

602+
uint64_t blockId = GetBlockId();
603+
590604
static_assert(G < 7, "gates acting on more than 6 qubits are not supported.");
591605

592606
constexpr unsigned gsize = 1 << G;
@@ -612,7 +626,7 @@ __global__ void ExpectationValueL_Kernel(
612626
double im = 0;
613627

614628
for (idx_type iter = 0; iter < num_iterations_per_block; ++iter) {
615-
idx_type i = 32 * (num_iterations_per_block * idx_type{blockIdx.x} + iter);
629+
idx_type i = 32 * (num_iterations_per_block * idx_type{blockId} + iter);
616630
idx_type ii = i & mss[0];
617631
for (unsigned j = 1; j <= G; ++j) {
618632
i *= 2;
@@ -673,8 +687,8 @@ __global__ void ExpectationValueL_Kernel(
673687
auto val = WarpReduce(partial[threadIdx.x], op);
674688

675689
if (threadIdx.x == 0) {
676-
result[blockIdx.x].re = val.re;
677-
result[blockIdx.x].im = val.im;
690+
result[blockId].re = val.re;
691+
result[blockId].im = val.im;
678692
}
679693
}
680694

0 commit comments

Comments
 (0)