================
@@ -0,0 +1,81 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux-gnu \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-host.bc \
+// RUN: -fopenmp-assume-teams-oversubscription \
+// RUN: -fopenmp-assume-threads-oversubscription \
+// RUN: -emit-llvm %s -o - | FileCheck %s \
+// RUN: --check-prefixes=NOLOOP
+
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-host.bc \
+// RUN: -emit-llvm %s -o - | FileCheck %s \
+// RUN: --check-prefix=SPMD
--implicit-check-not=__kmpc_distribute_for_static_loop_4u
+
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-host.bc \
+// RUN: -fopenmp-assume-teams-oversubscription \
+// RUN: -emit-llvm %s -o - | FileCheck %s \
+// RUN: --check-prefix=SPMD
--implicit-check-not=__kmpc_distribute_for_static_loop_4u
+
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgcn-amd-amdhsa \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-host.bc \
+// RUN: -fopenmp-assume-threads-oversubscription \
+// RUN: -emit-llvm %s -o - | FileCheck %s \
+// RUN: --check-prefix=SPMD
--implicit-check-not=__kmpc_distribute_for_static_loop_4u
+
+// expected-no-diagnostics
+
+void no_loop(int *array) {
+#pragma omp target teams distribute parallel for
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void no_loop_simd(int *array) {
+#pragma omp target teams distribute parallel for simd
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void no_loop_nowait(int *array) {
+#pragma omp target teams distribute parallel for nowait
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+// NOLOOP: no_loop_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP: no_loop_simd_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_l{{[0-9]+}}{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: ret void
+// NOLOOP:
@__kmpc_distribute_for_static_loop_4u({{.*}}no_loop_l{{[0-9]+}}{{.*}}, i32 0,
i32 0, i8 1)
+// NOLOOP: omp_loop.after:
+// NOLOOP-NEXT: ret void
+
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_simd{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: store i32 1024, ptr %i
----------------
nicebert wrote:
You're right, the store isn't needed in this test as written, but I don't think
dropping the call to EmitOMPSimdFinal on the NoLoop path is the right fix, for
two reasons.
The regular (non no-loop) path emits the same unnecessary store for this loop,
guarded by the last-iteration check.
It would lead to a missing store for a case where the counter is declared
outside the loop and mapped, e.g.
```
int i = -1;
#pragma omp target teams distribute parallel for simd map(tofrom: i)
for (i = 0; i < 1024; ++i) ...
```
which means the host would never see `i` reach 1024.
I could change the simd test to the mapped version above, so the store it
checks is one that's actually needed, if you'd prefer that.
A possible solution would be to only emit the store when the counter outlives
the loop but that would change every simd loop, not just this one, so I'd
rather keep that out of this PR.
https://github.com/llvm/llvm-project/pull/224041
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
