[clang] [llvm] [clang][OpenMP] Add no-loop SPMD kernel promotion (PR #224041)

[Nicole Aschenbrenner via cfe-commits]

================
@@ -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