Fix missing sync in init (llvm#870) (llvm#911)

Salinas, David · web-flow · commit ce75949d8044 · 2025-02-28T12:30:27.000-05:00
diff --git a/amd/device-libs/ockl/src/dm.cl b/amd/device-libs/ockl/src/dm.cl
@@ -950,6 +950,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
`@@ -950,6 +950,9 @@ __ockl_dm_init_v1(ulong hp, ulong sp, uint hb, uint nis)`
`950`	`950`	`}`
`951`	`951`	`}`
`952`	`952`
	`953`	`+ __builtin_amdgcn_fence(__ATOMIC_RELEASE, "agent", "global");`
	`954`	`+ __builtin_amdgcn_s_barrier();`
	`955`	`+`
`953`	`956`	`if (lid == 0) {`
`954`	`957`	`__global heap_t thp = (__global heap_t )hp;`
`955`	`958`	`AS(&thp->initial_slabs, sp, memory_order_relaxed);`