Skip to content

Commit f34dd88

Browse files
uuudownuuudown
committed
update left shift
1 parent b960bf6 commit f34dd88

13 files changed

+105
-102
lines changed

‎svsim_qsharp/svsim_cpu_omp.hpp

+7-4
Original file line numberDiff line numberDiff line change
@@ -124,6 +124,7 @@ class Gate
124124
void exe_op(Simulation* sim, ValType* sv_real, ValType* sv_imag)
125125
{
126126
(*(this->op))(this, sim, sv_real, sv_imag);
127+
#pragma omp barrier
127128
}
128129
//for dumping the gate
129130
void gateToString(std::stringstream& ss)
@@ -453,8 +454,8 @@ class Simulation
453454
assert(_n_qubits < N_QUBIT_SLOT);
454455
this->n_qubits = _n_qubits;
455456
this->n_gates = _n_gates;
456-
this->dim = ((IdxType)1UL<<(_n_qubits));
457-
this->half_dim = (IdxType)1UL<<(_n_qubits-1UL);
457+
this->dim = ((IdxType)1<<(_n_qubits));
458+
this->half_dim = (IdxType)1<<(_n_qubits-1);
458459
this->sv_size = dim*(IdxType)sizeof(ValType);
459460
}
460461
std::string circuitToString()
@@ -585,7 +586,9 @@ void simulation_kernel(Simulation* sim)
585586
#ifndef USE_AVX512 //Without AVX512 Acceleration
586587

587588
//Define MG-BSP machine operation footer
588-
#define OP_TAIL } _Pragma("omp barrier")
589+
//#define OP_TAIL } _Pragma("omp barrier")
590+
#define OP_TAIL }
591+
589592

590593
//Define MG-BSP machine operation header (Optimized version)
591594
#define OP_HEAD \
@@ -1412,7 +1415,7 @@ inline void Measure_GATE(const Gate* g, const Simulation* sim, ValType* sv_real,
14121415
//printf("\n");
14131416

14141417
ValType * m_real = sim->m_real;
1415-
IdxType mask = (1UL<<qubit);
1418+
IdxType mask = ((IdxType)1<<qubit);
14161419

14171420
if (pauli == 1)
14181421
{

‎svsim_qsharp/svsim_cpu_sin.hpp

+3-3
Original file line numberDiff line numberDiff line change
@@ -453,8 +453,8 @@ class Simulation
453453
assert(_n_qubits < N_QUBIT_SLOT);
454454
this->n_qubits = _n_qubits;
455455
this->n_gates = _n_gates;
456-
this->dim = ((IdxType)1UL<<(_n_qubits));
457-
this->half_dim = (IdxType)1UL<<(_n_qubits-1UL);
456+
this->dim = ((IdxType)1<<(_n_qubits));
457+
this->half_dim = (IdxType)1<<(_n_qubits-1);
458458
this->sv_size = dim*(IdxType)sizeof(ValType);
459459
}
460460
std::string circuitToString()
@@ -1389,7 +1389,7 @@ inline void Measure_GATE(const Gate* g, const Simulation* sim, ValType* sv_real,
13891389
//printf("\n");
13901390

13911391
ValType * m_real = sim->m_real;
1392-
IdxType mask = (1UL<<qubit);
1392+
IdxType mask = ((IdxType)1<<qubit);
13931393

13941394
if (pauli == 1)
13951395
{

‎svsim_qsharp/svsim_nvgpu_mpi.cuh

+13-13
Original file line numberDiff line numberDiff line change
@@ -811,12 +811,12 @@ __global__ void simulation_kernel(Simulation* sim)
811811
//for (IdxType i=(sim->i_gpu)*per_pe_work+tid; i<(sim->i_gpu+1)*per_pe_work;\
812812
//i+=blockDim.x*gridDim.x){ \
813813
//IdxType outer = (i >> qubit); \
814-
//IdxType inner = (i & ((1UL<<qubit)-1UL)); \
815-
//IdxType offset = (outer << (qubit+1UL)); \
814+
//IdxType inner = (i & (((IdxType)1<<qubit)-1)); \
815+
//IdxType offset = (outer << (qubit+1)); \
816816
//IdxType pos0_gid = ((offset + inner) >> (sim->lg2_m_gpu)); \
817-
//IdxType pos0 = ((offset + inner) & (sim->m_gpu-1UL)); \
818-
//IdxType pos1_gid = ((offset + inner + (1<<qubit)) >> (sim->lg2_m_gpu));\
819-
//IdxType pos1 = ((offset + inner + (1<<qubit)) & (sim->m_gpu-1UL));
817+
//IdxType pos0 = ((offset + inner) & (sim->m_gpu-1)); \
818+
//IdxType pos1_gid = ((offset + inner + ((IdxType)1<<qubit)) >> (sim->lg2_m_gpu));\
819+
//IdxType pos1 = ((offset + inner + ((IdxType)1<<qubit)) & (sim->m_gpu-1));
820820

821821
#define OP_HEAD grid_group grid = this_grid(); \
822822
const IdxType tid = blockDim.x * blockIdx.x + threadIdx.x; \
@@ -838,14 +838,14 @@ __global__ void simulation_kernel(Simulation* sim)
838838
//for (IdxType i=(sim->i_gpu)*per_pe_work+tid; i<(sim->i_gpu+1)*per_pe_work;\
839839
//i+=blockDim.x*gridDim.x){ \
840840
//IdxType outer = (i >> qubit); \
841-
//IdxType inner = (i & ((1UL<<qubit)-1UL)); \
842-
//IdxType offset = (outer << (qubit+1UL)); \
841+
//IdxType inner = (i & (((IdxType)1<<qubit)-1)); \
842+
//IdxType offset = (outer << (qubit+1)); \
843843
//IdxType pos0_src = offset + inner; \
844844
//if (((~(pos0_src&mask))&mask) != 0) continue; \
845845
//IdxType pos0_gid = ((offset + inner) >> (sim->lg2_m_gpu)); \
846-
//IdxType pos0 = ((offset + inner) & (sim->m_gpu-1UL)); \
847-
//IdxType pos1_gid = ((offset + inner + (1<<qubit)) >> (sim->lg2_m_gpu));\
848-
//IdxType pos1 = ((offset + inner + (1<<qubit)) & (sim->m_gpu-1UL));
846+
//IdxType pos0 = ((offset + inner) & (sim->m_gpu-1)); \
847+
//IdxType pos1_gid = ((offset + inner + ((IdxType)1<<qubit)) >> (sim->lg2_m_gpu));\
848+
//IdxType pos1 = ((offset + inner + ((IdxType)1<<qubit)) & (sim->m_gpu-1));
849849

850850

851851
//Define MG-BSP machine operation header with a mask for multi-controlled gates
@@ -855,8 +855,8 @@ __global__ void simulation_kernel(Simulation* sim)
855855
for (IdxType i=(sim->i_gpu)*per_pe_work+tid; i<(sim->i_gpu+1)*per_pe_work;\
856856
i+=blockDim.x*gridDim.x){ \
857857
IdxType outer = (i >> qubit); \
858-
IdxType inner = (i & ((1UL<<qubit)-1UL)); \
859-
IdxType offset = (outer << (qubit+1UL)); \
858+
IdxType inner = (i & (((IdxType)1<<qubit)-1)); \
859+
IdxType offset = (outer << (qubit+1)); \
860860
IdxType pos0_src = offset + inner; \
861861
if (((~(pos0_src&mask))&mask) != 0) continue; \
862862
IdxType pos0_gid = ((offset + inner)&(sim->n_gpus-1));\
@@ -1754,7 +1754,7 @@ __device__ __inline__ void Measure_GATE(const Gate* g, const Simulation* sim, Va
17541754
#define BARRIER if(threadIdx.x==0 && blockIdx.x==0) nvshmem_barrier_all(); grid.sync();
17551755

17561756
ValType * m_real = sim->m_real;
1757-
IdxType mask = (1UL<<qubit);
1757+
IdxType mask = ((IdxType)1<<qubit);
17581758

17591759
if (pauli == 1)
17601760
{

‎svsim_qsharp/svsim_nvgpu_sin.cuh

+2-2
Original file line numberDiff line numberDiff line change
@@ -510,8 +510,8 @@ public:
510510
assert(_n_qubits < N_QUBIT_SLOT);
511511
this->n_qubits = _n_qubits;
512512
this->n_gates = _n_gates;
513-
this->dim = ((IdxType)1UL<<(_n_qubits));
514-
this->half_dim = (IdxType)1UL<<(_n_qubits-1UL);
513+
this->dim = ((IdxType)1<<(_n_qubits));
514+
this->half_dim = (IdxType)1<<(_n_qubits-1);
515515
this->sv_size = dim*(IdxType)sizeof(ValType);
516516
}
517517
std::string circuitToString()

‎svsim_qsharp/svsim_qsharp_wrapper.cpp

+5-5
Original file line numberDiff line numberDiff line change
@@ -256,7 +256,7 @@ class SVSimSimulator final : public Microsoft::Quantum::IRuntimeDriver,
256256
IdxType mask = 0;
257257
for (long i=0; i<numControls; i++)
258258
{
259-
mask = mask | (1UL<<to_slot(controls[i]));
259+
mask = mask | ((IdxType)1<<to_slot(controls[i]));
260260
}
261261
return mask;
262262
}
@@ -336,7 +336,7 @@ class SVSimSimulator final : public Microsoft::Quantum::IRuntimeDriver,
336336
if (paulis[i] != PauliId_I)
337337
{
338338
IdxType control = to_slot(targets[i]);
339-
IdxType mask = 1UL<<control;
339+
IdxType mask = (IdxType)1<<control;
340340
sim->ControlledX(target_0, mask);
341341
}
342342
}
@@ -348,7 +348,7 @@ class SVSimSimulator final : public Microsoft::Quantum::IRuntimeDriver,
348348
if (paulis[i] != PauliId_I)
349349
{
350350
IdxType control = to_slot(targets[i]);
351-
IdxType mask = 1UL<<control;
351+
IdxType mask = (IdxType)1<<control;
352352
sim->ControlledX(target_0, mask);
353353
}
354354
}
@@ -460,7 +460,7 @@ class SVSimSimulator final : public Microsoft::Quantum::IRuntimeDriver,
460460
if (paulis[i] != PauliId_I)
461461
{
462462
IdxType control = to_slot(Qtargets[i]);
463-
IdxType mask = 1UL<<control;
463+
IdxType mask = (IdxType)1<<control;
464464
sim->ControlledX(target_0, mask);
465465
}
466466
}
@@ -473,7 +473,7 @@ class SVSimSimulator final : public Microsoft::Quantum::IRuntimeDriver,
473473
if (paulis[i] != PauliId_I)
474474
{
475475
IdxType control = to_slot(Qtargets[i]);
476-
IdxType mask = 1UL<<control;
476+
IdxType mask = (IdxType)1<<control;
477477
sim->ControlledX(target_0, mask);
478478
}
479479
}

‎svsim_qsharp/svsim_wrapper_nvgpu_omp.cu

+1-1
Original file line numberDiff line numberDiff line change
@@ -102,7 +102,7 @@ public:
102102
IdxType mask = 0;
103103
for (long i=0; i<numControls; i++)
104104
{
105-
mask = mask | (1<<to_qubit(controls[i]));
105+
mask = mask | ((IdxType)1<<to_qubit(controls[i]));
106106
}
107107
return mask;
108108
}

‎svsim_qsharp/svsim_wrapper_nvgpu_sin.cu

+5-5
Original file line numberDiff line numberDiff line change
@@ -102,7 +102,7 @@ public:
102102
IdxType mask = 0;
103103
for (long i=0; i<numControls; i++)
104104
{
105-
mask = mask | (1UL<<to_qubit(controls[i]));
105+
mask = mask | ((IdxType)1<<to_qubit(controls[i]));
106106
}
107107
return mask;
108108
}
@@ -182,7 +182,7 @@ public:
182182
if (paulis[i] != PauliId_I)
183183
{
184184
IdxType control = to_qubit(targets[i]);
185-
IdxType mask = 1UL<<control;
185+
IdxType mask = (IdxType)1<<control;
186186
sim->ControlledX(target_0, mask);
187187
}
188188
}
@@ -194,7 +194,7 @@ public:
194194
if (paulis[i] != PauliId_I)
195195
{
196196
IdxType control = to_qubit(targets[i]);
197-
IdxType mask = 1UL<<control;
197+
IdxType mask = (IdxType)1<<control;
198198
sim->ControlledX(target_0, mask);
199199
}
200200
}
@@ -306,7 +306,7 @@ public:
306306
if (paulis[i] != PauliId_I)
307307
{
308308
IdxType control = to_qubit(Qtargets[i]);
309-
IdxType mask = 1UL<<control;
309+
IdxType mask = (IdxType)1<<control;
310310
sim->ControlledX(target_0, mask);
311311
}
312312
}
@@ -319,7 +319,7 @@ public:
319319
if (paulis[i] != PauliId_I)
320320
{
321321
IdxType control = to_qubit(Qtargets[i]);
322-
IdxType mask = 1UL<<control;
322+
IdxType mask = (IdxType)1<<control;
323323
sim->ControlledX(target_0, mask);
324324
}
325325
}

‎svsim_qsharp_noise/adder_n10_nvgpu_sin.cu

+7-7
Original file line numberDiff line numberDiff line change
@@ -21,15 +21,15 @@ using namespace SVSim;
2121
//You can define circuit module functions as below.
2222
void majority(Simulation &sim, const IdxType a, const IdxType b, const IdxType c)
2323
{
24-
sim.ControlledX(b,(1UL<<c));
25-
sim.ControlledX(a,(1UL<<c));
26-
sim.ControlledX(c,((1UL<<a)|(1UL<<b)));
24+
sim.ControlledX(b,((IdxType)1<<c));
25+
sim.ControlledX(a,((IdxType)1<<c));
26+
sim.ControlledX(c,(((IdxType)1<<a)|((IdxType)1<<b)));
2727
}
2828
void unmaj(Simulation &sim, const IdxType a, const IdxType b, const IdxType c)
2929
{
30-
sim.ControlledX(c,((1UL<<a)|(1UL<<b)));
31-
sim.ControlledX(a,(1UL<<c));
32-
sim.ControlledX(b,(1UL<<a));
30+
sim.ControlledX(c,(((IdxType)1<<a)|((IdxType)1<<b)));
31+
sim.ControlledX(a,((IdxType)1<<c));
32+
sim.ControlledX(b,((IdxType)1<<a));
3333
}
3434

3535
int org_test()
@@ -57,7 +57,7 @@ int org_test()
5757
majority(sim, 1, 6, 2);
5858
majority(sim, 2, 7, 3);
5959
majority(sim, 3, 8, 4);
60-
sim.ControlledX(9, (1UL<<4));
60+
sim.ControlledX(9, ((IdxType)1<<4));
6161
unmaj(sim, 3, 8, 4);
6262
unmaj(sim, 2, 7, 3);
6363
unmaj(sim, 1, 6, 2);

‎svsim_qsharp_noise/svsim_nvgpu_mpi.cuh

+1-1
Original file line numberDiff line numberDiff line change
@@ -686,7 +686,7 @@ public:
686686

687687
IdxType* res_state = new IdxType[repetition];
688688
memset(res_state, 0, (repetition*sizeof(IdxType)));
689-
IdxType sv_num = (1<<n_qubits);
689+
IdxType sv_num = ((IdxType)1<<n_qubits);
690690
ValType* sv_diag_real = NULL;
691691
ValType* sv_diag_imag = NULL;
692692
if (i_gpu == 0) SAFE_ALOC_HOST(sv_diag_real, sv_num*sizeof(ValType));

‎svsim_qsharp_noise/svsim_nvgpu_sin.cuh

+11-11
Original file line numberDiff line numberDiff line change
@@ -31,7 +31,7 @@
3131
//#include "noise_gate_BCSZ_0.98.cuh"
3232
//#include "noise_gate_BCSZ_0.99.cuh"
3333

34-
#include "noise_gate_BCSZ_1_array.cuh"
34+
//#include "noise_gate_BCSZ_1_array.cuh"
3535

3636
#include "config.hpp"
3737

@@ -524,8 +524,8 @@ public:
524524
//For density matrix, we need double the qubits
525525
this->n_qubits = _n_qubits;
526526
this->n_gates = _n_gates;
527-
this->dim = ((IdxType)1UL<<(2*n_qubits));
528-
this->half_dim = (IdxType)1UL<<(2*n_qubits-1UL);
527+
this->dim = ((IdxType)1<<(2*n_qubits));
528+
this->half_dim = (IdxType)1<<(2*n_qubits-1);
529529
this->sv_size = dim*(IdxType)sizeof(ValType);
530530
}
531531
std::string circuitToString()
@@ -682,21 +682,21 @@ __global__ void simulation_kernel(Simulation* sim)
682682
for (IdxType i=tid; i<(sim->half_dim);\
683683
i+=blockDim.x*gridDim.x){ \
684684
IdxType outer = (i >> qubit); \
685-
IdxType inner = (i & ((1UL<<qubit)-1UL)); \
686-
IdxType offset = (outer << (qubit+1UL)); \
685+
IdxType inner = (i & (((IdxType)1<<qubit)-1)); \
686+
IdxType offset = (outer << (qubit+1)); \
687687
IdxType pos0 = offset + inner; \
688-
IdxType pos1 = pos0 + (1UL<<qubit);
688+
IdxType pos1 = pos0 + ((IdxType)1<<qubit);
689689

690690
//Define MG-BSP machine operation header with a mask for multi-controlled gates
691691
#define OP_HEAD_MASK grid_group grid = this_grid(); \
692692
const IdxType tid = blockDim.x * blockIdx.x + threadIdx.x; \
693693
for (IdxType i=tid; i<(sim->half_dim);\
694694
i+=blockDim.x*gridDim.x){ \
695695
IdxType outer = (i >> qubit); \
696-
IdxType inner = (i & ((1UL<<qubit)-1UL)); \
697-
IdxType offset = (outer << (qubit+1UL)); \
696+
IdxType inner = (i & (((IdxType)1<<qubit)-1)); \
697+
IdxType offset = (outer << (qubit+1)); \
698698
IdxType pos0 = offset + inner; \
699-
IdxType pos1 = pos0 + (1UL<<qubit); \
699+
IdxType pos1 = pos0 + ((IdxType)1<<qubit); \
700700
if (((~(pos0&mask))&mask) != 0) continue;
701701

702702
//Define MG-BSP machine operation footer
@@ -1965,7 +1965,7 @@ __device__ __inline__ void Measure_GATE(const Gate* g, const Simulation* sim, Va
19651965
const int tid = blockDim.x * blockIdx.x + threadIdx.x;
19661966

19671967
ValType * m_real = sim->m_real;
1968-
IdxType mask = (1UL<<qubit);
1968+
IdxType mask = ((IdxType)1<<qubit);
19691969

19701970
if (pauli == 1)
19711971
{
@@ -1980,7 +1980,7 @@ __device__ __inline__ void Measure_GATE(const Gate* g, const Simulation* sim, Va
19801980
H_GATE(sim, sv_real, sv_imag, sim->n_qubits+qubit);
19811981
}
19821982

1983-
for (IdxType i = tid; i<(1UL<<(sim->n_qubits)); i+=blockDim.x*gridDim.x)
1983+
for (IdxType i = tid; i<((IdxType)1<<(sim->n_qubits)); i+=blockDim.x*gridDim.x)
19841984
{
19851985
if ( (i & mask) == 0) //for all conditions with qubit=0, we set it to 0, so we sum up all prob that qubit=1
19861986
{

‎svsim_qsharp_noise/svsim_wrapper.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -124,7 +124,7 @@ class SVSimSimulator final : public Microsoft::Quantum::IRuntimeDriver,
124124
IdxType mask = 0;
125125
for (long i=0; i<numControls; i++)
126126
{
127-
mask = mask | (1UL<<to_qubit(controls[i]));
127+
mask = mask | ((IdxType)1<<to_qubit(controls[i]));
128128
}
129129
return mask;
130130
}

‎svsim_qsharp_noise/svsim_wrapper_nvgpu_sin.cu

+1-1
Original file line numberDiff line numberDiff line change
@@ -102,7 +102,7 @@ public:
102102
IdxType mask = 0;
103103
for (long i=0; i<numControls; i++)
104104
{
105-
mask = mask | (1UL<<to_qubit(controls[i]));
105+
mask = mask | ((IdxType)1<<to_qubit(controls[i]));
106106
}
107107
return mask;
108108
}

0 commit comments

Comments
 (0)