xref: /llvm-project/offload/test/mapping/array_section_implicit_capture.c (revision 16bb7e89a995dfe8b14a8ddbce3da79c217e884a)
1*330d8983SJohannes Doerfert // RUN: %libomptarget-compile-generic
2*330d8983SJohannes Doerfert // RUN: %libomptarget-run-generic 2>&1 \
3*330d8983SJohannes Doerfert // RUN: | %fcheck-generic
4*330d8983SJohannes Doerfert 
5*330d8983SJohannes Doerfert #include <stdio.h>
6*330d8983SJohannes Doerfert #include <stdlib.h>
7*330d8983SJohannes Doerfert 
8*330d8983SJohannes Doerfert #define N 1024
9*330d8983SJohannes Doerfert #define FROM 64
10*330d8983SJohannes Doerfert #define LENGTH 128
11*330d8983SJohannes Doerfert 
main()12*330d8983SJohannes Doerfert int main() {
13*330d8983SJohannes Doerfert   float *A = (float *)malloc(N * sizeof(float));
14*330d8983SJohannes Doerfert   float *B = (float *)malloc(N * sizeof(float));
15*330d8983SJohannes Doerfert   float *C = (float *)malloc(N * sizeof(float));
16*330d8983SJohannes Doerfert 
17*330d8983SJohannes Doerfert   for (int i = 0; i < N; i++) {
18*330d8983SJohannes Doerfert     C[i] = 0.0;
19*330d8983SJohannes Doerfert   }
20*330d8983SJohannes Doerfert 
21*330d8983SJohannes Doerfert   for (int i = 0; i < N; i++) {
22*330d8983SJohannes Doerfert     A[i] = i;
23*330d8983SJohannes Doerfert     B[i] = 2 * i;
24*330d8983SJohannes Doerfert   }
25*330d8983SJohannes Doerfert 
26*330d8983SJohannes Doerfert #pragma omp target enter data map(to : A[FROM : LENGTH], B[FROM : LENGTH])
27*330d8983SJohannes Doerfert #pragma omp target enter data map(alloc : C[FROM : LENGTH])
28*330d8983SJohannes Doerfert 
29*330d8983SJohannes Doerfert // A, B and C have been mapped starting at index FROM, but inside the kernel
30*330d8983SJohannes Doerfert // they are captured implicitly so the library must look them up using their
31*330d8983SJohannes Doerfert // base address.
32*330d8983SJohannes Doerfert #pragma omp target
33*330d8983SJohannes Doerfert   {
34*330d8983SJohannes Doerfert     for (int i = FROM; i < FROM + LENGTH; i++) {
35*330d8983SJohannes Doerfert       C[i] = A[i] + B[i];
36*330d8983SJohannes Doerfert     }
37*330d8983SJohannes Doerfert   }
38*330d8983SJohannes Doerfert 
39*330d8983SJohannes Doerfert #pragma omp target exit data map(from : C[FROM : LENGTH])
40*330d8983SJohannes Doerfert #pragma omp target exit data map(delete : A[FROM : LENGTH], B[FROM : LENGTH])
41*330d8983SJohannes Doerfert 
42*330d8983SJohannes Doerfert   int errors = 0;
43*330d8983SJohannes Doerfert   for (int i = FROM; i < FROM + LENGTH; i++)
44*330d8983SJohannes Doerfert     if (C[i] != A[i] + B[i])
45*330d8983SJohannes Doerfert       ++errors;
46*330d8983SJohannes Doerfert 
47*330d8983SJohannes Doerfert   // CHECK: Success
48*330d8983SJohannes Doerfert   if (errors)
49*330d8983SJohannes Doerfert     fprintf(stderr, "Failure\n");
50*330d8983SJohannes Doerfert   else
51*330d8983SJohannes Doerfert     fprintf(stderr, "Success\n");
52*330d8983SJohannes Doerfert 
53*330d8983SJohannes Doerfert   free(A);
54*330d8983SJohannes Doerfert   free(B);
55*330d8983SJohannes Doerfert   free(C);
56*330d8983SJohannes Doerfert 
57*330d8983SJohannes Doerfert   return 0;
58*330d8983SJohannes Doerfert }
59