Skip to content

Conversation

@abhinavgaba
Copy link

The test was mapping a pointee array-section s.data[0:N], without mapping the base-pointer s.data itself, and then expecting s.data to be accessible on the device.

The fix is to map s.data as well.

The test was mapping a pointee array-section `s.data[0:N]`, without mapping
the base-pointer `s.data` itself, and then expecting `s.data` to be
accessible on the device.

The fix is to map `s.data` as well.
// the always, the mapper force the update of the declared mapper
// to be seen outside the "target data region"

#pragma omp declare mapper(default:myvec_t v) map(always, present, from: v, v.len, v.data[0:v.len])
Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We could add map(present, alloc: v.data) on the mapper as well, but it's not necessary for correctness.

@abhinavgaba abhinavgaba marked this pull request as ready for review December 22, 2025 22:18
@abhinavgaba
Copy link
Author

@spophale, @tob2, please take a look.

@colleeneb
Copy link
Collaborator

Out of curiosity, is this change specific to s being a struct and so for structs we will need to allocate members like s.data on the device? Or is this general to any array section mapping, like for:

   int * s_array = (int *) malloc(N*sizeof(int)) 

  #pragma omp target enter data map(to: s_array[0:s.len]) # do we now also need map(alloc: s_array)?
   ….
  // code uses s_array[0] in a target region
  ….

Do we now also need map(alloc: s_array)?

And would map(to: s.data) have worked instead of map(alloc: s.data)? I’m asking in case this comes up in some other code in the future.

Thanks!

@abhinavgaba
Copy link
Author

abhinavgaba commented Jan 6, 2026

Or is this general to any array section mapping, like for

Yes, this is applicable to all array-sections with pointer bases. For any map(s_array[...]), the pointer s_array is not supposed to be mapped, only the pointee array-section. The user has control over whether s_array itself should be mapped.

  int * s_array = (int *) malloc(N*sizeof(int));
  #pragma omp target enter data map(to: s_array[0:s.len]) // needs map(alloc: s_array)

  printf("%d\n", omp_target_is_present(&s_array, omp_get_default_device()); // Should print 0 when s_array is not mapped.

And would map(to: s.data) have worked instead of map(alloc: s.data)?

Yes, map(to: s.data) would have worked, but the to is redundant, as it would have been overwritten by the pointer-attachment that happens between s.data and s.data[...] because they are being newly mapped on the directive.

@abhinavgaba
Copy link
Author

abhinavgaba commented Jan 6, 2026

Note that for scalar pointers, there is an implicit assumed-size-array-map on target constructs, so even if the pointer is not mapped, it may still be legal to access the pointee through it. e.g.

  int * s_array = (int *) malloc(N*sizeof(int));
  #pragma omp target enter data map(to: s_array[0:s.len])

  #pragma omp target // implies map(s_array[:]) (or s_array[0:0] as per OpenMP 5.2)
  {
    s_array[0] = 111; // This is OK
  }

  #pragma omp target map(present, alloc: s_array) // But this present-check should fail
  {}

@colleeneb
Copy link
Collaborator

Got it, thank you!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants