[OpenMP] Add lit test for metadirective device arch inspired from sollve

[animeshk-amd]

Summary

This lit test is added based upon the tests present in the
tests/5.0/metadirective directory of the SOLLVE repo
https://github.com/SOLLVE/sollve_vv

[jdoerfert]

This doesn't actually test much, only once case/compilation is covered. In the second function nothing specific to LLVM as impl is checked.

[saiislam]

In D131763#3719132, @jdoerfert wrote:

This doesn't actually test much, only once case/compilation is covered. In the second function nothing specific to LLVM as impl is checked.

The second function, is the only place in llvm-project where vendor(llvm) is being tested for a non-error test.

[jdoerfert]

In D131763#3719140, @saiislam wrote:

In D131763#3719132, @jdoerfert wrote:

This doesn't actually test much, only once case/compilation is covered. In the second function nothing specific to LLVM as impl is checked.

The second function, is the only place in llvm-project where vendor(llvm) is being tested for a non-error test.

Really?

ag 'vendor\(llvm\)' clang/test/OpenMP --files-with-matches                                                                                                                                                                

clang/test/OpenMP/begin_declare_variant_messages.c
clang/test/OpenMP/begin_declare_variant_using_messages.cpp
clang/test/OpenMP/declare_variant_ast_print.c
clang/test/OpenMP/declare_variant_ast_print.cpp
clang/test/OpenMP/declare_variant_implementation_vendor_codegen.cpp
clang/test/OpenMP/declare_variant_messages.cpp
clang/test/OpenMP/declare_variant_mixed_codegen.cpp
clang/test/OpenMP/metadirective_ast_print.c
clang/test/OpenMP/metadirective_implementation_codegen.c
clang/test/OpenMP/nvptx_declare_variant_implementation_vendor_codegen.cpp
clang/test/OpenMP/declare_variant_messages.c
clang/test/OpenMP/metadirective_empty.cpp
clang/test/OpenMP/metadirective_implementation_codegen.cpp

That said, the above function still doesn't test anything wrt. llvm as impl anyway. We could just as well match amd or nvidia and the check lines still match just fine.

[saiislam]

In D131763#3719323, @jdoerfert wrote:

In D131763#3719140, @saiislam wrote:

In D131763#3719132, @jdoerfert wrote:

This doesn't actually test much, only once case/compilation is covered. In the second function nothing specific to LLVM as impl is checked.

The second function, is the only place in llvm-project where vendor(llvm) is being tested for a non-error test.

Really?

ag 'vendor\(llvm\)' clang/test/OpenMP --files-with-matches                                                                                                                                                                

clang/test/OpenMP/begin_declare_variant_messages.c
clang/test/OpenMP/begin_declare_variant_using_messages.cpp
clang/test/OpenMP/declare_variant_ast_print.c
clang/test/OpenMP/declare_variant_ast_print.cpp
clang/test/OpenMP/declare_variant_implementation_vendor_codegen.cpp
clang/test/OpenMP/declare_variant_messages.cpp
clang/test/OpenMP/declare_variant_mixed_codegen.cpp
clang/test/OpenMP/metadirective_ast_print.c
clang/test/OpenMP/metadirective_implementation_codegen.c
clang/test/OpenMP/nvptx_declare_variant_implementation_vendor_codegen.cpp
clang/test/OpenMP/declare_variant_messages.c
clang/test/OpenMP/metadirective_empty.cpp
clang/test/OpenMP/metadirective_implementation_codegen.cpp

That said, the above function still doesn't test anything wrt. llvm as impl anyway. We could just as well match amd or nvidia and the check lines still match just fine.

My mistake. You are right about more tests being present. I searched in the wrong branch locally :-)

[animeshk-amd]

The second function was testing the implementation trait set for
the vendor trait with respect to the metadirective directive.
After Johannes's comment, on further searching it turns out that
this trait and trait set are already being tested in other tests.
Therfore second function need not be there. Also, the ast_print
test for the first function is added.

Updating D131763: [OpenMP] Add lit test for metadirective device arch inspired

from sollve

[saiislam]

LGTM. Thanks!

[animeshk-amd]

Thanks @saiislam. Since I don't have the commit access yet, I request you to kindly commit this on my behalf.

clang/test/OpenMP/metadirective_ast_print.c

 // RUN: %clang_cc1 -verify -fopenmp -triple x86_64-unknown-linux-gnu -x c -std=c99 -ast-print %s -o - | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -triple x86_64-unknown-linux-gnu -x c -std=c99 -ast-print %s -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=CHECK-AMDGCN
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=CHECK-AMDGCN
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 void bar(void);
 
 #define N 10
// Test metadirective with nested OpenMP directive.
   int array[16];
   #pragma omp metadirective when(user = {condition(1)} \
                                  : parallel for)
   for (int i = 0; i < 16; i++) {
     #pragma omp simd
     for (int j = 0; j < 16; j++)
       array[i] = i;
   }
+
+#pragma omp metadirective when(device={arch("amdgcn")}: \
+                                teams distribute parallel for)\
+                                default(parallel for)
+  for (int i = 0; i < 100; i++)
+  ;
 }
 
 // CHECK: void bar(void);
 // CHECK: void foo(void)
 // CHECK-NEXT: #pragma omp parallel
 // CHECK-NEXT: bar()
 // CHECK-NEXT: #pragma omp parallel
 // CHECK-NEXT: bar()

 // CHECK: #pragma omp parallel for
 // CHECK-NEXT: for (int i = 0; i < 100; i++)
 // CHECK: #pragma omp parallel
 // CHECK-NEXT: for (int i = 0; i < 100; i++)
 // CHECK: #pragma omp parallel for
 // CHECK-NEXT: for (int i = 0; i < 16; i++) {
 // CHECK-NEXT: #pragma omp simd
 // CHECK-NEXT: for (int j = 0; j < 16; j++)
+// CHECK-AMDGCN: #pragma omp teams distribute parallel for
+// CHECK-AMDGCN-NEXT: for (int i = 0; i < 100; i++)
 
 #endif

clang/test/OpenMP/metadirective_device_arch_codegen.cpp

+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -target-cpu gfx906 -o - | FileCheck %s
+// expected-no-diagnostics
+
+
+/*===-----------------------------------------------------------------------=== 
+
+Inspired from SOLLVE tests:
+ - 5.0/metadirective/test_metadirective_arch_is_nvidia.c
+
+
+===------------------------------------------------------------------------===*/
+
+
+#define N 1024
+
+int metadirective1() {
+   
+   int v1[N], v2[N], v3[N];
+
+   int target_device_num, host_device_num, default_device;
+   int errors = 0;
+
+   #pragma omp target map(to:v1,v2) map(from:v3, target_device_num) device(default_device)
+   {
+      #pragma omp metadirective \
+                   when(device={arch("amdgcn")}: teams distribute parallel for) \
+                   default(parallel for)
+
+         for (int i = 0; i < N; i++) {
+      #pragma omp atomic write
+            v3[i] = v1[i] * v2[i];
+         }
+   }
+
+   return errors;
+}
+
+// CHECK-LABEL: define weak_odr amdgpu_kernel void {{.+}}metadirective1
+// CHECK: entry:
+// CHECK: %{{[0-9]}} = call i32 @__kmpc_target_init
+// CHECK: user_code.entry:
+// CHECK: call void @__omp_outlined__
+// CHECK-NOT: call void @__kmpc_parallel_51
+// CHECK: ret void
+
+
+// CHECK-LABEL: define internal void @__omp_outlined__
+// CHECK: entry:
+// CHECK: call void @__kmpc_distribute_static_init
+// CHECK: omp.loop.exit:  
+// CHECK: call void @__kmpc_distribute_static_fini
+
+
+// CHECK-LABEL: define internal void @__omp_outlined__.{{[0-9]+}}
+// CHECK: entry:
+// CHECK: call void @__kmpc_for_static_init_4
+// CHECK: omp.inner.for.body:
+// CHECK: store atomic {{.*}} monotonic
+// CHECK: omp.loop.exit:                                    
+// CHECK-NEXT: call void @__kmpc_distribute_static_fini
+// CHECK-NEXT: ret void
+