xref: /llvm-project/offload/test/mapping/is_device_ptr.cpp (revision 330d8983d25d08580fc1642fea48b2473f47a9da)
1*330d8983SJohannes Doerfert // RUN: %libomptarget-compilexx-run-and-check-generic
2*330d8983SJohannes Doerfert 
3*330d8983SJohannes Doerfert #include <assert.h>
4*330d8983SJohannes Doerfert #include <iostream>
5*330d8983SJohannes Doerfert #include <omp.h>
6*330d8983SJohannes Doerfert 
7*330d8983SJohannes Doerfert struct view {
8*330d8983SJohannes Doerfert   const int size = 10;
9*330d8983SJohannes Doerfert   int *data_host;
10*330d8983SJohannes Doerfert   int *data_device;
fooview11*330d8983SJohannes Doerfert   void foo() {
12*330d8983SJohannes Doerfert     std::size_t bytes = size * sizeof(int);
13*330d8983SJohannes Doerfert     const int host_id = omp_get_initial_device();
14*330d8983SJohannes Doerfert     const int device_id = omp_get_default_device();
15*330d8983SJohannes Doerfert     data_host = (int *)malloc(bytes);
16*330d8983SJohannes Doerfert     data_device = (int *)omp_target_alloc(bytes, device_id);
17*330d8983SJohannes Doerfert #pragma omp target teams distribute parallel for is_device_ptr(data_device)
18*330d8983SJohannes Doerfert     for (int i = 0; i < size; ++i)
19*330d8983SJohannes Doerfert       data_device[i] = i;
20*330d8983SJohannes Doerfert     omp_target_memcpy(data_host, data_device, bytes, 0, 0, host_id, device_id);
21*330d8983SJohannes Doerfert     for (int i = 0; i < size; ++i)
22*330d8983SJohannes Doerfert       assert(data_host[i] == i);
23*330d8983SJohannes Doerfert   }
24*330d8983SJohannes Doerfert };
25*330d8983SJohannes Doerfert 
main()26*330d8983SJohannes Doerfert int main() {
27*330d8983SJohannes Doerfert   view a;
28*330d8983SJohannes Doerfert   a.foo();
29*330d8983SJohannes Doerfert   // CHECK: PASSED
30*330d8983SJohannes Doerfert   printf("PASSED\n");
31*330d8983SJohannes Doerfert }
32