Merge branch 'fix_grout_norm' into '24.04.1-fastpt'

fixed a bug regarding PTX commands

See merge request !7

apex/contrib/csrc/group_norm/group_norm_nhwc_bwd_one_pass_kernel.cuh

@@ -328,7 +328,8 @@ __global__ __launch_bounds__(THREADS_PER_BLOCK_)
 
   // Mark the completion of the threadblock.
   if( threadIdx.x == 0 ) {
-    asm volatile("red.release.gpu.global.add.s32 [%0], 1;" :: "l"(barrier));
+    int one = 1;
+    asm volatile("red.release.gpu.global.add.s32 [%0], 1;" :: "l"(barrier), "r"(one));
   }
 
   // Exit if that's not the last thread block.