Block Orthogonalize broken since kernel ops was applied

[maddyscientist]

Report from @cjmorningstar10: A student in Juelich is trying to use quda_laph....he has found the multigrid solver giving back an error message I have never seen before:

Initializing the fabulous Multigrid preconditioner!
MG level 0 (GPU): Functor: N4quda11BlockOrtho_INS_14BlockKernelArgILj512ENS_13BlockOrthoArgILb1EsNS_11colorspinor12FieldOrderCBIfLi4ELi3ELi24EL16QudaFieldOrder_s8EssLb1ELb0EEENS4_IfLi4ELi3ELi1ELS5_8EssLb1ELb1EEELi4ELi3ELi2ELi24EEEEEEE
MG level 0 (GPU): block: 512 1 1
MG level 0 (GPU): ERROR: Shared bytes mismatch KernelOps: 1024  cu: 0
 (rank 0, host jrc0203, tune_quda.h:379 in void quda::Tunable::checkSharedBytes(const quda::TuneParam&, const Arg&) const [with Functor = quda::BlockOrtho_; Arg = quda::BlockKernelArg<512, quda::BlockOrthoArg<true, short int, quda::colorspinor::FieldOrderCB<float, 4, 3, 24, QUDA_FLOAT8_FIELD_ORDER, short int, short int, true, false>, quda::colorspinor::FieldOrderCB<float, 4, 3, 1, QUDA_FLOAT8_FIELD_ORDER, short int, short int, true, true>, 4, 3, 2, 24> >]())
MG level 0 (GPU):        last kernel called was (name=N4quda10BlockOrthoIssLi4ELi2ELi3ELi2ELi24EEE,volume=32x32x8x16,aux=GPU-offline,large_kernel_arg,vol=131072,parity=2,precision=2,order=8,Ns=4,Nc=72,nVec=24,block_size=8x4x4x4,n_block_ortho=2,mVec=4)

@jcosborn this looks like coming from your KernelOps addition. Do you know where this is coming from?

[jcosborn]

Sorry, I guess I didn't test MG on CUDA after moving the checking code around. This should be fixed in #1581. This was actually related to my question about __shared__ accounting, but I think it is handled correctly now.