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]

TODO: use a MAX_NUM_QUBITS = 64 or something constant in constants.hpp

[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!

[TysonRayJones]

Capture of simple struct in kernel launch tested on AMD so we might loose the benefit of __grid_constant__ to make the device data read only but we will still remove the overhead of implicity or explicit copy of data to the device outside of the kernel launch

Aw shoot! Does removing __grid_constant__ still achieve correct behaviour? Could do a regrettable but defensible macro (similar to what inliner.hpp does), like:

#if defined(__NVCC__)

    #define _NO_COPY_OPT __grid_constant__

#elif defined(__HIP__)

    #define _NO_COPY_OPT

#endif

Btw the reference you link indicates:

The __grid_constant__ attribute can be applied to a const-qualified kernel function argument and allows compiler to take the address of that argument without making a copy. The argument applies to sm_70 or newer GPUs, during compilation with CUDA-11.7(PTX 7.7) or newer, and is ignored otherwise.

Use of __grid_constant__ appears to require we bump the minimum compute capability. Right now, I'm fairly sure the minimum is like 3.0 (or whatever is imposed by Thrust)! My main QuEST development was on a Quadro P6000, so a CC of 6.1 is tattooed into my brain ehehe. Do we object to a bump?? I'm not personally too offended because the min CC appearing on the NVIDIA site is now 7.5, but it is kind of cool to try to support old GPUs - especially if one believes consumer-end GPUs are becoming harder to obtain. Could we retain support for CC < 7.0 using a similar macro trick as above?

[JPRichings]

I agree macro guarding this will probably be the solution I need to implement this change in one of the kernels and retest to guarantee the performance holds and then confirm this works for Nvidia and AMD.

I agree keeping the GPU capability support as old as possible is best to support a wider set of users so I will explore adding this into the solution so we don't have to bump the compute capability if we don't need to.

[TysonRayJones]

Heads up that SmallList has reached devel

[TysonRayJones]

Btw I incidentally made these same bitwise changes in #750, but where alle existing bitwise functions receive const pointers, as was necessary for SmallView in #754.

[TysonRayJones]

QubitList_t should be made entirely redundant by SmallList (which we can still rename)

[TysonRayJones]

Is there a blocker to this beyond further performance investigation and/or updating given my upstream changes? I'm happy to make a PR to this migrating it over to List64, and handle the aforementioned macros.

(Are you seeing a performance difference between ctrl_device_t and QubitList_t? I can switch all over to ConstList64 where possible)

[JPRichings]

No blocker technical just human time constraints I will get to this on Monday and have agreed a strategy with Oliver for merging this with everything else.

[otbrown]

Just noting here that we defer this to v4.3.1.

(Sidenote: semver says that the patch index is only for changes that fix "incorrect behaviour", but the minor version is only a MAY change if there are significant changes to the private behaviour, so where does that leave a performance patch? I suppose one could declare being much slower than it needs to be on GPU a bug, and therefore it's a bug fix...)

[TysonRayJones]

Hehe a patch index bump is inoffensive to me - they're "free" in my mental model ¯\(ツ)/¯

[otbrown]

I think a post Unitary Hack 4.3.1 makes sense then (assuming we've already got 4.3.0 out by then!)

[JPRichings]

I agree we push this to 4.3.1 so we get 4.3 out the door.

[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.