Fix cuda::atomic_thread_fence example

docs/extended_api/synchronization_primitives/atomic/atomic_thread_fence.md

@@ -20,17 +20,39 @@ It has the same semantics as [`cuda::std::atomic_thread_fence`].
 
 ## Example
 
+The following code is an example of the [Message Passing] pattern:
+
 ```cuda
+#include <cstdio>
 #include <cuda/atomic>
-
-__global__ void example_kernel(int* data) {
-  *data = 42;
-  cuda::atomic_thread_fence(cuda::std::memory_order_release,
-                            cuda::thread_scope_device);
+#include <cooperative_groups.h>
+
+namespace cg = cooperative_groups;
+
+__global__ void example_kernel(int* data, cuda::std::atomic_flag* flag) {
+  assert(cg::grid_group::size() == 2);
+  assert(cg::thread_block::size() == 1);
+
+  if (blockIdx.x == 0) {
+    *data = 42;
+    cuda::atomic_thread_fence(cuda::memory_order_release,
+                              cuda::thread_scope_device);
+    flag->test_and_set(cuda::std::memory_order_relaxed);
+    flag->notify_one();
+  }
+  else {
+    // an atomic operation is required to set up the synchronization
+    flag->wait(false, cuda::std::memory_order_relaxed);
+    cuda::atomic_thread_fence(cuda::memory_order_acquire,
+                              cuda::thread_scope_device);
+    std::printf("%d\n", *data); // Prints 42
+  }
 }
 ```
 
 
-[See it on Godbolt](https://godbolt.org/z/nfcoTW1Kz){: .btn }
+[See it on Godbolt](https://godbolt.org/z/aG37o5qxx){: .btn }
 
 [`cuda::std::atomic_thread_fence`]: https://en.cppreference.com/w/cpp/atomic/atomic_thread_fence
+
+[Message Passing]: ../../../extended_api/memory_model.md#example-message-passing