[Openmp-commits] [PATCH] D115279: [OpenMP][AMDGPU] Switch host-device memory copy to asynchronous version

[Ye Luo via Phabricator via Openmp-commits]

ye-luo added a comment.

I failed to verify your first lock behaves as intended.

  #include <hsa/hsa.h>
  #include <hsa/hsa_ext_amd.h>
  #include <omp.h>
  #include <stdio.h>
  
  #define N 100293
  
  int checkLocked(void *ptr) {
    hsa_amd_pointer_info_t info;
    hsa_status_t herr;
  
    herr = hsa_amd_pointer_info(ptr, &info, NULL, NULL, NULL);
    if (herr != HSA_STATUS_SUCCESS) {
      printf("  hsa_amd_pointer_info failed\n");
      return 1;
    }
  
    if (info.type != HSA_EXT_POINTER_TYPE_LOCKED) {
      printf("  pointer is noooooooooooot locked\n");
      return 1;
    } else
      printf("  pointer is locked\n");
  
    return 0;
  }
  
  int main() {
    int n = N;
    int *a = new int[n];
    for (int i = 0; i < n; i++)
      a[i] = 0;
  
    int *a_locked = nullptr;
    hsa_status_t herr =
        hsa_amd_memory_lock(a, n * sizeof(int), nullptr, 0, (void **)&a_locked);
    if (herr != HSA_STATUS_SUCCESS) {
      printf("Locking failed\n");
      return 1;
    }
  
    checkLocked(a);
  
  #pragma omp target parallel for map(tofrom : a_locked[:n])
    for (int i = 0; i < n; i++)
      a_locked[i] = i;
  
    herr = hsa_amd_memory_unlock(a);
    if (herr != HSA_STATUS_SUCCESS) {
      printf("Unlocking failed\n");
      return 1;
    }
  
    int err = 0;
    for (int i = 0; i < n; i++)
      if (a[i] != i) {
        err++;
        printf("Err at %d, got %d expected %d\n", i, a[i], i);
        if (err > 10)
          break;
      }
  
    delete[] a;
  
    return err;
  }

I got  failure at the first check with "hsa_amd_pointer_info failed". Could you take a look?


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D115279/new/

https://reviews.llvm.org/D115279