[AUTOGENERATED] [release/2.8] [ROCm] Use opportunistic fastatomics based on heuristics (#2441)

okakarpa · tvukovic-amd · commit af7b538f71a5 · 2025-08-20T20:19:00.000+02:00
diff --git a/aten/src/ATen/native/cuda/KernelUtils.cuh b/aten/src/ATen/native/cuda/KernelUtils.cuh
@@ -281,6 +281,13 @@ __device__ __forceinline__ void opportunistic_fastAtomicAdd(
         }
     }
 
+    if (numel > 16 /*<-hueristic threshold*/ * 64 ) {
+      // well shucks, unlikely to capture same-dest atomics in a wave.
+      // fall back to direct fastAtomic...
+      fastAtomicAdd(self_ptr, index, numel, value, true);
+      return;
+    }
+
     // not coalsced, so now let try to capture lane-matches...
     // __activemask() -- finds the set of threads in the warp that are about to perform atomicAdd
     // __match_any_sync() -- returns bit mask of the threads that have same dest addr

Original file line number	Diff line number	Diff line change
`@@ -281,6 +281,13 @@ __device__ __forceinline__ void opportunistic_fastAtomicAdd(`
`281`	`281`	`}`
`282`	`282`	`}`
`283`	`283`
	`284`	`+ if (numel > 16 /<-hueristic threshold/ * 64 ) {`
	`285`	`+ // well shucks, unlikely to capture same-dest atomics in a wave.`
	`286`	`+ // fall back to direct fastAtomic...`
	`287`	`+ fastAtomicAdd(self_ptr, index, numel, value, true);`
	`288`	`+ return;`
	`289`	`+ }`
	`290`	`+`
`284`	`291`	`// not coalsced, so now let try to capture lane-matches...`
`285`	`292`	`// __activemask() -- finds the set of threads in the warp that are about to perform atomicAdd`
`286`	`293`	`// __match_any_sync() -- returns bit mask of the threads that have same dest addr`