Ctrls performance

[JPRichings]

Initial pass on performance changed to remove thrust device_vector that is used to move ctrls to device that is causing a performance impact due to the thrust device_vector moving data from host to device on construction.

[JPRichings]

Todo:

• Confirm no horrible race condition is introduced by cudamemcpyToSymbol
• Set the ctrls buffer size programmatically or to some sensible limit, say 50? (at least 64KB of constant memory on device so no worries giving ourselves some extra room
• Apply fix to all kernels in QuEST

Assumptions in this code:

• Single operation on quantum register at a time which allows the assumption that ctrls can be overwritten before each subsequent gate application.

[otbrown]

Alloc size 64 qubits -- add as macro somewhere if not done already MAX_QUREG_SIZE

[TysonRayJones]

Reminder of other stuff from meeting:

• storing targets also in constant mem
• retain passing ptrs to kernels; dispatcher passes constant mem ptr

[JPRichings]

Way easier to do this instead:

Beginning with CUDA 12.1, you can now pass up to 32,764 bytes as kernel parameters on NVIDIA Volta and above

#define TOTAL_PARAMS (8000) // ints

typedef struct {
    int param[TOTAL_PARAMS];
} param_large_t;

__global__ void kernelLargeParam(__grid_constant__ const param_large_t p,...) {
    // access all parameters from p
}

int main() {
    param_large_t p_large;
    kernelLargeParam<<<GRIDDIM,BLOCKDIM>>>(p_large,...);
    cudaDeviceSynchronize();
}

Note that in both preceding examples, kernel parameters are annotated with the grid_constant qualifier to indicate they are read-only.

reference: https://developer.nvidia.com/blog/cuda-12-1-supports-large-kernel-parameters/

I think this solves our concerns about multi-qureg operations both accessing a common ctrls cache in future.

(this is also much better than the variadic kernel idea I had earlier today)

Other benefits:

1. Performance: sweet sweet 200ns more like a few micro seconds now ive checked my profiling data (could be another factor of two here) performance save possible by removing cudaMemcpyToSymbol call.
2. Removes any risk of race condition

[JPRichings]

Profiling to confirm but here is an extra factor of 2 over #729 (comment):

Total number of gates: 210
Measured probability amplitude of |0..0> state: 9.53674e-07
Calculated probability amplitude of |0..0>, C0 = 1 / 2^20: 9.53674e-07
Measuring final state: (all probabilities should be 0.5)
Qubit 0 measured in state 1 with probability 0.5
Qubit 1 measured in state 1 with probability 0.5
Qubit 2 measured in state 1 with probability 0.5
Qubit 3 measured in state 0 with probability 0.5
Qubit 4 measured in state 0 with probability 0.5
Qubit 5 measured in state 0 with probability 0.5
Qubit 6 measured in state 1 with probability 0.5
Qubit 7 measured in state 0 with probability 0.5
Qubit 8 measured in state 1 with probability 0.5
Qubit 9 measured in state 1 with probability 0.5
Qubit 10 measured in state 0 with probability 0.5
Qubit 11 measured in state 0 with probability 0.5
Qubit 12 measured in state 0 with probability 0.5
Qubit 13 measured in state 1 with probability 0.5
Qubit 14 measured in state 1 with probability 0.5
Qubit 15 measured in state 1 with probability 0.5
Qubit 16 measured in state 0 with probability 0.5
Qubit 17 measured in state 0 with probability 0.5
Qubit 18 measured in state 1 with probability 0.5
Qubit 19 measured in state 0 with probability 0.5

Final state:
|11100010110001110010>
QFT run time: 0.00141512s
Total run time: 2.36306s

[JPRichings]

profile to confirm no data movement outside of kernel launch at all:

[TysonRayJones]

Is it possible to pass multiple, distinct arguments to the one function that way? E.g. so that the target qubits of kernel_statevec_anyCtrlAnyTargDiagMatr_sub are also __grid_constant__? It'd be gross (but not the end of the world) if each kernel must manually isolate the targets array from a single ctrlsAndTargs array (although sneakily, kernel_statevec_anyCtrlOneTargDenseMatr_subA already performs such a trick).

The pattern of gpu_subroutines.cpp "secretly" passing __device__ __constant__ to unchanged kernels which still receive a ctrls ptr (and separately, a targs ptr) is still more pleasant to me, because then only gpu_subroutines.cpp needs to know about the optimisation being done, at copy-time. But a 2x speedup is nothing to sneeze at (if I interpreted that right)!

Orthogonally, do we have consternations about bumping min CUDA to 12.1 (released 2023)?

[JPRichings]

I think it is possible to have multiple arguments passed this way. All this is (and its what I tried to do first with this optimisation) capture by value in the kernel launch a c style array. The __grid_const__ is just to make sure its treated as read only on the device.

I think this is much cleaner than the data movement to const memory as instead of pointers to device constant memory we just pass the array directly. Only change to the original kernel (before all this optimisation stuff) is then we need to access a data member of a struct. There is no other change inside the kernel and we aren't accessing a global device variable out of the blue mid kernel. In gpu_subroutines.cpp if I can get util_getSorted to return directly to the struct we are using for hiding the c array then there will be no change to that code apart from a function returning to a new small struct opposed to a vector<int>.

We probably don't need to to change cuda versions unless we think we are passing over more than 4096 bytes in the kernel. I don't think we are anywhere near this but no objection to moving to CUDA 12.1.

[JPRichings]

Just to hammer home the performance improvement here. The red is the explicit cudamemcpytosymbol is the previous ctrls buffer which is now eliminated in this version.

[JPRichings]

Finally usual caveats that correctness checking needs to take place with a full run of the test suite!

[JPRichings]

Devel merged in and List64 applied to all the gpu subroutines and kernels.

Simple QFT performance test gives slightly improved results:

Total number of gates: 210
Measured probability amplitude of |0..0> state: 9.53674e-07
Calculated probability amplitude of |0..0>, C0 = 1 / 2^20: 9.53674e-07
Measuring final state: (all probabilities should be 0.5)
Qubit 0 measured in state 1 with probability 0.5
Qubit 1 measured in state 0 with probability 0.5
Qubit 2 measured in state 1 with probability 0.5
Qubit 3 measured in state 1 with probability 0.5
Qubit 4 measured in state 0 with probability 0.5
Qubit 5 measured in state 0 with probability 0.5
Qubit 6 measured in state 0 with probability 0.5
Qubit 7 measured in state 0 with probability 0.5
Qubit 8 measured in state 0 with probability 0.5
Qubit 9 measured in state 0 with probability 0.5
Qubit 10 measured in state 1 with probability 0.5
Qubit 11 measured in state 1 with probability 0.5
Qubit 12 measured in state 1 with probability 0.5
Qubit 13 measured in state 0 with probability 0.5
Qubit 14 measured in state 0 with probability 0.5
Qubit 15 measured in state 0 with probability 0.5
Qubit 16 measured in state 0 with probability 0.5
Qubit 17 measured in state 1 with probability 0.5
Qubit 18 measured in state 1 with probability 0.5
Qubit 19 measured in state 1 with probability 0.5

Final state:
|10110000001110000111>
QFT run time: 0.00120811s
Total run time: 2.37806s

Tested 20 qubits on grace-hopper.

Todo:

• Fix compilation on AMD hardware as the __grid_constant__ type modifier is not currently supported in HiP as discussed elsewhere.

[JPRichings]

I am definitely seeing a memory leak in the density evolution test but it can be removed by changing the number of repeats and number of qubits used in the test so is a bit odd really.

Have thrown nvidia compute-sanitizer over it the test:

compute-sanitizer --tool memcheck --leak-check=full ./tests/tests "density evolution"

And produced the following and many other errors:

========= Invalid __global__ read of size 8 bytes
=========     at void kernel_statevec_anyCtrlOneTargDenseMatr_subA<(int)-1>(base_qcomp *, long long, List64, long long, int, base_qcomp, base_qcomp, base_qcomp, base_qcomp)+0x650
=========     by thread (64,0,0) in block (1,0,0)
=========     Access to 0xec56a69c3248 is out of bounds
=========     and is 249272 bytes before the nearest allocation at 0xec56a6a00000 of size 1 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========         Host Frame: cudaLaunchKernel [0x3391e3] in libQuEST.so.4
=========         Host Frame: void gpu_statevec_anyCtrlOneTargDenseMatr_subA<-1>(Qureg, List64 const&, List64 const&, int, CompMatr1) [0x277993] in libQuEST.so.4
=========         Host Frame: accel_statevec_anyCtrlOneTargDenseMatr_subA(Qureg, List64 const&, List64 const&, int, CompMatr1) [0xed68b] in libQuEST.so.4
=========         Host Frame: localiser_statevec_anyCtrlOneTargDenseMatr(Qureg, List64 const&, List64 const&, int, CompMatr1, bool, bool) [0xff0c3] in libQuEST.so.4
=========         Host Frame: void localiser_statevec_anyCtrlAnyTargAnyMatr<CompMatr1>(Qureg, List64 const&, List64 const&, List64 const&, CompMatr1, bool) [0x105e6b] in libQuEST.so.4
=========         Host Frame: void validateAndApplyAnyCtrlAnyTargUnitaryMatrix<CompMatr1>(Qureg, int*, int*, int, int*, int, CompMatr1, char const*) [0xdfef7] in libQuEST.so.4
=========         Host Frame: applyMultiStateControlledCompMatr1 [0xd9b47] in libQuEST.so.4
=========         Host Frame: testDensityMatrixEvolution(Qureg, Qureg) [0x3308cf] in tests
=========         Host Frame: CATCH2_INTERNAL_TEST_0() [0x334ccb] in tests
=========         Host Frame: Catch::RunContext::invokeActiveTestCase() [0xa2f47] in libCatch2.so.3.8.0
=========         Host Frame: Catch::RunContext::runCurrentTest() [0xa3d07] in libCatch2.so.3.8.0
=========         Host Frame: Catch::RunContext::runTest(Catch::TestCaseHandle const&) [0xa417f] in libCatch2.so.3.8.0
=========         Host Frame: Catch::Session::runInternal() [0x7e007] in libCatch2.so.3.8.0
=========         Host Frame: Catch::Session::run() [0x7e4f3] in libCatch2.so.3.8.0
=========         Host Frame: main [0xeca3] in tests

Test run on grace-hopper.

Exact kernel that error is not fixed and changes based on qubit and repeat counts set in test. One concern I have is that this only got caught as gpu_sync() has a cuda error check call in it and threw the last error seen but if not for this being in a later call the errors would not have been caught.

Action: Review GPU error handling in QuEST just so we catch device side errors more consistently.

[TysonRayJones]

😆

[TysonRayJones]

Flagging an alarming bug here; ctrlsAndTarg.size() is one bigger than its previous numCtrls size. The final parameter to SET_VAR_AT_COMPILE_TIME is only ever consulted when it exceeds the NumCtrls template parameter, which has a maximum of 5 (iirc). So this bug is only triggered when passing 6 control qubits - which is impossible of our 6-qubit Qureg unit tests, since that leaves insufficient available qubits to be targets. Ergo, our tests do not uncover this bug.

Easy patch, but indicates that our tests do not test sufficiently large Qureg! I chose Qureg <= 6 deliberately to exceed the max parameter of 5, but I was thinking of target qubits which can reach 6, whereas control qubits cannot. Eep!

[TysonRayJones]

Added to my backlog to clean up the macro mischief responsible for the pitfall, described in #802.

[TysonRayJones]

This is a brilliant PR! 🎉 🎉 I note that it changes sensitive and error-prone kernel logic which (as mentioned above) is actually not covered by our existing unit tests. However, that problem already indicts the integrity of the devel branch, so I advocate for merging this PR anyway. I guess it will be a priority to get devel rigorously tested before release - I'll address the gap in the unit tests 🙏