[Libomptarget] Implement the host memory allocator with fine grained memory

[jhuber6]

Summary

This patch should enable the "Host" allocation using fine-grained
memory. As far as I understand, this is HSA managed memory that is
availible to the host, but can be accessed by the device as well.
The original patch that introduced these extensions just stipulated that
it's "non-migratable" memory, which is most likely true because it's
managed by the host but accessible by the device. This should work
sufficiently well for what we expect the "host" allocation to do.

Depends on D143771

[kevinsala]

LGTM

[JonChesterfield]

I believe coarse/fine grain and accessibility are orthogonal concepts. If the coarse grain pool happens to have the accessibility properties of "target_host", then it'll work. That may be target specific and even order of pool iteration specific.

Much safer would be to identify a memory pool with the properties that the given allocator specifies and return that. Coarse grain means synchronised on kernel boundaries iiuc.

So what is a target_host_mem allocator?

[jhuber6]

In D143775#4122420, @JonChesterfield wrote:

I believe coarse/fine grain and accessibility are orthogonal concepts. If the coarse grain pool happens to have the accessibility properties of "target_host", then it'll work. That may be target specific and even order of pool iteration specific.

Much safer would be to identify a memory pool with the properties that the given allocator specifies and return that. Coarse grain means synchronised on kernel boundaries iiuc.

So what is a target_host_mem allocator?

Yeah, I think the problem is that target_host_mem_alloc was never fully defined by @grokos. I interpreted it to mean that it's pinned memory that the device can access. I'm pretty sure coarse memory fits the bill there, but I don't know what the exact intention was with the "non-migratable" sentence.

[JonChesterfield]

In D143775#4122610, @jhuber6 wrote:

I interpreted it to mean that it's pinned memory that the device can access. I'm pretty sure coarse memory fits the bill there

I don't think that's what coarse grain means.

[jhuber6]

In D143775#4122646, @JonChesterfield wrote:

In D143775#4122610, @jhuber6 wrote:

I interpreted it to mean that it's pinned memory that the device can access. I'm pretty sure coarse memory fits the bill there

I don't think that's what coarse grain means.

I think the CUDA plugin right now just allocates pinned memory. We could do the same here just using new and the hsa_amd_memory_lock. Would be consistent at least.

[jhuber6]

Changing this to just go back to fine-grained. It worked before so it's less controversial.

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

struct AMDHostDeviceTy : public AMDGenericDeviceTy {
 
   /// Get a memory pool for fine-grained allocations.
   AMDGPUMemoryPoolTy &getFineGrainedMemoryPool() {
     assert(!FineGrainedMemoryPools.empty() && "No fine-grained mempool");
     // Retrive any memory pool.
     return *FineGrainedMemoryPools[0];
   }
 
+  AMDGPUMemoryPoolTy &getCoarseGrainedMemoryPool() {
+    assert(!CoarseGrainedMemoryPools.empty() && "No coarse-grained mempool");
+    // Retrive any memory pool.
+    return *CoarseGrainedMemoryPools[0];
+  }
+
   /// Get a memory pool for kernel args allocations.
   AMDGPUMemoryPoolTy &getArgsMemoryPool() {
     assert(!ArgsMemoryPools.empty() && "No kernelargs mempool");
     // Retrieve any memory pool.
     return *ArgsMemoryPools[0];
   }
 
   /// Getters for kernel args and host pinned memory managers.
int free(void *TgtPtr, TargetAllocTy Kind) override {
 
     AMDGPUMemoryPoolTy *MemoryPool = nullptr;
     switch (Kind) {
     case TARGET_ALLOC_DEFAULT:
     case TARGET_ALLOC_DEVICE:
       MemoryPool = CoarseGrainedMemoryPools[0];
       break;
     case TARGET_ALLOC_HOST:
+      MemoryPool = &HostDevice.getCoarseGrainedMemoryPool();
       break;
     case TARGET_ALLOC_SHARED:
       MemoryPool = &HostDevice.getFineGrainedMemoryPool();
       break;
     }
 
     if (!MemoryPool) {
       REPORT("No memory pool for the specified allocation kind\n");
void *AMDGPUDeviceTy::allocate(size_t Size, void *, TargetAllocTy Kind) {
   // Find the correct memory pool.
   AMDGPUMemoryPoolTy *MemoryPool = nullptr;
   switch (Kind) {
   case TARGET_ALLOC_DEFAULT:
   case TARGET_ALLOC_DEVICE:
     MemoryPool = CoarseGrainedMemoryPools[0];
     break;
   case TARGET_ALLOC_HOST:
+    MemoryPool = &HostDevice.getCoarseGrainedMemoryPool();
     break;
   case TARGET_ALLOC_SHARED:
     MemoryPool = &HostDevice.getFineGrainedMemoryPool();
     break;
   }
 
   if (!MemoryPool) {
     REPORT("No memory pool for the specified allocation kind\n");

openmp/libomptarget/test/api/omp_host_pinned_memory.c

 // RUN: %libomptarget-compile-run-and-check-generic
-// UNSUPPORTED: amdgcn-amd-amdhsa
 
 #include <omp.h>
 #include <stdio.h>
 
 // Allocate pinned memory on the host
 void *llvm_omp_target_alloc_host(size_t, int);
 void llvm_omp_target_free_host(void *, int);
 

openmp/libomptarget/test/api/omp_host_pinned_memory_alloc.c

 // RUN: %libomptarget-compile-run-and-check-generic
-// UNSUPPORTED: amdgcn-amd-amdhsa
 
 #include <omp.h>
 #include <stdio.h>
 
 int main() {
   const int N = 64;
 
   int *hst_ptr = omp_alloc(N * sizeof(int), llvm_omp_target_host_mem_alloc);
 
   for (int i = 0; i < N; ++i)
     hst_ptr[i] = 2;
 
 #pragma omp target teams distribute parallel for map(tofrom : hst_ptr[0 : N])
   for (int i = 0; i < N; ++i)
     hst_ptr[i] -= 1;
 
   int sum = 0;
   for (int i = 0; i < N; ++i)
     sum += hst_ptr[i];
 
-  omp_free(hst_ptr, llvm_omp_target_shared_mem_alloc);
+  omp_free(hst_ptr, llvm_omp_target_host_mem_alloc);
   // CHECK: PASS
   if (sum == N)
     printf("PASS\n");
 }