330d8983d2
In a nutshell, this moves our libomptarget code to populate the offload subproject. With this commit, users need to enable the new LLVM/Offload subproject as a runtime in their cmake configuration. No further changes are expected for downstream code. Tests and other components still depend on OpenMP and have also not been renamed. The results below are for a build in which OpenMP and Offload are enabled runtimes. In addition to the pure `git mv`, we needed to adjust some CMake files. Nothing is intended to change semantics. ``` ninja check-offload ``` Works with the X86 and AMDGPU offload tests ``` ninja check-openmp ``` Still works but doesn't build offload tests anymore. ``` ls install/lib ``` Shows all expected libraries, incl. - `libomptarget.devicertl.a` - `libomptarget-nvptx-sm_90.bc` - `libomptarget.rtl.amdgpu.so` -> `libomptarget.rtl.amdgpu.so.18git` - `libomptarget.so` -> `libomptarget.so.18git` Fixes: https://github.com/llvm/llvm-project/issues/75124 --------- Co-authored-by: Saiyedul Islam <Saiyedul.Islam@amd.com>
231 lines
7.7 KiB
C
231 lines
7.7 KiB
C
// RUN: %libomptarget-compile-generic -fopenmp-extensions
|
|
// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
|
|
|
|
#include <omp.h>
|
|
#include <stdio.h>
|
|
|
|
#pragma omp begin declare target
|
|
#define MAX_NAME_SIZE 100
|
|
char N1[MAX_NAME_SIZE], N2[MAX_NAME_SIZE];
|
|
int V1, V2;
|
|
void copy_name(char *dst, char *src) {
|
|
int i;
|
|
for (i = 0; i < MAX_NAME_SIZE - 1 && src[i]; ++i)
|
|
dst[i] = src[i];
|
|
dst[i] = 0;
|
|
}
|
|
#pragma omp end declare target
|
|
|
|
#define CHECK_PRESENCE(Var1, Var2, Var3) \
|
|
printf(" presence of %s, %s, %s: %d, %d, %d\n", #Var1, #Var2, #Var3, \
|
|
omp_target_is_present(&(Var1), omp_get_default_device()), \
|
|
omp_target_is_present(&(Var2), omp_get_default_device()), \
|
|
omp_target_is_present(&(Var3), omp_get_default_device()))
|
|
|
|
#define CHECK_VALUES_HELPER(N1, N2, Var1, Var2) \
|
|
printf(" values of %s, %s: %d, %d\n", N1, N2, (Var1), (Var2))
|
|
|
|
#define CHECK_VALUES_DELAYED(Var1, Var2) \
|
|
copy_name(N1, #Var1); \
|
|
copy_name(N2, #Var2); \
|
|
V1 = (Var1); \
|
|
V2 = (Var2);
|
|
|
|
#define CHECK_DELAYED_VALUES() \
|
|
_Pragma("omp target update from(N1, N2, V1, V2)") \
|
|
CHECK_VALUES_HELPER(N1, N2, V1, V2)
|
|
|
|
#define CHECK_VALUES(Var1, Var2) \
|
|
CHECK_VALUES_HELPER(#Var1, #Var2, (Var1), (Var2))
|
|
|
|
int main() {
|
|
struct S {
|
|
int i;
|
|
int j;
|
|
} s;
|
|
// CHECK: presence of s, s.i, s.j: 0, 0, 0
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
|
|
// =======================================================================
|
|
// Check that ompx_hold keeps entire struct present.
|
|
|
|
// -----------------------------------------------------------------------
|
|
// CHECK-LABEL: check:{{.*}}
|
|
printf("check: ompx_hold only on first member\n");
|
|
s.i = 20;
|
|
s.j = 30;
|
|
#pragma omp target data map(tofrom : s) map(ompx_hold, tofrom : s.i) \
|
|
map(tofrom : s.j)
|
|
{
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
#pragma omp target map(tofrom : s)
|
|
{
|
|
s.i = 21;
|
|
s.j = 31;
|
|
}
|
|
#pragma omp target exit data map(delete : s, s.i)
|
|
// ompx_hold on s.i applies to all of s.
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
// CHECK-NEXT: values of s.i, s.j: 20, 30
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
}
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
|
|
// CHECK-NEXT: values of s.i, s.j: 21, 31
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
|
|
// -----------------------------------------------------------------------
|
|
// CHECK-LABEL: check:{{.*}}
|
|
printf("check: ompx_hold only on last member\n");
|
|
s.i = 20;
|
|
s.j = 30;
|
|
#pragma omp target data map(tofrom : s) map(tofrom : s.i) \
|
|
map(ompx_hold, tofrom : s.j)
|
|
{
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
#pragma omp target map(tofrom : s)
|
|
{
|
|
s.i = 21;
|
|
s.j = 31;
|
|
}
|
|
#pragma omp target exit data map(delete : s, s.i)
|
|
// ompx_hold on s.j applies to all of s.
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
// CHECK-NEXT: values of s.i, s.j: 20, 30
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
}
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
|
|
// CHECK-NEXT: values of s.i, s.j: 21, 31
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
|
|
// -----------------------------------------------------------------------
|
|
// CHECK-LABEL: check:{{.*}}
|
|
printf("check: ompx_hold only on struct\n");
|
|
s.i = 20;
|
|
s.j = 30;
|
|
#pragma omp target data map(ompx_hold, tofrom : s) map(tofrom : s.i) \
|
|
map(tofrom : s.j)
|
|
{
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
#pragma omp target map(tofrom : s)
|
|
{
|
|
s.i = 21;
|
|
s.j = 31;
|
|
}
|
|
#pragma omp target exit data map(delete : s, s.i)
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
// CHECK-NEXT: values of s.i, s.j: 20, 30
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
}
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
|
|
// CHECK-NEXT: values of s.i, s.j: 21, 31
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
|
|
// =======================================================================
|
|
// Check that transfer to/from host checks reference count correctly.
|
|
|
|
// -----------------------------------------------------------------------
|
|
// CHECK-LABEL: check:{{.*}}
|
|
printf("check: parent DynRefCount=1 is not sufficient for transfer\n");
|
|
s.i = 20;
|
|
s.j = 30;
|
|
#pragma omp target data map(ompx_hold, tofrom : s)
|
|
#pragma omp target data map(ompx_hold, tofrom : s)
|
|
{
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
#pragma omp target map(from : s.i, s.j)
|
|
{
|
|
s.i = 21;
|
|
s.j = 31;
|
|
} // No transfer here even though parent's DynRefCount=1.
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
// CHECK-NEXT: values of s.i, s.j: 20, 30
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
#pragma omp target map(to : s.i, s.j)
|
|
{ // No transfer here even though parent's DynRefCount=1.
|
|
// CHECK-NEXT: values of s.i, s.j: 21, 31
|
|
CHECK_VALUES_DELAYED(s.i, s.j);
|
|
}
|
|
CHECK_DELAYED_VALUES();
|
|
}
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
|
|
// CHECK-NEXT: values of s.i, s.j: 21, 31
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
|
|
// -----------------------------------------------------------------------
|
|
// CHECK-LABEL: check:{{.*}}
|
|
printf("check: parent HoldRefCount=1 is not sufficient for transfer\n");
|
|
s.i = 20;
|
|
s.j = 30;
|
|
#pragma omp target data map(tofrom : s)
|
|
#pragma omp target data map(tofrom : s)
|
|
{
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
#pragma omp target map(ompx_hold, from : s.i, s.j)
|
|
{
|
|
s.i = 21;
|
|
s.j = 31;
|
|
} // No transfer here even though parent's HoldRefCount=1.
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
// CHECK-NEXT: values of s.i, s.j: 20, 30
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
#pragma omp target map(ompx_hold, to : s.i, s.j)
|
|
{ // No transfer here even though parent's HoldRefCount=1.
|
|
// CHECK-NEXT: values of s.i, s.j: 21, 31
|
|
CHECK_VALUES_DELAYED(s.i, s.j);
|
|
}
|
|
CHECK_DELAYED_VALUES();
|
|
}
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
|
|
// CHECK-NEXT: values of s.i, s.j: 21, 31
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
|
|
// -----------------------------------------------------------------------
|
|
// CHECK-LABEL: check:{{.*}}
|
|
//
|
|
// At the beginning of a region, if the parent's TotalRefCount=1, then the
|
|
// transfer should happen.
|
|
//
|
|
// At the end of a region, it also must be true that the reference count being
|
|
// decremented is the reference count that is 1.
|
|
printf("check: parent TotalRefCount=1 is not sufficient for transfer\n");
|
|
s.i = 20;
|
|
s.j = 30;
|
|
#pragma omp target data map(ompx_hold, tofrom : s)
|
|
{
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
#pragma omp target map(ompx_hold, tofrom : s.i, s.j)
|
|
{
|
|
s.i = 21;
|
|
s.j = 31;
|
|
}
|
|
#pragma omp target exit data map(from : s.i, s.j)
|
|
// No transfer here even though parent's TotalRefCount=1.
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
|
|
// CHECK-NEXT: values of s.i, s.j: 20, 30
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
}
|
|
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
|
|
// CHECK-NEXT: values of s.i, s.j: 21, 31
|
|
CHECK_PRESENCE(s, s.i, s.j);
|
|
CHECK_VALUES(s.i, s.j);
|
|
|
|
return 0;
|
|
}
|