Fix missing sync in init (llvm#870)

b-sumner · web-flow · commit cebd37351b9b · 2025-02-26T19:56:03.000-08:00
diff --git a/amd/device-libs/ockl/src/dm.cl b/amd/device-libs/ockl/src/dm.cl
@@ -942,6 +942,9 @@ __ockl_dm_init_v1(ulong hp, ulong sp, uint hb, uint nis)
         }
     }
 
+    __builtin_amdgcn_fence(__ATOMIC_RELEASE, "agent", "global");
+    __builtin_amdgcn_s_barrier();
+
     if (lid == 0) {
         __global heap_t *thp = (__global heap_t *)hp;
         AS(&thp->initial_slabs, sp, memory_order_relaxed);

Original file line number	Diff line number	Diff line change
`@@ -942,6 +942,9 @@ __ockl_dm_init_v1(ulong hp, ulong sp, uint hb, uint nis)`
`942`	`942`	`}`
`943`	`943`	`}`
`944`	`944`
	`945`	`+ __builtin_amdgcn_fence(__ATOMIC_RELEASE, "agent", "global");`
	`946`	`+ __builtin_amdgcn_s_barrier();`
	`947`	`+`
`945`	`948`	`if (lid == 0) {`
`946`	`949`	`__global heap_t thp = (__global heap_t )hp;`
`947`	`950`	`AS(&thp->initial_slabs, sp, memory_order_relaxed);`