Home | History | Annotate | Download | only in OpenMP
      1 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
      2 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
      3 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
      4 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
      5 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
      6 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
      7 // expected-no-diagnostics
      8 #ifndef HEADER
      9 #define HEADER
     10 
     11 
     12 // CHECK-DAG: [[GA:@.+]] = global double 1.000000e+00
     13 // CHECK-DAG: [[GB:@.+]] = global double 2.000000e+00
     14 // CHECK-DAG: [[GC:@.+]] = global double 3.000000e+00
     15 // CHECK-DAG: [[GD:@.+]] = global double 4.000000e+00
     16 // CHECK-DAG: [[FA:@.+]] = internal global float 5.000000e+00
     17 // CHECK-DAG: [[FB:@.+]] = internal global float 6.000000e+00
     18 // CHECK-DAG: [[FC:@.+]] = internal global float 7.000000e+00
     19 // CHECK-DAG: [[FD:@.+]] = internal global float 8.000000e+00
     20 // CHECK-DAG: [[BA:@.+]] = internal global float 9.000000e+00
     21 // CHECK-DAG: [[BB:@.+]] = internal global float 1.000000e+01
     22 // CHECK-DAG: [[BC:@.+]] = internal global float 1.100000e+01
     23 // CHECK-DAG: [[BD:@.+]] = internal global float 1.200000e+01
     24 double Ga = 1.0;
     25 double Gb = 2.0;
     26 double Gc = 3.0;
     27 double Gd = 4.0;
     28 
     29 // CHECK: define {{.*}} @{{.*}}foo{{.*}}(
     30 // CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
     31 // CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
     32 // CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
     33 // CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
     34 // CHECK: [[LA:%.+]] = alloca i16
     35 // CHECK: [[LB:%.+]] = alloca i16
     36 // CHECK: [[LC:%.+]] = alloca i16
     37 // CHECK: [[LD:%.+]] = alloca i16
     38 int foo(short a, short b, short c, short d){
     39   static float Sa = 5.0;
     40   static float Sb = 6.0;
     41   static float Sc = 7.0;
     42   static float Sd = 8.0;
     43 
     44   // CHECK-DAG:    [[VALLB:%.+]] = load i16, i16* [[LB]],
     45   // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb,
     46   // CHECK-DAG:    [[VALFB:%.+]] = load float, float* @_ZZ3foossssE2Sb,
     47   // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc,
     48   // CHECK-DAG:    [[VALLC:%.+]] = load i16, i16* [[LC]],
     49   // CHECK-DAG:    [[VALFC:%.+]] = load float, float* @_ZZ3foossssE2Sc,
     50   // CHECK-DAG:    [[VALLD:%.+]] = load i16, i16* [[LD]],
     51   // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd,
     52   // CHECK-DAG:    [[VALFD:%.+]] = load float, float* @_ZZ3foossssE2Sd,
     53 
     54   // 3 local vars being captured.
     55 
     56   // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
     57   // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
     58   // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
     59   // CHECK-DAG: [[CPTRLB:%.+]] = inttoptr i[[sz]] [[CVALLB]] to i8*
     60   // CHECK-DAG: store i8* [[CPTRLB]], i8** [[GEPLB:%.+]],
     61   // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
     62 
     63   // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
     64   // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
     65   // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
     66   // CHECK-DAG: [[CPTRLC:%.+]] = inttoptr i[[sz]] [[CVALLC]] to i8*
     67   // CHECK-DAG: store i8* [[CPTRLC]], i8** [[GEPLC:%.+]],
     68   // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
     69 
     70   // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
     71   // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
     72   // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
     73   // CHECK-DAG: [[CPTRLD:%.+]] = inttoptr i[[sz]] [[CVALLD]] to i8*
     74   // CHECK-DAG: store i8* [[CPTRLD]], i8** [[GEPLD:%.+]],
     75   // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
     76 
     77   // 3 static vars being captured.
     78 
     79   // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
     80   // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
     81   // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
     82   // CHECK-DAG: [[CPTRFB:%.+]] = inttoptr i[[sz]] [[CVALFB]] to i8*
     83   // CHECK-DAG: store i8* [[CPTRFB]], i8** [[GEPFB:%.+]],
     84   // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
     85 
     86   // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
     87   // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
     88   // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
     89   // CHECK-DAG: [[CPTRFC:%.+]] = inttoptr i[[sz]] [[CVALFC]] to i8*
     90   // CHECK-DAG: store i8* [[CPTRFC]], i8** [[GEPFC:%.+]],
     91   // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
     92 
     93   // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
     94   // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
     95   // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
     96   // CHECK-DAG: [[CPTRFD:%.+]] = inttoptr i[[sz]] [[CVALFD]] to i8*
     97   // CHECK-DAG: store i8* [[CPTRFD]], i8** [[GEPFD:%.+]],
     98   // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
     99 
    100   // 3 static global vars being captured.
    101 
    102   // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
    103   // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
    104   // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
    105   // CHECK-64-DAG: [[CPTRGB:%.+]] = inttoptr i[[sz]] [[CVALGB]] to i8*
    106   // CHECK-64-DAG: store i8* [[CPTRGB]], i8** [[GEPGB:%.+]],
    107   // CHECK-32-DAG: store i8* bitcast (double* @Gb to i8*), i8** [[GEPGB:%.+]],
    108   // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    109 
    110   // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
    111   // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
    112   // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
    113   // CHECK-64-DAG: [[CPTRGC:%.+]] = inttoptr i[[sz]] [[CVALGC]] to i8*
    114   // CHECK-64-DAG: store i8* [[CPTRGC]], i8** [[GEPGC:%.+]],
    115   // CHECK-32-DAG: store i8* bitcast (double* @Gc to i8*), i8** [[GEPGC:%.+]],
    116   // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    117 
    118   // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
    119   // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
    120   // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
    121   // CHECK-64-DAG: [[CPTRGD:%.+]] = inttoptr i[[sz]] [[CVALGD]] to i8*
    122   // CHECK-64-DAG: store i8* [[CPTRGD]], i8** [[GEPGD:%.+]],
    123   // CHECK-32-DAG: store i8* bitcast (double* @Gd to i8*), i8** [[GEPGD:%.+]],
    124   // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    125 
    126   // CHECK: call i32 @__tgt_target
    127   // CHECK: call void [[OFFLOADF:@.+]](
    128   // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
    129   #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
    130   {
    131     b += 1;
    132     Gb += 1.0;
    133     Sb += 1.0;
    134 
    135     // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
    136     // The parallel region only uses 3 captures.
    137     // CHECK:     call {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), {{.+}}* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}})
    138     // CHECK:     call void @.omp_outlined.(i32* %{{.+}}, i32* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}})
    139     // Capture d, Gd, Sd,
    140 
    141     // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}},
    142     #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
    143     {
    144       d += 1;
    145       Gd += 1.0;
    146       Sd += 1.0;
    147     }
    148   }
    149   return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
    150 }
    151 
    152 // CHECK: define {{.*}} @{{.*}}bar{{.*}}(
    153 // CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
    154 // CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
    155 // CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
    156 // CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
    157 // CHECK: [[LA:%.+]] = alloca i16
    158 // CHECK: [[LB:%.+]] = alloca i16
    159 // CHECK: [[LC:%.+]] = alloca i16
    160 // CHECK: [[LD:%.+]] = alloca i16
    161 int bar(short a, short b, short c, short d){
    162   static float Sa = 9.0;
    163   static float Sb = 10.0;
    164   static float Sc = 11.0;
    165   static float Sd = 12.0;
    166 
    167   // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}})
    168   // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, i16* dereferenceable(2) [[A:%.+]], i16* dereferenceable(2) [[B:%.+]], i16* dereferenceable(2) [[C:%.+]], i16* dereferenceable(2) [[D:%.+]])
    169   // Capture a, b, c, d
    170   // CHECK: [[ALLOCLA:%.+]] = alloca i16
    171   // CHECK: [[ALLOCLB:%.+]] = alloca i16
    172   // CHECK: [[ALLOCLC:%.+]] = alloca i16
    173   // CHECK: [[ALLOCLD:%.+]] = alloca i16
    174   // CHECK: [[LLA:%.+]] = load i16*, i16** [[ALLOCLA]],
    175   // CHECK: [[LLB:%.+]] = load i16*, i16** [[ALLOCLB]],
    176   // CHECK: [[LLC:%.+]] = load i16*, i16** [[ALLOCLC]],
    177   // CHECK: [[LLD:%.+]] = load i16*, i16** [[ALLOCLD]],
    178   #pragma omp parallel
    179   {
    180     // CHECK-DAG:    [[VALLB:%.+]] = load i16, i16* [[LLB]],
    181     // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb,
    182     // CHECK-DAG:    [[VALFB:%.+]] = load float, float* @_ZZ3barssssE2Sb,
    183     // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc,
    184     // CHECK-DAG:    [[VALLC:%.+]] = load i16, i16* [[LLC]],
    185     // CHECK-DAG:    [[VALFC:%.+]] = load float, float* @_ZZ3barssssE2Sc,
    186     // CHECK-DAG:    [[VALLD:%.+]] = load i16, i16* [[LLD]],
    187     // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd,
    188     // CHECK-DAG:    [[VALFD:%.+]] = load float, float* @_ZZ3barssssE2Sd,
    189 
    190     // 3 local vars being captured.
    191 
    192     // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
    193     // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
    194     // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
    195     // CHECK-DAG: [[CPTRLB:%.+]] = inttoptr i[[sz]] [[CVALLB]] to i8*
    196     // CHECK-DAG: store i8* [[CPTRLB]], i8** [[GEPLB:%.+]],
    197     // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    198 
    199     // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
    200     // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
    201     // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
    202     // CHECK-DAG: [[CPTRLC:%.+]] = inttoptr i[[sz]] [[CVALLC]] to i8*
    203     // CHECK-DAG: store i8* [[CPTRLC]], i8** [[GEPLC:%.+]],
    204     // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    205 
    206     // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
    207     // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
    208     // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
    209     // CHECK-DAG: [[CPTRLD:%.+]] = inttoptr i[[sz]] [[CVALLD]] to i8*
    210     // CHECK-DAG: store i8* [[CPTRLD]], i8** [[GEPLD:%.+]],
    211     // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    212 
    213     // 3 static vars being captured.
    214 
    215     // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
    216     // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
    217     // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
    218     // CHECK-DAG: [[CPTRFB:%.+]] = inttoptr i[[sz]] [[CVALFB]] to i8*
    219     // CHECK-DAG: store i8* [[CPTRFB]], i8** [[GEPFB:%.+]],
    220     // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    221 
    222     // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
    223     // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
    224     // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
    225     // CHECK-DAG: [[CPTRFC:%.+]] = inttoptr i[[sz]] [[CVALFC]] to i8*
    226     // CHECK-DAG: store i8* [[CPTRFC]], i8** [[GEPFC:%.+]],
    227     // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    228 
    229     // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
    230     // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
    231     // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
    232     // CHECK-DAG: [[CPTRFD:%.+]] = inttoptr i[[sz]] [[CVALFD]] to i8*
    233     // CHECK-DAG: store i8* [[CPTRFD]], i8** [[GEPFD:%.+]],
    234     // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    235 
    236     // 3 static global vars being captured.
    237 
    238     // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
    239     // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
    240     // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
    241     // CHECK-64-DAG: [[CPTRGB:%.+]] = inttoptr i[[sz]] [[CVALGB]] to i8*
    242     // CHECK-64-DAG: store i8* [[CPTRGB]], i8** [[GEPGB:%.+]],
    243     // CHECK-32-DAG: store i8* bitcast (double* @Gb to i8*), i8** [[GEPGB:%.+]],
    244     // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    245 
    246     // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
    247     // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
    248     // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
    249     // CHECK-64-DAG: [[CPTRGC:%.+]] = inttoptr i[[sz]] [[CVALGC]] to i8*
    250     // CHECK-64-DAG: store i8* [[CPTRGC]], i8** [[GEPGC:%.+]],
    251     // CHECK-32-DAG: store i8* bitcast (double* @Gc to i8*), i8** [[GEPGC:%.+]],
    252     // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    253 
    254     // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
    255     // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
    256     // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
    257     // CHECK-64-DAG: [[CPTRGD:%.+]] = inttoptr i[[sz]] [[CVALGD]] to i8*
    258     // CHECK-64-DAG: store i8* [[CPTRGD]], i8** [[GEPGD:%.+]],
    259     // CHECK-32-DAG: store i8* bitcast (double* @Gd to i8*), i8** [[GEPGD:%.+]],
    260     // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
    261 
    262     // CHECK: call i32 @__tgt_target
    263     // CHECK: call void [[OFFLOADF:@.+]](
    264     // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
    265     #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
    266     {
    267       b += 1;
    268       Gb += 1.0;
    269       Sb += 1.0;
    270 
    271       // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
    272       // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}})
    273 
    274       // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}})
    275       // Capture d, Gd, Sd
    276       #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
    277       {
    278         d += 1;
    279         Gd += 1.0;
    280         Sd += 1.0;
    281       }
    282     }
    283   }
    284   return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
    285 }
    286 
    287 #endif
    288