[libomptarget] [amdgpu] After a kernel dispatch packet is published, its contents must not be accessed.

[dhruvachak]

Summary

Fixes: SWDEV-275232 (With contributions from Ammar Elwazir, Laurent Morichetti, and Tony Tye)

The current code is racy. After the packet is submitted, the GPU will increment the read index. If this wraps around before the memory is read from it'll refer to a signal from an unrelated packet. Change avoids reading from the packet post-submission.

[JonChesterfield]

Change looks non-functional to me. Lifetime of signal s increased, accesses to signal and kernarg memory are made through variables instead of lifting back out of the queue. So that's fine.

What's the problem with the current code? SWDEV-275232 looks like an internal link, and I don't know of anything in the HSA spec to say the contents of the queue buffer can't be touched after the signal store + atomic write to the top combination. Is this requirement documented somewhere?

[JonChesterfield]

Some discussion occurred offline. I don't have a reference to somewhere in the HSA spec saying that this memory cannot safely be read after publishing but it follows directly from another thread potentially reusing the same memory before it is accessed here.

[JonChesterfield]

@t-tye found the reference
http://hsa.glossner.org/wp-content/uploads/2021/02/HSA-SysArch-1.2.pdf

Section 2.8.3 Queue Mechanics
When an agent assigns an AQL packet to the packet processor, the ownership of the AQL
packet is transferred. The packet processor may update the AQL packet at any time after
packet submission, and the submitting agent should not rely on the content of the AQL packet
after submission.

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

// Run on the device.
     packet->private_segment_size = KernelInfoEntry.private_segment_size;
     packet->group_segment_size = KernelInfoEntry.group_segment_size;
     packet->kernel_object = KernelInfoEntry.kernel_object;
     packet->kernarg_address = 0;     // use the block allocator
     packet->reserved2 = 0;           // impl writes id_ here
     packet->completion_signal = {0}; // may want a pool of signals
 
     KernelArgPool *ArgPool = nullptr;
+    void *kernarg = nullptr;
     {
       auto it = KernelArgPoolMap.find(std::string(KernelInfo->Name));
       if (it != KernelArgPoolMap.end()) {
         ArgPool = (it->second).get();
       }
     }
     if (!ArgPool) {
       DP("Warning: No ArgPool for %s on device %d\n", KernelInfo->Name,
          device_id);
     }
     {
-      void *kernarg = nullptr;
       if (ArgPool) {
         assert(ArgPool->kernarg_segment_size == (arg_num * sizeof(void *)));
         kernarg = ArgPool->allocate(arg_num);
       }
       if (!kernarg) {
         DP("Allocate kernarg failed\n");
         return OFFLOAD_FAIL;
       }
if (!ArgPool) {
              "error\n");
           return OFFLOAD_FAIL;
         }
       }
 
       packet->kernarg_address = kernarg;
     }
 
-    {
-      hsa_signal_t s = DeviceInfo.FreeSignalPool.pop();
-      if (s.handle == 0) {
-        DP("Failed to get signal instance\n");
-        return OFFLOAD_FAIL;
-      }
-      packet->completion_signal = s;
-      hsa_signal_store_relaxed(packet->completion_signal, 1);
-    }
 
+    // Publish the packet indicating it is ready to be processed
     core::packet_store_release(reinterpret_cast<uint32_t *>(packet),
                                core::create_header(), packet->setup);
 
+    // Since the packet is already published, its contents must not be
+    // accessed any more
     hsa_signal_store_relaxed(queue->doorbell_signal, packet_id);
 
-    while (hsa_signal_wait_scacquire(packet->completion_signal,
-                                     HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX,
+    while (hsa_signal_wait_scacquire(s, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX,
                                      HSA_WAIT_STATE_BLOCKED) != 0)
       ;
 
     assert(ArgPool);
-    ArgPool->deallocate(packet->kernarg_address);
-    DeviceInfo.FreeSignalPool.push(packet->completion_signal);
+    ArgPool->deallocate(kernarg);
+    DeviceInfo.FreeSignalPool.push(s);
   }
 
   DP("Kernel completed\n");
   return OFFLOAD_SUCCESS;
 }
 
 int32_t __tgt_rtl_run_target_region(int32_t device_id, void *tgt_entry_ptr,
                                     void **tgt_args, ptrdiff_t *tgt_offsets,