Device pointers not implicitly mapped firstprivate

In OpenMP 5.2, §5.8.6, page 160 line 32-33, when a device pointer allocated by omp_target_alloc has implicitly been included on a target construct as a zero-length array, the pointer initialisation should not find a matching mapped list item, and so should retain its value as a firstprivate variable. I.e., the pointer should just be passed into the target region for use by the device. However, the attached code segfaults on the device without using the is_device_ptr clause. I think the runtime is following the old 5.1 behaviour, rather than the newer 5.2 behaviour.

#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
#include <omp.h>

int main(void) {

  int N = 80000000;

  int dev = omp_get_default_device();

  // Allocate
  double *A = omp_target_alloc(sizeof(double)*N, dev);
  double *B = omp_target_alloc(sizeof(double)*N, dev);
  double *C = omp_target_alloc(sizeof(double)*N, dev);
  
  // Initalize
  // segfaults without is_device_ptr(A,B,C) clause
  #pragma omp target
  #pragma omp loop
  for (int i = 0; i < N; ++i) {
    A[i] = 1.0; B[i] = 2.0; C[i] = 0.0;
  }

  // Vector add
  // segfaults without is_device_ptr(A,B,C) clause
  #pragma omp target
  #pragma omp loop
  for (int i = 0; i < N; ++i) {
    C[i] = A[i] + B[i];
  }

  double *h_C = malloc(sizeof(double) * N);

  int host = omp_get_initial_device();

  omp_target_memcpy(h_C, C, sizeof(double)*N, 0, 0, host, dev);

  omp_target_free(A, dev);
  omp_target_free(B, dev);
  omp_target_free(C, dev);

  for (int i = 0; i < N; ++i) {
    assert(h_C[i] == 3.0);
  }
  printf("Success!\n");

}

IIRC, the new behavior is not implemented yet. I assume we still zero the pointer, right? We can make an issue and see if someone can pick it up. Should not be too hard to change it from zero to keeping the value. Question is if we should have a way to retain the 5.0/1 behavior, but that could also be done easily.

@jhuber6 Interested in taking this one?

Sure, I mentioned this earlier today. To my knowledge no one has implemented this, but it would require the runtime keeping either a separate map of pointers allocated via omp_target_alloc or simply making the mapping dev_ptr -> dev_ptr in the existing map. This will require some slight overhead given the extra lookup, but users can avoid that by manually specifying is_device_ptr as before.

@tomdeakin please open an issue and I’ll try to get to it.

Did 5.2 restrict the “keep the pointer value if not mapped” only for omp_target_alloc? I thought we keep any unmapped pointer:

If a matching mapped list item is not found, the pointer retains its original value as per the32
firstprivate semantics described in Section 5.4.4.

So why do we need the map? I assume we currently just null the pointers above, and for 5.2. we would simply not do that.

What am I missing?

You’re right, I didn’t read the standard closely and though this was only applying to pointers allocated via omp_target_alloc. In this case all we need to do is check if the pointer has an associated device entry or not. I’m assuming this won’t conflict with the existing implicit mappings we do, I’ll need to look at it more closely.

1 Like

I read it as the change should be to just the base case of the implicit mappings. If all the other criteria fail, keep (5.2) rather than zero (<=5.1) the value.

Support landed in rG6e8d93e5c235