ShflSegScanStepB write to an input-only operand #54

Activity

1. ChrisKitching commented on Jul 3, 2026

   

   ChrisKitching

   on Jul 3, 2026

   More actions

   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.