ShflSegScanStepB write to an input-only operand

[griffi-gh]

In dev_approximate_gravity_warp.cu/dev_approximate_gravity_warp_new.cu, ShflSegScanStepB binds partial as an input-only operand "r"(partial)

but the asm both reads/writes to it:

"@p add.u32 %1, r0, %1;"
"mov.u32 %0, %1;}"
: "=r"(partial) : "r"(partial), "r"(up_offset), "r"(distance));

nvcc seems to tolerate it, but this is non-conforming code/UB.

this caused it to fail compilation under SCALE which is more strict about this sort of violation:
also see: spectral-compute/scale-validation#25

[ChrisKitching]

The code could even have one fewer instructions in the asm block by fixing this problem:

static __device__ __forceinline__ int ShflSegScanStepB(
    int partial,
    uint32_t distance,
    uint32_t up_offset
) {
  asm(
      "{.reg .u32 r0;"
      ".reg .pred p;"
      "shfl.up.b32 r0, %1, %2, 0;"
      "setp.le.u32 p, %2, %3;"
      "@p add.u32 %0, r0, %1;}"
      : "=r"(partial) : "r"(partial), "r"(up_offset), "r"(distance));
   return partial;
}

The key problem is that the add was writing to the read-only binding %1 instead of the write binding %0. Fixing that makes the final mov unnecessary.