target_parallel_for_codegen.cpp 52 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540541542543544545546547548549550551552553554555556557558559560561562563564565566567568569570571572573574575576577578579580581582583584585586587588589590591592593594595596597598599600601602603604605606607608609610611612613614615616617618619620621622623624625626627628629630631632633634635636637638639640641642643644645646647648649650651652653654655656657658659660661662663664665666667668669670671672673674675676677678679680681682683684685686687688689690691692693694695696697698699700701702703704705706707708709710711712713714715716717718719720721722723724725726727728729730731732733734735736737738739740741742743744745746747748749750751752753754755756757758759760761762763764765766767768769770771772773774775776777778779780781782783784785786787788789790791792793794795796797798799800801802803804805806807808809810811812813814815816817818819820821822823824825826827828829830831832833834835836837838839840841
  1. // Test host codegen.
  2. // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CHECK --check-prefix CHECK-64
  3. // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
  4. // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CHECK --check-prefix CHECK-64
  5. // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CHECK --check-prefix CHECK-32
  6. // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
  7. // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CHECK --check-prefix CHECK-32
  8. // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s
  9. // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
  10. // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s
  11. // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s
  12. // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
  13. // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s
  14. // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
  15. // Test target codegen - host bc file has to be created first.
  16. // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
  17. // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix TCHECK --check-prefix TCHECK-64
  18. // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
  19. // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix TCHECK --check-prefix TCHECK-64
  20. // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
  21. // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix TCHECK --check-prefix TCHECK-32
  22. // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
  23. // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix TCHECK --check-prefix TCHECK-32
  24. // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
  25. // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1 %s
  26. // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
  27. // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1 %s
  28. // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
  29. // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1 %s
  30. // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
  31. // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1 %s
  32. // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
  33. // expected-no-diagnostics
  34. #ifndef HEADER
  35. #define HEADER
  36. // CHECK-DAG: %struct.ident_t = type { i32, i32, i32, i32, i8* }
  37. // CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
  38. // CHECK-DAG: [[DEF_LOC:@.+]] = private unnamed_addr global %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
  39. // CHECK-DAG: [[TT:%.+]] = type { i64, i8 }
  40. // CHECK-DAG: [[S1:%.+]] = type { double }
  41. // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
  42. // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 }
  43. // We have 8 target regions, but only 7 that actually will generate offloading
  44. // code, only 6 will have mapped arguments, and only 4 have all-constant map
  45. // sizes.
  46. // CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [3 x i64] [i64 2, i64 4, i64 4]
  47. // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 800]
  48. // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 2]
  49. // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
  50. // CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [10 x i64] [i64 800, i64 547, i64 800, i64 547, i64 547, i64 800, i64 800, i64 547, i64 547, i64 800]
  51. // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 2, i64 40]
  52. // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547]
  53. // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40]
  54. // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547]
  55. // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547]
  56. // CHECK-DAG: @{{.*}} = weak constant i8 0
  57. // CHECK-DAG: @{{.*}} = weak constant i8 0
  58. // CHECK-DAG: @{{.*}} = weak constant i8 0
  59. // CHECK-DAG: @{{.*}} = weak constant i8 0
  60. // CHECK-DAG: @{{.*}} = weak constant i8 0
  61. // CHECK-DAG: @{{.*}} = weak constant i8 0
  62. // CHECK-DAG: @{{.*}} = weak constant i8 0
  63. // TCHECK: @{{.+}} = weak constant [[ENTTY]]
  64. // TCHECK: @{{.+}} = weak constant [[ENTTY]]
  65. // TCHECK: @{{.+}} = weak constant [[ENTTY]]
  66. // TCHECK: @{{.+}} = weak constant [[ENTTY]]
  67. // TCHECK: @{{.+}} = weak constant [[ENTTY]]
  68. // TCHECK: @{{.+}} = weak constant [[ENTTY]]
  69. // TCHECK: @{{.+}} = weak constant [[ENTTY]]
  70. // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
  71. // Check target registration is registered as a Ctor.
  72. // CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @.omp_offloading.requires_reg, i8* null }]
  73. template<typename tx, typename ty>
  74. struct TT{
  75. tx X;
  76. ty Y;
  77. };
  78. // CHECK-LABEL: get_val
  79. long long get_val() { return 0; }
  80. // CHECK: define {{.*}}[[FOO:@.+]](
  81. int foo(int n) {
  82. int a = 0;
  83. short aa = 0;
  84. float b[10];
  85. float bn[n];
  86. double c[5][10];
  87. double cn[5][n];
  88. TT<long long, char> d;
  89. // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i64* null, i64* null, i32 1, i32 0)
  90. // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
  91. // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
  92. // CHECK: [[FAIL]]
  93. // CHECK: call void [[HVT0:@.+]]()
  94. // CHECK-NEXT: br label %[[END]]
  95. // CHECK: [[END]]
  96. #pragma omp target parallel for
  97. for (int i = 3; i < 32; i += 5) {
  98. #pragma omp cancel for
  99. #pragma omp cancellation point for
  100. }
  101. // CHECK: call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}}, i{{32|64}}{{[*]*}} {{[^)]+}})
  102. long long k = get_val();
  103. #pragma omp target parallel for if(target: 0) linear(k : 3) schedule(dynamic)
  104. for (int i = 10; i > 1; i--) {
  105. a += 1;
  106. }
  107. // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_nowait(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZET2]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT2]], i32 0, i32 0), i32 1, i32 0)
  108. // CHECK-DAG: [[BP]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR:%[^,]+]], i32 0, i32 0
  109. // CHECK-DAG: [[P]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR:%[^,]+]], i32 0, i32 0
  110. // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 0
  111. // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR]], i32 0, i32 0
  112. // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]*
  113. // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]*
  114. // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], i[[SZ]]* [[CBPADDR0]],
  115. // CHECK-DAG: store i[[SZ]] [[VAL0]], i[[SZ]]* [[CPADDR0]],
  116. // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 1
  117. // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR]], i32 0, i32 1
  118. // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*
  119. // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*
  120. // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], i[[SZ]]* [[CBPADDR1]],
  121. // CHECK-DAG: store i[[SZ]] [[VAL1]], i[[SZ]]* [[CPADDR1]],
  122. // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 1
  123. // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR]], i32 0, i32 1
  124. // CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to i[[SZ]]*
  125. // CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to i[[SZ]]*
  126. // CHECK-DAG: store i[[SZ]] [[VAL2:%.+]], i[[SZ]]* [[CBPADDR2]],
  127. // CHECK-DAG: store i[[SZ]] [[VAL2]], i[[SZ]]* [[CPADDR2]],
  128. // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
  129. // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
  130. // CHECK: [[FAIL]]
  131. // CHECK: call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^)]+}})
  132. // CHECK-NEXT: br label %[[END]]
  133. // CHECK: [[END]]
  134. int lin = 12;
  135. #pragma omp target parallel for if(target: 1) linear(lin, a : get_val()) nowait
  136. for (unsigned long long it = 2000; it >= 600; it-=400) {
  137. aa += 1;
  138. }
  139. // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 10
  140. // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
  141. // CHECK: [[IFTHEN]]
  142. // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[SIZET3]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT3]], i32 0, i32 0), i32 1, i32 0)
  143. // CHECK-DAG: [[BPR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%[^,]+]], i32 0, i32 0
  144. // CHECK-DAG: [[PR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%[^,]+]], i32 0, i32 0
  145. // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 0
  146. // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 0
  147. // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]*
  148. // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]*
  149. // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], i[[SZ]]* [[CBPADDR0]],
  150. // CHECK-DAG: store i[[SZ]] [[VAL0]], i[[SZ]]* [[CPADDR0]],
  151. // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 1
  152. // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 1
  153. // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*
  154. // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*
  155. // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], i[[SZ]]* [[CBPADDR1]],
  156. // CHECK-DAG: store i[[SZ]] [[VAL1]], i[[SZ]]* [[CPADDR1]],
  157. // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
  158. // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
  159. // CHECK: [[FAIL]]
  160. // CHECK: call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}})
  161. // CHECK-NEXT: br label %[[END]]
  162. // CHECK: [[END]]
  163. // CHECK-NEXT: br label %[[IFEND:.+]]
  164. // CHECK: [[IFELSE]]
  165. // CHECK: call void [[HVT3]]({{[^,]+}}, {{[^,]+}})
  166. // CHECK-NEXT: br label %[[IFEND]]
  167. // CHECK: [[IFEND]]
  168. #pragma omp target parallel for if(target: n>10)
  169. for (short it = 6; it <= 20; it-=-4) {
  170. a += 1;
  171. aa += 1;
  172. }
  173. // We capture 3 VLA sizes in this target region
  174. // CHECK: [[A_VAL:%.+]] = load i32, i32* %{{.+}},
  175. // CHECK: store i32 [[A_VAL]], i32* [[A_CADDR:%.+]],
  176. // CHECK-64: [[A_VAL:%.+]] = load i32, i32* %{{.+}},
  177. // CHECK-64: [[A_ADDR:%.+]] = bitcast i[[SZ]]* [[A_CADDR:%.+]] to i32*
  178. // CHECK-64: store i32 [[A_VAL]], i32* [[A_ADDR]],
  179. // CHECK-64: [[A_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[A_CADDR]],
  180. // CHECK-32: [[A_VAL:%.+]] = load i32, i32* %{{.+}},
  181. // CHECK-32: store i32 [[A_VAL]], i32* [[A_CADDR:%.+]],
  182. // CHECK-32: [[A_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[A_CADDR]],
  183. // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 20
  184. // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]
  185. // CHECK: [[TRY]]
  186. // CHECK-64: [[BNSIZE:%.+]] = mul nuw i64 [[VLA0:%.+]], 4
  187. // CHECK-32: [[BNSZSIZE:%.+]] = mul nuw i32 [[VLA0:%.+]], 4
  188. // CHECK-32: [[BNSIZE:%.+]] = sext i32 [[BNSZSIZE]] to i64
  189. // CHECK: [[CNELEMSIZE2:%.+]] = mul nuw i[[SZ]] 5, [[VLA1:%.+]]
  190. // CHECK-64: [[CNSIZE:%.+]] = mul nuw i64 [[CNELEMSIZE2]], 8
  191. // CHECK-32: [[CNSZSIZE:%.+]] = mul nuw i32 [[CNELEMSIZE2]], 8
  192. // CHECK-32: [[CNSIZE:%.+]] = sext i32 [[CNSZSIZE]] to i64
  193. // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 10, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([10 x i64], [10 x i64]* [[MAPT4]], i32 0, i32 0), i32 1, i32 0)
  194. // CHECK-DAG: [[BPR]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP:%[^,]+]], i32 0, i32 0
  195. // CHECK-DAG: [[PR]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P:%[^,]+]], i32 0, i32 0
  196. // CHECK-DAG: [[SR]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S:%[^,]+]], i32 0, i32 0
  197. // CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S]], i32 0, i32 [[IDX0:[0-9]+]]
  198. // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP]], i32 0, i32 [[IDX0]]
  199. // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P]], i32 0, i32 [[IDX0]]
  200. // CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S]], i32 0, i32 [[IDX1:[0-9]+]]
  201. // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP]], i32 0, i32 [[IDX1]]
  202. // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P]], i32 0, i32 [[IDX1]]
  203. // CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S]], i32 0, i32 [[IDX2:[0-9]+]]
  204. // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP]], i32 0, i32 [[IDX2]]
  205. // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P]], i32 0, i32 [[IDX2]]
  206. // CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S]], i32 0, i32 [[IDX3:[0-9]+]]
  207. // CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP]], i32 0, i32 [[IDX3]]
  208. // CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P]], i32 0, i32 [[IDX3]]
  209. // CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S]], i32 0, i32 [[IDX4:[0-9]+]]
  210. // CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP]], i32 0, i32 [[IDX4]]
  211. // CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P]], i32 0, i32 [[IDX4]]
  212. // CHECK-DAG: [[SADDR5:%.+]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S]], i32 0, i32 [[IDX5:[0-9]+]]
  213. // CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP]], i32 0, i32 [[IDX5]]
  214. // CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P]], i32 0, i32 [[IDX5]]
  215. // CHECK-DAG: [[SADDR6:%.+]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S]], i32 0, i32 [[IDX6:[0-9]+]]
  216. // CHECK-DAG: [[BPADDR6:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP]], i32 0, i32 [[IDX6]]
  217. // CHECK-DAG: [[PADDR6:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P]], i32 0, i32 [[IDX6]]
  218. // CHECK-DAG: [[SADDR7:%.+]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S]], i32 0, i32 [[IDX7:[0-9]+]]
  219. // CHECK-DAG: [[BPADDR7:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP]], i32 0, i32 [[IDX7]]
  220. // CHECK-DAG: [[PADDR7:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P]], i32 0, i32 [[IDX7]]
  221. // CHECK-DAG: [[SADDR8:%.+]] = getelementptr inbounds [10 x i64], [10 x i64]* [[S]], i32 0, i32 [[IDX8:[0-9]+]]
  222. // CHECK-DAG: [[BPADDR8:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[BP]], i32 0, i32 [[IDX8]]
  223. // CHECK-DAG: [[PADDR8:%.+]] = getelementptr inbounds [10 x i8*], [10 x i8*]* [[P]], i32 0, i32 [[IDX8]]
  224. // The names below are not necessarily consistent with the names used for the
  225. // addresses above as some are repeated.
  226. // CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CBPADDR0:%.+]],
  227. // CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CPADDR0:%.+]],
  228. // CHECK-DAG: [[CBPADDR0]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  229. // CHECK-DAG: [[CPADDR0]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  230. // CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}}
  231. // CHECK-DAG: store i[[SZ]] [[VLA1]], i[[SZ]]* [[CBPADDR1:%.+]],
  232. // CHECK-DAG: store i[[SZ]] [[VLA1]], i[[SZ]]* [[CPADDR1:%.+]],
  233. // CHECK-DAG: [[CBPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  234. // CHECK-DAG: [[CPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  235. // CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}}
  236. // CHECK-DAG: store i[[SZ]] 5, i[[SZ]]* [[CBPADDR2:%.+]],
  237. // CHECK-DAG: store i[[SZ]] 5, i[[SZ]]* [[CPADDR2:%.+]],
  238. // CHECK-DAG: [[CBPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  239. // CHECK-DAG: [[CPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  240. // CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}}
  241. // CHECK-DAG: store i[[SZ]] [[A_CVAL]], i[[SZ]]* [[CBPADDR3:%.+]],
  242. // CHECK-DAG: store i[[SZ]] [[A_CVAL]], i[[SZ]]* [[CPADDR3:%.+]],
  243. // CHECK-DAG: [[CBPADDR3]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  244. // CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  245. // CHECK-DAG: store i64 4, i64* {{%[^,]+}}
  246. // CHECK-DAG: store [10 x float]* %{{.+}}, [10 x float]** [[CBPADDR4:%.+]],
  247. // CHECK-DAG: store [10 x float]* %{{.+}}, [10 x float]** [[CPADDR4:%.+]],
  248. // CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to [10 x float]**
  249. // CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to [10 x float]**
  250. // CHECK-DAG: store i64 40, i64* {{%[^,]+}}
  251. // CHECK-DAG: store float* %{{.+}}, float** [[CBPADDR5:%.+]],
  252. // CHECK-DAG: store float* %{{.+}}, float** [[CPADDR5:%.+]],
  253. // CHECK-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to float**
  254. // CHECK-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to float**
  255. // CHECK-DAG: store i64 [[BNSIZE]], i64* {{%[^,]+}}
  256. // CHECK-DAG: store [5 x [10 x double]]* %{{.+}}, [5 x [10 x double]]** [[CBPADDR6:%.+]],
  257. // CHECK-DAG: store [5 x [10 x double]]* %{{.+}}, [5 x [10 x double]]** [[CPADDR6:%.+]],
  258. // CHECK-DAG: [[CBPADDR6]] = bitcast i8** {{%[^,]+}} to [5 x [10 x double]]**
  259. // CHECK-DAG: [[CPADDR6]] = bitcast i8** {{%[^,]+}} to [5 x [10 x double]]**
  260. // CHECK-DAG: store i64 400, i64* {{%[^,]+}}
  261. // CHECK-DAG: store double* %{{.+}}, double** [[CBPADDR7:%.+]],
  262. // CHECK-DAG: store double* %{{.+}}, double** [[CPADDR7:%.+]],
  263. // CHECK-DAG: [[CBPADDR7]] = bitcast i8** {{%[^,]+}} to double**
  264. // CHECK-DAG: [[CPADDR7]] = bitcast i8** {{%[^,]+}} to double**
  265. // CHECK-DAG: store i64 [[CNSIZE]], i64* {{%[^,]+}}
  266. // CHECK-DAG: store [[TT]]* %{{.+}}, [[TT]]** [[CBPADDR8:%.+]],
  267. // CHECK-DAG: store [[TT]]* %{{.+}}, [[TT]]** [[CPADDR8:%.+]],
  268. // CHECK-DAG: [[CBPADDR8]] = bitcast i8** {{%[^,]+}} to [[TT]]**
  269. // CHECK-DAG: [[CPADDR8]] = bitcast i8** {{%[^,]+}} to [[TT]]**
  270. // CHECK-DAG: store i64 {{12|16}}, i64* {{%[^,]+}}
  271. // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
  272. // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
  273. // CHECK: [[FAIL]]
  274. // CHECK: call void [[HVT4:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^)]+}})
  275. // CHECK-NEXT: br label %[[END]]
  276. // CHECK: [[END]]
  277. #pragma omp target parallel for if(target: n>20) schedule(static, a)
  278. for (unsigned char it = 'z'; it >= 'a'; it+=-1) {
  279. a += 1;
  280. b[2] += 1.0;
  281. bn[3] += 1.0;
  282. c[1][2] += 1.0;
  283. cn[1][3] += 1.0;
  284. d.X += 1;
  285. d.Y += 1;
  286. }
  287. return a;
  288. }
  289. // Check that the offloading functions are emitted and that the arguments are
  290. // correct and loaded correctly for the target regions in foo().
  291. // CHECK: define internal void [[HVT0]]()
  292. // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OMP_OUTLINED:@.+]] to void (i32*, i32*, ...)*))
  293. //
  294. //
  295. // CHECK: define internal {{.*}}void [[OMP_OUTLINED]](i32* noalias %.global_tid., i32* noalias %.bound_tid.)
  296. // CHECK: call i32 @__kmpc_cancel(%struct.ident_t* @
  297. // CHECK: call i32 @__kmpc_cancellationpoint(%struct.ident_t* @
  298. // CHECK: ret void
  299. // CHECK: }
  300. // CHECK: define internal void [[HVT1]](i[[SZ]] %{{.+}}, i{{32|64}}{{[*]*.*}} %{{.+}})
  301. // Create stack storage and store argument in there.
  302. // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
  303. // CHECK: alloca i{{32|64}}{{[*]*}}, align
  304. // CHECK: [[AA_CASTED:%.+]] = alloca i[[SZ]], align
  305. // CHECK: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align
  306. // CHECK-64: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i32*
  307. // CHECK-64: [[AA:%.+]] = load i32, i32* [[AA_CADDR]], align
  308. // CHECK-32: [[AA:%.+]] = load i32, i32* [[AA_ADDR]], align
  309. // CHECK-64: [[AA_C:%.+]] = bitcast i[[SZ]]* [[AA_CASTED]] to i32*
  310. // CHECK-64: store i32 [[AA]], i32* [[AA_C]], align
  311. // CHECK-32: store i32 [[AA]], i32* [[AA_CASTED]], align
  312. // CHECK: [[PARAM:%.+]] = load i[[SZ]], i[[SZ]]* [[AA_CASTED]], align
  313. // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]], i{{32|64}}{{[*]*}})* [[OMP_OUTLINED1:@.+]] to void (i32*, i32*, ...)*), i[[SZ]] [[PARAM]], i{{32|64}}{{[*]*}} %{{.+}})
  314. //
  315. //
  316. // CHECK: define internal {{.*}}void [[OMP_OUTLINED1]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}})
  317. // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
  318. // CHECK: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align
  319. // CHECK-64: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i32*
  320. // CHECK-64: [[AA:%.+]] = load i32, i32* [[AA_CADDR]], align
  321. // CHECK-32: [[AA:%.+]] = load i32, i32* [[AA_ADDR]], align
  322. // CHECK: ret void
  323. // CHECK-NEXT: }
  324. // CHECK: define internal void [[HVT2]](i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}})
  325. // Create stack storage and store argument in there.
  326. // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
  327. // CHECK: alloca i[[SZ]], align
  328. // CHECK: alloca i[[SZ]], align
  329. // CHECK: [[AA_CASTED:%.+]] = alloca i[[SZ]], align
  330. // CHECK: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align
  331. // CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
  332. // CHECK: [[AA:%.+]] = load i16, i16* [[AA_CADDR]], align
  333. // CHECK: [[AA_C:%.+]] = bitcast i[[SZ]]* [[AA_CASTED]] to i16*
  334. // CHECK: store i16 [[AA]], i16* [[AA_C]], align
  335. // CHECK: [[PARAM:%.+]] = load i[[SZ]], i[[SZ]]* [[AA_CASTED]], align
  336. // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 3, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]], i[[SZ]], i[[SZ]])* [[OMP_OUTLINED2:@.+]] to void (i32*, i32*, ...)*), i[[SZ]] [[PARAM]], i[[SZ]] {{.+}}, i[[SZ]] {{.+}})
  337. //
  338. //
  339. // CHECK: define internal {{.*}}void [[OMP_OUTLINED2]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}})
  340. // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
  341. // CHECK: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align
  342. // CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
  343. // CHECK: [[AA:%.+]] = load i16, i16* [[AA_CADDR]], align
  344. // CHECK: ret void
  345. // CHECK-NEXT: }
  346. // CHECK: define internal void [[HVT3]]
  347. // Create stack storage and store argument in there.
  348. // CHECK: [[A_ADDR:%.+]] = alloca i[[SZ]], align
  349. // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
  350. // CHECK: [[A_CASTED:%.+]] = alloca i[[SZ]], align
  351. // CHECK: [[AA_CASTED:%.+]] = alloca i[[SZ]], align
  352. // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[A_ADDR]], align
  353. // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align
  354. // CHECK-64-DAG:[[A_CADDR:%.+]] = bitcast i[[SZ]]* [[A_ADDR]] to i32*
  355. // CHECK-DAG: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
  356. // CHECK-64-DAG:[[A:%.+]] = load i32, i32* [[A_CADDR]], align
  357. // CHECK-32-DAG:[[A:%.+]] = load i32, i32* [[A_ADDR]], align
  358. // CHECK-64-DAG:[[A_C:%.+]] = bitcast i[[SZ]]* [[A_CASTED]] to i32*
  359. // CHECK-64-DAG:store i32 [[A]], i32* [[A_C]], align
  360. // CHECK-32-DAG:store i32 [[A]], i32* [[A_CASTED]], align
  361. // CHECK-DAG: [[AA:%.+]] = load i16, i16* [[AA_CADDR]], align
  362. // CHECK-DAG: [[AA_C:%.+]] = bitcast i[[SZ]]* [[AA_CASTED]] to i16*
  363. // CHECK-DAG: store i16 [[AA]], i16* [[AA_C]], align
  364. // CHECK-DAG: [[PARAM1:%.+]] = load i[[SZ]], i[[SZ]]* [[A_CASTED]], align
  365. // CHECK-DAG: [[PARAM2:%.+]] = load i[[SZ]], i[[SZ]]* [[AA_CASTED]], align
  366. // CHECK-DAG: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]], i[[SZ]])* [[OMP_OUTLINED3:@.+]] to void (i32*, i32*, ...)*), i[[SZ]] [[PARAM1]], i[[SZ]] [[PARAM2]])
  367. //
  368. //
  369. // CHECK: define internal {{.*}}void [[OMP_OUTLINED3]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}})
  370. // CHECK: [[A_ADDR:%.+]] = alloca i[[SZ]], align
  371. // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
  372. // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[A_ADDR]], align
  373. // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align
  374. // CHECK-64-DAG:[[A_CADDR:%.+]] = bitcast i[[SZ]]* [[A_ADDR]] to i32*
  375. // CHECK-DAG: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
  376. // CHECK: ret void
  377. // CHECK-NEXT: }
  378. // CHECK: define internal void [[HVT4]]
  379. // Create local storage for each capture.
  380. // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
  381. // CHECK: [[LOCAL_B:%.+]] = alloca [10 x float]*
  382. // CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
  383. // CHECK: [[LOCAL_BN:%.+]] = alloca float*
  384. // CHECK: [[LOCAL_C:%.+]] = alloca [5 x [10 x double]]*
  385. // CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
  386. // CHECK: [[LOCAL_VLA3:%.+]] = alloca i[[SZ]]
  387. // CHECK: [[LOCAL_CN:%.+]] = alloca double*
  388. // CHECK: [[LOCAL_D:%.+]] = alloca [[TT]]*
  389. // CHECK: alloca i[[SZ]]
  390. // CHECK: [[LOCAL_A_CASTED:%.+]] = alloca i[[SZ]]
  391. // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
  392. // CHECK-DAG: store [10 x float]* [[ARG_B:%.+]], [10 x float]** [[LOCAL_B]]
  393. // CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], i[[SZ]]* [[LOCAL_VLA1]]
  394. // CHECK-DAG: store float* [[ARG_BN:%.+]], float** [[LOCAL_BN]]
  395. // CHECK-DAG: store [5 x [10 x double]]* [[ARG_C:%.+]], [5 x [10 x double]]** [[LOCAL_C]]
  396. // CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], i[[SZ]]* [[LOCAL_VLA2]]
  397. // CHECK-DAG: store i[[SZ]] [[ARG_VLA3:%.+]], i[[SZ]]* [[LOCAL_VLA3]]
  398. // CHECK-DAG: store double* [[ARG_CN:%.+]], double** [[LOCAL_CN]]
  399. // CHECK-DAG: store [[TT]]* [[ARG_D:%.+]], [[TT]]** [[LOCAL_D]]
  400. // CHECK-64-DAG:[[CONV_AP:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
  401. // CHECK-DAG: [[REF_B:%.+]] = load [10 x float]*, [10 x float]** [[LOCAL_B]],
  402. // CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA1]],
  403. // CHECK-DAG: [[REF_BN:%.+]] = load float*, float** [[LOCAL_BN]],
  404. // CHECK-DAG: [[REF_C:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[LOCAL_C]],
  405. // CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA2]],
  406. // CHECK-DAG: [[VAL_VLA3:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA3]],
  407. // CHECK-DAG: [[REF_CN:%.+]] = load double*, double** [[LOCAL_CN]],
  408. // CHECK-DAG: [[REF_D:%.+]] = load [[TT]]*, [[TT]]** [[LOCAL_D]],
  409. // CHECK-64-DAG:[[CONV_A:%.+]] = load i32, i32* [[CONV_AP]]
  410. // CHECK-64-DAG:[[CONV:%.+]] = bitcast i[[SZ]]* [[LOCAL_A_CASTED]] to i32*
  411. // CHECK-64-DAG:store i32 [[CONV_A]], i32* [[CONV]], align
  412. // CHECK-32-DAG:[[LOCAL_AV:%.+]] = load i32, i32* [[LOCAL_A]]
  413. // CHECK-32-DAG:store i32 [[LOCAL_AV]], i32* [[LOCAL_A_CASTED]], align
  414. // CHECK-DAG: [[REF_A:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_A_CASTED]],
  415. // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 10, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]], [10 x float]*, i[[SZ]], float*, [5 x [10 x double]]*, i[[SZ]], i[[SZ]], double*, [[TT]]*, i[[SZ]])* [[OMP_OUTLINED4:@.+]] to void (i32*, i32*, ...)*), i[[SZ]] [[REF_A]], [10 x float]* [[REF_B]], i[[SZ]] [[VAL_VLA1]], float* [[REF_BN]], [5 x [10 x double]]* [[REF_C]], i[[SZ]] [[VAL_VLA2]], i[[SZ]] [[VAL_VLA3]], double* [[REF_CN]], [[TT]]* [[REF_D]], i[[SZ]] %{{.+}})
  416. //
  417. //
  418. // CHECK: define internal {{.*}}void [[OMP_OUTLINED4]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}}, [10 x float]* {{.+}}, i[[SZ]] %{{.+}}, float* {{.+}}, [5 x [10 x double]]* {{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, double* {{.+}}, [[TT]]* {{.+}})
  419. // To reduce complexity, we're only going as far as validating the signature of the outlined parallel function.
  420. template<typename tx>
  421. tx ftemplate(int n) {
  422. tx a = 0;
  423. short aa = 0;
  424. tx b[10];
  425. #pragma omp target parallel for if(target: n>40)
  426. for (long long i = -10; i < 10; i += 3) {
  427. a += 1;
  428. aa += 1;
  429. b[2] += 1;
  430. }
  431. return a;
  432. }
  433. static
  434. int fstatic(int n) {
  435. int a = 0;
  436. short aa = 0;
  437. char aaa = 0;
  438. int b[10];
  439. #pragma omp target parallel for if(target: n>50)
  440. for (unsigned i=100; i<10; i+=10) {
  441. a += 1;
  442. aa += 1;
  443. aaa += 1;
  444. b[2] += 1;
  445. }
  446. return a;
  447. }
  448. struct S1 {
  449. double a;
  450. int r1(int n){
  451. int b = n+1;
  452. short int c[2][n];
  453. #pragma omp target parallel for if(target: n>60)
  454. for (unsigned long long it = 2000; it >= 600; it -= 400) {
  455. this->a = (double)b + 1.5;
  456. c[1][1] = ++a;
  457. }
  458. return c[1][1] + (int)b;
  459. }
  460. };
  461. // CHECK: define {{.*}}@{{.*}}bar{{.*}}
  462. int bar(int n){
  463. int a = 0;
  464. // CHECK: call {{.*}}i32 [[FOO]](i32 {{.*}})
  465. a += foo(n);
  466. S1 S;
  467. // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}})
  468. a += S.r1(n);
  469. // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})
  470. a += fstatic(n);
  471. // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})
  472. a += ftemplate<int>(n);
  473. return a;
  474. }
  475. //
  476. // CHECK: define {{.*}}[[FS1]]
  477. //
  478. // CHECK: i8* @llvm.stacksave()
  479. // CHECK-64: [[B_ADDR:%.+]] = bitcast i[[SZ]]* [[B_CADDR:%.+]] to i32*
  480. // CHECK-64: store i32 %{{.+}}, i32* [[B_ADDR]],
  481. // CHECK-64: [[B_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[B_CADDR]],
  482. // CHECK-32: store i32 %{{.+}}, i32* %__vla_expr
  483. // CHECK-32: store i32 %{{.+}}, i32* [[B_ADDR:%.+]],
  484. // CHECK-32: [[B_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[B_ADDR]],
  485. // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60
  486. // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]
  487. // CHECK: [[TRY]]
  488. // We capture 2 VLA sizes in this target region
  489. // CHECK: [[CELEMSIZE2:%.+]] = mul nuw i[[SZ]] 2, [[VLA0:%.+]]
  490. // CHECK-64: [[CSIZE:%.+]] = mul nuw i64 [[CELEMSIZE2]], 2
  491. // CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2
  492. // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64
  493. // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT7]], i32 0, i32 0), i32 1, i32 0)
  494. // CHECK-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0
  495. // CHECK-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0
  496. // CHECK-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0
  497. // CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX0:[0-9]+]]
  498. // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]]
  499. // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]]
  500. // CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX1:[0-9]+]]
  501. // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]]
  502. // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]]
  503. // CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX2:[0-9]+]]
  504. // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]]
  505. // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]]
  506. // CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]]
  507. // CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]]
  508. // CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]]
  509. // CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX4:[0-9]+]]
  510. // CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX4]]
  511. // CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX4]]
  512. // CHECK-DAG: [[SADDR5:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX5:[0-9]+]]
  513. // CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX5]]
  514. // CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX5]]
  515. // The names below are not necessarily consistent with the names used for the
  516. // addresses above as some are repeated.
  517. // CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CBPADDR0:%.+]],
  518. // CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CPADDR0:%.+]],
  519. // CHECK-DAG: [[CBPADDR0]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  520. // CHECK-DAG: [[CPADDR0]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  521. // CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}}
  522. // CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CBPADDR1:%.+]],
  523. // CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CPADDR1:%.+]],
  524. // CHECK-DAG: [[CBPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  525. // CHECK-DAG: [[CPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  526. // CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}}
  527. // CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CBPADDR2:%.+]],
  528. // CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CPADDR2:%.+]],
  529. // CHECK-DAG: [[CBPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  530. // CHECK-DAG: [[CPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]*
  531. // CHECK-DAG: store i64 4, i64* {{%[^,]+}}
  532. // CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR3:%.+]],
  533. // CHECK-DAG: store double* %{{.+}}, double** [[CPADDR3:%.+]],
  534. // CHECK-DAG: [[CBPADDR3]] = bitcast i8** {{%[^,]+}} to [[S1]]**
  535. // CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to double**
  536. // CHECK-DAG: store i64 8, i64* {{%[^,]+}}
  537. // CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR4:%.+]],
  538. // CHECK-DAG: store double* %{{.+}}, double** [[CPADDR4:%.+]],
  539. // CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to [[S1]]**
  540. // CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to double**
  541. // CHECK-DAG: store i64 8, i64* {{%[^,]+}}
  542. // CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR5:%.+]],
  543. // CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR5:%.+]],
  544. // CHECK-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to i16**
  545. // CHECK-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to i16**
  546. // CHECK-DAG: store i64 [[CSIZE]], i64* {{%[^,]+}}
  547. // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
  548. // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
  549. // CHECK: [[FAIL]]
  550. // CHECK: call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
  551. // CHECK-NEXT: br label %[[END]]
  552. // CHECK: [[END]]
  553. //
  554. // CHECK: define {{.*}}[[FSTATIC]]
  555. //
  556. // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 50
  557. // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
  558. // CHECK: [[IFTHEN]]
  559. // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* getelementptr inbounds ([4 x i64], [4 x i64]* [[SIZET6]], i32 0, i32 0), i64* getelementptr inbounds ([4 x i64], [4 x i64]* [[MAPT6]], i32 0, i32 0), i32 1, i32 0)
  560. // CHECK-DAG: [[BPR]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP:%.+]], i32 0, i32 0
  561. // CHECK-DAG: [[PR]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P:%.+]], i32 0, i32 0
  562. // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 0
  563. // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 0
  564. // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]*
  565. // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]*
  566. // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], i[[SZ]]* [[CBPADDR0]],
  567. // CHECK-DAG: store i[[SZ]] [[VAL0]], i[[SZ]]* [[CPADDR0]],
  568. // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 1
  569. // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 1
  570. // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*
  571. // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*
  572. // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], i[[SZ]]* [[CBPADDR1]],
  573. // CHECK-DAG: store i[[SZ]] [[VAL1]], i[[SZ]]* [[CPADDR1]],
  574. // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 2
  575. // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 2
  576. // CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to i[[SZ]]*
  577. // CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to i[[SZ]]*
  578. // CHECK-DAG: store i[[SZ]] [[VAL2:%.+]], i[[SZ]]* [[CBPADDR2]],
  579. // CHECK-DAG: store i[[SZ]] [[VAL2]], i[[SZ]]* [[CPADDR2]],
  580. // CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 3
  581. // CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 3
  582. // CHECK-DAG: [[CBPADDR3:%.+]] = bitcast i8** [[BPADDR3]] to [10 x i32]**
  583. // CHECK-DAG: [[CPADDR3:%.+]] = bitcast i8** [[PADDR3]] to [10 x i32]**
  584. // CHECK-DAG: store [10 x i32]* [[VAL3:%.+]], [10 x i32]** [[CBPADDR3]],
  585. // CHECK-DAG: store [10 x i32]* [[VAL3]], [10 x i32]** [[CPADDR3]],
  586. // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
  587. // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
  588. // CHECK: [[FAIL]]
  589. // CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
  590. // CHECK-NEXT: br label %[[END]]
  591. // CHECK: [[END]]
  592. // CHECK-NEXT: br label %[[IFEND:.+]]
  593. // CHECK: [[IFELSE]]
  594. // CHECK: call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
  595. // CHECK-NEXT: br label %[[IFEND]]
  596. // CHECK: [[IFEND]]
  597. //
  598. // CHECK: define {{.*}}[[FTEMPLATE]]
  599. //
  600. // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 40
  601. // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
  602. // CHECK: [[IFTHEN]]
  603. // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZET5]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT5]], i32 0, i32 0), i32 1, i32 0)
  604. // CHECK-DAG: [[BPR]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP:%.+]], i32 0, i32 0
  605. // CHECK-DAG: [[PR]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P:%.+]], i32 0, i32 0
  606. // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 0
  607. // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 0
  608. // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]*
  609. // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]*
  610. // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], i[[SZ]]* [[CBPADDR0]],
  611. // CHECK-DAG: store i[[SZ]] [[VAL0]], i[[SZ]]* [[CPADDR0]],
  612. // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 1
  613. // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 1
  614. // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*
  615. // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*
  616. // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], i[[SZ]]* [[CBPADDR1]],
  617. // CHECK-DAG: store i[[SZ]] [[VAL1]], i[[SZ]]* [[CPADDR1]],
  618. // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 2
  619. // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 2
  620. // CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to [10 x i32]**
  621. // CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to [10 x i32]**
  622. // CHECK-DAG: store [10 x i32]* [[VAL2:%.+]], [10 x i32]** [[CBPADDR2]],
  623. // CHECK-DAG: store [10 x i32]* [[VAL2]], [10 x i32]** [[CPADDR2]],
  624. // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
  625. // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
  626. // CHECK: [[FAIL]]
  627. // CHECK: call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
  628. // CHECK-NEXT: br label %[[END]]
  629. // CHECK: [[END]]
  630. // CHECK-NEXT: br label %[[IFEND:.+]]
  631. // CHECK: [[IFELSE]]
  632. // CHECK: call void [[HVT:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
  633. // CHECK-NEXT: br label %[[IFEND]]
  634. // CHECK: [[IFEND]]
  635. // Check that the offloading functions are emitted and that the arguments are
  636. // correct and loaded correctly for the target regions of the callees of bar().
  637. // CHECK: define internal void [[HVT7]]
  638. // Create local storage for each capture.
  639. // CHECK: [[LOCAL_THIS:%.+]] = alloca [[S1]]*
  640. // CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]]
  641. // CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
  642. // CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
  643. // CHECK: [[LOCAL_C:%.+]] = alloca i16*
  644. // CHECK: [[LOCAL_B_CASTED:%.+]] = alloca i[[SZ]]
  645. // CHECK-DAG: store [[S1]]* [[ARG_THIS:%.+]], [[S1]]** [[LOCAL_THIS]]
  646. // CHECK-DAG: store i[[SZ]] [[ARG_B:%.+]], i[[SZ]]* [[LOCAL_B]]
  647. // CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], i[[SZ]]* [[LOCAL_VLA1]]
  648. // CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], i[[SZ]]* [[LOCAL_VLA2]]
  649. // CHECK-DAG: store i16* [[ARG_C:%.+]], i16** [[LOCAL_C]]
  650. // Store captures in the context.
  651. // CHECK-DAG: [[REF_THIS:%.+]] = load [[S1]]*, [[S1]]** [[LOCAL_THIS]],
  652. // CHECK-64-DAG:[[CONV_BP:%.+]] = bitcast i[[SZ]]* [[LOCAL_B]] to i32*
  653. // CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA1]],
  654. // CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA2]],
  655. // CHECK-DAG: [[REF_C:%.+]] = load i16*, i16** [[LOCAL_C]],
  656. // CHECK-64-DAG:[[CONV_B:%.+]] = load i32, i32* [[CONV_BP]]
  657. // CHECK-64-DAG:[[CONV:%.+]] = bitcast i[[SZ]]* [[LOCAL_B_CASTED]] to i32*
  658. // CHECK-64-DAG:store i32 [[CONV_B]], i32* [[CONV]], align
  659. // CHECK-32-DAG:[[LOCAL_BV:%.+]] = load i32, i32* [[LOCAL_B]]
  660. // CHECK-32-DAG:store i32 [[LOCAL_BV]], i32* [[LOCAL_B_CASTED]], align
  661. // CHECK-DAG: [[REF_B:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_B_CASTED]],
  662. // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 5, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, [[S1]]*, i[[SZ]], i[[SZ]], i[[SZ]], i16*)* [[OMP_OUTLINED5:@.+]] to void (i32*, i32*, ...)*), [[S1]]* [[REF_THIS]], i[[SZ]] [[REF_B]], i[[SZ]] [[VAL_VLA1]], i[[SZ]] [[VAL_VLA2]], i16* [[REF_C]])
  663. //
  664. //
  665. // CHECK: define internal {{.*}}void [[OMP_OUTLINED5]](i32* noalias %.global_tid., i32* noalias %.bound_tid., [[S1]]* %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i16* {{.+}})
  666. // To reduce complexity, we're only going as far as validating the signature of the outlined parallel function.
  667. // CHECK: define internal void [[HVT6]]
  668. // Create local storage for each capture.
  669. // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
  670. // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
  671. // CHECK: [[LOCAL_AAA:%.+]] = alloca i[[SZ]]
  672. // CHECK: [[LOCAL_B:%.+]] = alloca [10 x i32]*
  673. // CHECK: [[LOCAL_A_CASTED:%.+]] = alloca i[[SZ]]
  674. // CHECK: [[LOCAL_AA_CASTED:%.+]] = alloca i[[SZ]]
  675. // CHECK: [[LOCAL_AAA_CASTED:%.+]] = alloca i[[SZ]]
  676. // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
  677. // CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]
  678. // CHECK-DAG: store i[[SZ]] [[ARG_AAA:%.+]], i[[SZ]]* [[LOCAL_AAA]]
  679. // CHECK-DAG: store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
  680. // Store captures in the context.
  681. // CHECK-64-DAG:[[CONV_AP:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
  682. // CHECK-DAG: [[CONV_AAP:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16*
  683. // CHECK-DAG: [[CONV_AAAP:%.+]] = bitcast i[[SZ]]* [[LOCAL_AAA]] to i8*
  684. // CHECK-DAG: [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
  685. // CHECK-64-DAG:[[CONV_A:%.+]] = load i32, i32* [[CONV_AP]]
  686. // CHECK-64-DAG:[[CONV:%.+]] = bitcast i[[SZ]]* [[LOCAL_A_CASTED]] to i32*
  687. // CHECK-64-DAG:store i32 [[CONV_A]], i32* [[CONV]], align
  688. // CHECK-32-DAG:[[LOCAL_AV:%.+]] = load i32, i32* [[LOCAL_A]]
  689. // CHECK-32-DAG:store i32 [[LOCAL_AV]], i32* [[LOCAL_A_CASTED]], align
  690. // CHECK-DAG: [[REF_A:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_A_CASTED]],
  691. // CHECK-DAG: [[CONV_AA:%.+]] = load i16, i16* [[CONV_AAP]]
  692. // CHECK-DAG: [[CONV:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA_CASTED]] to i16*
  693. // CHECK-DAG: store i16 [[CONV_AA]], i16* [[CONV]], align
  694. // CHECK-DAG: [[REF_AA:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_AA_CASTED]],
  695. // CHECK-DAG: [[CONV_AAA:%.+]] = load i8, i8* [[CONV_AAAP]]
  696. // CHECK-DAG: [[CONV:%.+]] = bitcast i[[SZ]]* [[LOCAL_AAA_CASTED]] to i8*
  697. // CHECK-DAG: store i8 [[CONV_AAA]], i8* [[CONV]], align
  698. // CHECK-DAG: [[REF_AAA:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_AAA_CASTED]],
  699. // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 4, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]], i[[SZ]], i[[SZ]], [10 x i32]*)* [[OMP_OUTLINED6:@.+]] to void (i32*, i32*, ...)*), i[[SZ]] [[REF_A]], i[[SZ]] [[REF_AA]], i[[SZ]] [[REF_AAA]], [10 x i32]* [[REF_B]])
  700. //
  701. //
  702. // CHECK: define internal {{.*}}void [[OMP_OUTLINED6]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, [10 x i32]* {{.+}})
  703. // To reduce complexity, we're only going as far as validating the signature of the outlined parallel function.
  704. // CHECK: define internal void [[HVT5]]
  705. // Create local storage for each capture.
  706. // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
  707. // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
  708. // CHECK: [[LOCAL_B:%.+]] = alloca [10 x i32]*
  709. // CHECK: [[LOCAL_A_CASTED:%.+]] = alloca i[[SZ]]
  710. // CHECK: [[LOCAL_AA_CASTED:%.+]] = alloca i[[SZ]]
  711. // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
  712. // CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]
  713. // CHECK-DAG: store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
  714. // Store captures in the context.
  715. // CHECK-64-DAG:[[CONV_AP:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
  716. // CHECK-DAG: [[CONV_AAP:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16*
  717. // CHECK-DAG: [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
  718. // CHECK-64-DAG:[[CONV_A:%.+]] = load i32, i32* [[CONV_AP]]
  719. // CHECK-64-DAG:[[CONV:%.+]] = bitcast i[[SZ]]* [[LOCAL_A_CASTED]] to i32*
  720. // CHECK-64-DAG:store i32 [[CONV_A]], i32* [[CONV]], align
  721. // CHECK-32-DAG:[[LOCAL_AV:%.+]] = load i32, i32* [[LOCAL_A]]
  722. // CHECK-32-DAG:store i32 [[LOCAL_AV]], i32* [[LOCAL_A_CASTED]], align
  723. // CHECK-DAG: [[REF_A:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_A_CASTED]],
  724. // CHECK-DAG: [[CONV_AA:%.+]] = load i16, i16* [[CONV_AAP]]
  725. // CHECK-DAG: [[CONV:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA_CASTED]] to i16*
  726. // CHECK-DAG: store i16 [[CONV_AA]], i16* [[CONV]], align
  727. // CHECK-DAG: [[REF_AA:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_AA_CASTED]],
  728. // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 3, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]], i[[SZ]], [10 x i32]*)* [[OMP_OUTLINED7:@.+]] to void (i32*, i32*, ...)*), i[[SZ]] [[REF_A]], i[[SZ]] [[REF_AA]], [10 x i32]* [[REF_B]])
  729. //
  730. //
  731. // CHECK: define internal {{.*}}void [[OMP_OUTLINED7]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, [10 x i32]* {{.+}})
  732. // To reduce complexity, we're only going as far as validating the signature of the outlined parallel function.
  733. // CHECK: define internal void @.omp_offloading.requires_reg()
  734. // CHECK: call void @__tgt_register_requires(i64 1)
  735. // CHECK: ret void
  736. #endif