Commit c52dee52 authored by Joel E. Denny's avatar Joel E. Denny
Browse files

[Clacc][OpenACC] Reconsider object map extension tests 

As discussed in D149685, extending partially mapped objects (e.g.,
structs or arrays) is invalid in OpenMP.  This patch removes many
tests that do that ("Concat2" tests), adds comments explaining tests
that do that only in some compilation modes ("Disjoint" tests), and
adds a new test where D149685 fixes the no_create presence check
("DisjointDevicePad").
parent 24bc227d

clang/test/OpenACC/directives/Tests/no-create.c

+132 −259

File changed.

Preview size limit exceeded, changes collapsed.

clang/test/OpenACC/directives/Tests/present.c

+18 −87
Original line number Diff line number Diff line
@@ -65,13 +65,11 @@ 
// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataMemberPresent            %, data      %,                                        %,                                  %,                                  %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataMemberAbsent             %, data      %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataMembersDisjoint          %, data      %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataMembersConcat2           %, data      %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataMemberFullStruct         %, data      %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataSubarrayPresent          %, data      %,                                        %,                                  %,                                  %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataSubarrayDisjoint         %, data      %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataSubarrayOverlapStart     %, data      %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataSubarrayOverlapEnd       %, data      %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataSubarrayConcat2          %, data      %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseDataSubarrayNonSubarray      %, data      %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelScalarPresent        %, parallel  %,                                        %,                                  %,                                  %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelScalarAbsent         %, parallel  %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %,                    %) && \
@@ -82,13 +80,11 @@ 
// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelMemberPresent        %, parallel  %,                                        %,                                  %,                                  %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelMemberAbsent         %, parallel  %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelMembersDisjoint      %, parallel  %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelMembersConcat2       %, parallel  %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelMemberFullStruct     %, parallel  %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelSubarrayPresent      %, parallel  %,                                        %,                                  %,                                  %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelSubarrayDisjoint     %, parallel  %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelSubarrayOverlapStart %, parallel  %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelSubarrayOverlapEnd   %, parallel  %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelSubarrayConcat2      %, parallel  %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelSubarrayNonSubarray  %, parallel  %, %if-tgt-host<|not --crash>             %, %if-tgt-host<|not>               %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|not> %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelLoopScalarPresent    %, parallel  %,                                        %,                                  %,                                  %,                    %) && \

// DEFINE:   %{check-case}( %{LINE}, %(line) %, caseParallelLoopScalarAbsent     %, parallel  %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %,                    %) && \
@@ -460,6 +456,11 @@ CASE(caseDataMemberAbsent) { 
//

//    PRT-NEXT:   ;

//    PRT-NEXT: }

//

// This case is also invalid when -fopenacc-present-omp=no-present because the

// present clause then attempts to extend a partial mapping of an object.  For

// now, the runtime doesn't complain.  If it does in the future, we'll have to

// change the expected results here.

CASE(caseDataMembersDisjoint) {

  struct S { int x; int y; } s;

  PRINT_VAR_INFO(s.x);
@@ -469,38 +470,6 @@ CASE(caseDataMembersDisjoint) { 
  USE_VAR(s.y = 1);

}



//   PRT-LABEL: {{.*}}caseDataMembersConcat2{{.*}} {

//    PRT-NEXT:   struct S {

//         PRT:   } s;

//    PRT-NEXT:   {{(PRINT_VAR_INFO|fprintf)\(.*\);}}

//    PRT-NEXT:   {{(PRINT_VAR_INFO|fprintf)\(.*\);}}

//

//  PRT-A-NEXT:   #pragma acc data copyout(s.x){{$}}

// PRT-AO-NEXT:   // #pragma omp target data map(ompx_hold,from: s.x){{$}}

//  PRT-A-NEXT:   #pragma acc data copy(s.y){{$}}

// PRT-AO-NEXT:   // #pragma omp target data map(ompx_hold,tofrom: s.y){{$}}

//  PRT-A-NEXT:   #pragma acc data present(s){{$}}

// PRT-AO-NEXT:   // #pragma omp target data map([[PRESENT_MT]]: s){{$}}

//

//  PRT-O-NEXT:   #pragma omp target data map(ompx_hold,from: s.x){{$}}

// PRT-OA-NEXT:   // #pragma acc data copyout(s.x){{$}}

//  PRT-O-NEXT:   #pragma omp target data map(ompx_hold,tofrom: s.y){{$}}

// PRT-OA-NEXT:   // #pragma acc data copy(s.y){{$}}

//  PRT-O-NEXT:   #pragma omp target data map([[PRESENT_MT]]: s){{$}}

// PRT-OA-NEXT:   // #pragma acc data present(s){{$}}

//

//    PRT-NEXT:   ;

//    PRT-NEXT: }

CASE(caseDataMembersConcat2) {

  struct S { int x; int y; } s;

  PRINT_VAR_INFO(s.x);

  PRINT_VAR_INFO(s);

  #pragma acc data copyout(s.x)

  #pragma acc data copy(s.y)

  #pragma acc data present(s)

  USE_VAR(s.x = 1);

}



CASE(caseDataMemberFullStruct) {

  struct S { int x; int y; int z; } s;

  PRINT_VAR_INFO(s.y);
@@ -576,6 +545,11 @@ CASE(caseDataSubarrayPresent) { 
//

//    PRT-NEXT:   ;

//    PRT-NEXT: }

//

// This case is also invalid when -fopenacc-present-omp=no-present because the

// present clause then attempts to extend a partial mapping of an object.  For

// now, the runtime doesn't complain.  If it does in the future, we'll have to

// change the expected results here.

CASE(caseDataSubarrayDisjoint) {

  int arr[4];

  PRINT_SUBARRAY_INFO(arr, 0, 2);
@@ -637,37 +611,6 @@ CASE(caseDataSubarrayOverlapEnd) { 
  USE_VAR(arr[2] = 1);

}



//   PRT-LABEL: {{.*}}caseDataSubarrayConcat2{{.*}} {

//    PRT-NEXT:   int arr[4];

//    PRT-NEXT:   {{(PRINT_SUBARRAY_INFO|fprintf)\(.*\);}}

//    PRT-NEXT:   {{(PRINT_SUBARRAY_INFO|fprintf)\(.*\);}}

//

//  PRT-A-NEXT:   #pragma acc data copyout(arr[0:2]){{$}}

// PRT-AO-NEXT:   // #pragma omp target data map(ompx_hold,from: arr[0:2]){{$}}

//  PRT-A-NEXT:   #pragma acc data copy(arr[2:2]){{$}}

// PRT-AO-NEXT:   // #pragma omp target data map(ompx_hold,tofrom: arr[2:2]){{$}}

//  PRT-A-NEXT:   #pragma acc data present(arr[0:4]){{$}}

// PRT-AO-NEXT:   // #pragma omp target data map([[PRESENT_MT]]: arr[0:4]){{$}}

//

//  PRT-O-NEXT:   #pragma omp target data map(ompx_hold,from: arr[0:2]){{$}}

// PRT-OA-NEXT:   // #pragma acc data copyout(arr[0:2]){{$}}

//  PRT-O-NEXT:   #pragma omp target data map(ompx_hold,tofrom: arr[2:2]){{$}}

// PRT-OA-NEXT:   // #pragma acc data copy(arr[2:2]){{$}}

//  PRT-O-NEXT:   #pragma omp target data map([[PRESENT_MT]]: arr[0:4]){{$}}

// PRT-OA-NEXT:   // #pragma acc data present(arr[0:4]){{$}}

//

//    PRT-NEXT:   ;

//    PRT-NEXT: }

CASE(caseDataSubarrayConcat2) {

  int arr[4];

  PRINT_SUBARRAY_INFO(arr, 0, 2);

  PRINT_SUBARRAY_INFO(arr, 0, 4);

  #pragma acc data copyout(arr[0:2])

  #pragma acc data copy(arr[2:2])

  #pragma acc data present(arr[0:4])

  USE_VAR(arr[0] = 1);

}



CASE(caseDataSubarrayNonSubarray) {

  int arr[5];

  PRINT_SUBARRAY_INFO(arr, 1, 2);
@@ -784,6 +727,10 @@ CASE(caseParallelMemberAbsent) { 
  USE_VAR(s.y = 1);

}



// This case is also invalid when -fopenacc-present-omp=no-present because the

// present clause then attempts to extend a partial mapping of an object.  For

// now, the runtime doesn't complain.  If it does in the future, we'll have to

// change the expected results here.

CASE(caseParallelMembersDisjoint) {

  struct S { int x; int y; } s;

  PRINT_VAR_INFO(s.x);
@@ -793,16 +740,6 @@ CASE(caseParallelMembersDisjoint) { 
  USE_VAR(s.y = 1);

}



CASE(caseParallelMembersConcat2) {

  struct S { int x; int y; } s;

  PRINT_VAR_INFO(s.x);

  PRINT_VAR_INFO(s);

  #pragma acc data copyout(s.x)

  #pragma acc data copy(s.y)

  #pragma acc parallel present(s)

  USE_VAR(s.x = 1);

}



CASE(caseParallelMemberFullStruct) {

  struct S { int x; int y; int z; } s;

  PRINT_VAR_INFO(s.y);
@@ -828,6 +765,10 @@ CASE(caseParallelSubarrayPresent) { 
  }

}



// This case is also invalid when -fopenacc-present-omp=no-present because the

// present clause then attempts to extend a partial mapping of an object.  For

// now, the runtime doesn't complain.  If it does in the future, we'll have to

// change the expected results here.

CASE(caseParallelSubarrayDisjoint) {

  int arr[4];

  PRINT_SUBARRAY_INFO(arr, 0, 2);
@@ -855,16 +796,6 @@ CASE(caseParallelSubarrayOverlapEnd) { 
  USE_VAR(arr[4] = 1);

}



CASE(caseParallelSubarrayConcat2) {

  int arr[4];

  PRINT_SUBARRAY_INFO(arr, 0, 2);

  PRINT_SUBARRAY_INFO(arr, 0, 4);

  #pragma acc data copyout(arr[0:2])

  #pragma acc data copy(arr[2:2])

  #pragma acc parallel present(arr[0:4])

  USE_VAR(arr[0] = 1);

}



CASE(caseParallelSubarrayNonSubarray) {

  int arr[10];

  PRINT_SUBARRAY_INFO(arr, 3, 4);

clang/test/OpenACC/directives/Tests/update.c

+0 −170
Original line number Diff line number Diff line
@@ -50,14 +50,12 @@ 
// DEFINE:   %{check-case}( caseMemberPresent        %,                                        %,                                  %) && \

// DEFINE:   %{check-case}( caseMemberAbsent         %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %) && \

// DEFINE:   %{check-case}( caseMembersDisjoint      %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %) && \

// DEFINE:   %{check-case}( caseMembersConcat2       %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %) && \

// DEFINE:   %{check-case}( caseMemberFullStruct     %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %) && \

// DEFINE:   %{check-case}( caseSubarrayPresent      %,                                        %,                                  %) && \

// DEFINE:   %{check-case}( caseSubarrayPresent2     %,                                        %,                                  %) && \

// DEFINE:   %{check-case}( caseSubarrayDisjoint     %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %) && \

// DEFINE:   %{check-case}( caseSubarrayOverlapStart %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %) && \

// DEFINE:   %{check-case}( caseSubarrayOverlapEnd   %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %) && \

// DEFINE:   %{check-case}( caseSubarrayConcat2      %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %) && \

// DEFINE:   %{check-case}( caseSubarrayNonSubarray  %, %if-tgt-host<|%{NOT_CRASH_IF_PRESENT}> %, %if-tgt-host<|%{NOT_IF_PRESENT}> %)



// Generate the enum of cases.
@@ -832,68 +830,6 @@ CASE(caseMembersDisjoint) { 
  }

}



// DMP-LABEL: FunctionDecl {{.*}} prev {{.*}} caseMembersConcat2

// PRT-LABEL: {{.*}}caseMembersConcat2{{.*}} {

CASE(caseMembersConcat2) {

  struct S { int i; int j; } s, h, d;

  PRINT_VAR_INFO(h.i);

  PRINT_VAR_INFO(h);



  #pragma acc data create(s.i, h.i, d.i)

  #pragma acc data create(s.j, h.j, d.j)

  {

    setHostInt(s.i, 10);

    setHostInt(s.j, 20);

    setHostInt(h.i, 30);

    setHostInt(h.j, 40);

    setHostInt(d.i, 50);

    setHostInt(d.j, 60);

    setDeviceInt(s.i, 11);

    setDeviceInt(s.j, 21);

    setDeviceInt(h.i, 31);

    setDeviceInt(h.j, 41);

    setDeviceInt(d.i, 51);

    setDeviceInt(d.j, 61);



    // We need multiple directives here so we can control which clause

    // produces the runtime error.  We vary which clause produces the runtime

    // error across the various CASE_* that produce it.

    #pragma acc update host(h) IF_PRESENT

    #pragma acc update self(s) device(d) IF_PRESENT



    //  EXE-OFF-caseMembersConcat2-PASS-NEXT:   host s.i=10{{$}}

    //  EXE-OFF-caseMembersConcat2-PASS-NEXT:   host s.j=20{{$}}

    //  EXE-OFF-caseMembersConcat2-PASS-NEXT:   host h.i=30{{$}}

    //  EXE-OFF-caseMembersConcat2-PASS-NEXT:   host h.j=40{{$}}

    //  EXE-OFF-caseMembersConcat2-PASS-NEXT:   host d.i=50{{$}}

    //  EXE-OFF-caseMembersConcat2-PASS-NEXT:   host d.j=60{{$}}

    // EXE-HOST-caseMembersConcat2-PASS-NEXT:   host s.i=11{{$}}

    // EXE-HOST-caseMembersConcat2-PASS-NEXT:   host s.j=21{{$}}

    // EXE-HOST-caseMembersConcat2-PASS-NEXT:   host h.i=31{{$}}

    // EXE-HOST-caseMembersConcat2-PASS-NEXT:   host h.j=41{{$}}

    // EXE-HOST-caseMembersConcat2-PASS-NEXT:   host d.i=51{{$}}

    // EXE-HOST-caseMembersConcat2-PASS-NEXT:   host d.j=61{{$}}

    //      EXE-caseMembersConcat2-PASS-NEXT: device s.i=11{{$}}

    //      EXE-caseMembersConcat2-PASS-NEXT: device s.j=21{{$}}

    //      EXE-caseMembersConcat2-PASS-NEXT: device h.i=31{{$}}

    //      EXE-caseMembersConcat2-PASS-NEXT: device h.j=41{{$}}

    //      EXE-caseMembersConcat2-PASS-NEXT: device d.i=51{{$}}

    //      EXE-caseMembersConcat2-PASS-NEXT: device d.j=61{{$}}

    printHostInt(s.i);

    printHostInt(s.j);

    printHostInt(h.i);

    printHostInt(h.j);

    printHostInt(d.i);

    printHostInt(d.j);

    printDeviceInt(s.i);

    printDeviceInt(s.j);

    printDeviceInt(h.i);

    printDeviceInt(h.j);

    printDeviceInt(d.i);

    printDeviceInt(d.j);

  }

}



// DMP-LABEL: FunctionDecl {{.*}} prev {{.*}} caseMemberFullStruct

// PRT-LABEL: {{.*}}caseMemberFullStruct{{.*}} {

CASE(caseMemberFullStruct) {
@@ -1560,112 +1496,6 @@ CASE(caseSubarrayOverlapEnd) { 
  }

}



// DMP-LABEL: FunctionDecl {{.*}} prev {{.*}} caseSubarrayConcat2

// PRT-LABEL: {{.*}}caseSubarrayConcat2{{.*}} {

CASE(caseSubarrayConcat2) {

  int s[4];

  int h[4];

  int d[4];

  PRINT_SUBARRAY_INFO(s, 0, 2);

  PRINT_SUBARRAY_INFO(s, 0, 4);



  #pragma acc data create(s[0:2], h[0:2], d[0:2])

  #pragma acc data create(s[2:2], h[2:2], d[2:2])

  {

    setHostInt(s[0], 10);

    setHostInt(s[1], 20);

    setHostInt(s[2], 30);

    setHostInt(s[3], 40);

    setHostInt(h[0], 50);

    setHostInt(h[1], 60);

    setHostInt(h[2], 70);

    setHostInt(h[3], 80);

    setHostInt(d[0], 90);

    setHostInt(d[1], 100);

    setHostInt(d[2], 110);

    setHostInt(d[3], 120);

    setDeviceInt(s[0], 11);

    setDeviceInt(s[1], 21);

    setDeviceInt(s[2], 31);

    setDeviceInt(s[3], 41);

    setDeviceInt(h[0], 51);

    setDeviceInt(h[1], 61);

    setDeviceInt(h[2], 71);

    setDeviceInt(h[3], 81);

    setDeviceInt(d[0], 91);

    setDeviceInt(d[1], 101);

    setDeviceInt(d[2], 111);

    setDeviceInt(d[3], 121);



    // We need multiple directives here so we can control which clause

    // produces the runtime error.  We vary which clause produces the runtime

    // error across the various CASE_* that produce it.

    #pragma acc update self(s[0:4]) IF_PRESENT

    #pragma acc update host(h[0:4]) device(d[0:4]) IF_PRESENT



    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host s[0]=10{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host s[1]=20{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host s[2]=30{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host s[3]=40{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host h[0]=50{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host h[1]=60{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host h[2]=70{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host h[3]=80{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host d[0]=90{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host d[1]=100{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host d[2]=110{{$}}

    //  EXE-OFF-caseSubarrayConcat2-PASS-NEXT:   host d[3]=120{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host s[0]=11{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host s[1]=21{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host s[2]=31{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host s[3]=41{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host h[0]=51{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host h[1]=61{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host h[2]=71{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host h[3]=81{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host d[0]=91{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host d[1]=101{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host d[2]=111{{$}}

    // EXE-HOST-caseSubarrayConcat2-PASS-NEXT:   host d[3]=121{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device s[0]=11{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device s[1]=21{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device s[2]=31{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device s[3]=41{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device h[0]=51{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device h[1]=61{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device h[2]=71{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device h[3]=81{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device d[0]=91{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device d[1]=101{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device d[2]=111{{$}}

    //      EXE-caseSubarrayConcat2-PASS-NEXT: device d[3]=121{{$}}

    printHostInt(s[0]);

    printHostInt(s[1]);

    printHostInt(s[2]);

    printHostInt(s[3]);

    printHostInt(h[0]);

    printHostInt(h[1]);

    printHostInt(h[2]);

    printHostInt(h[3]);

    printHostInt(d[0]);

    printHostInt(d[1]);

    printHostInt(d[2]);

    printHostInt(d[3]);

    printDeviceInt(s[0]);

    printDeviceInt(s[1]);

    printDeviceInt(s[2]);

    printDeviceInt(s[3]);

    printDeviceInt(h[0]);

    printDeviceInt(h[1]);

    printDeviceInt(h[2]);

    printDeviceInt(h[3]);

    printDeviceInt(d[0]);

    printDeviceInt(d[1]);

    printDeviceInt(d[2]);

    printDeviceInt(d[3]);

  }

}



// DMP-LABEL: FunctionDecl {{.*}} prev {{.*}} caseSubarrayNonSubarray

// PRT-LABEL: {{.*}}caseSubarrayNonSubarray{{.*}} {

CASE(caseSubarrayNonSubarray) {

openmp/libacc2omp/test/api/data.c

+1 −109

File changed.

Preview size limit exceeded, changes collapsed.

openmp/libacc2omp/test/directives/data-extension-errors.c

+56 −116

File changed.

Preview size limit exceeded, changes collapsed.