animeshk-amd updated this revision to Diff 455161.
animeshk-amd added a comment.
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
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D131763/new/
https://reviews.llvm.org/D131763
Files:
clang/test/OpenMP/metadirective_ast_print.c
clang/test/OpenMP/metadirective_device_arch_codegen.cpp
Index: clang/test/OpenMP/metadirective_device_arch_codegen.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/metadirective_device_arch_codegen.cpp
@@ -0,0 +1,65 @@
+// 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
+
Index: clang/test/OpenMP/metadirective_ast_print.c
===================================================================
--- clang/test/OpenMP/metadirective_ast_print.c
+++ clang/test/OpenMP/metadirective_ast_print.c
@@ -1,6 +1,10 @@
// 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
@@ -57,6 +61,12 @@
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);
@@ -83,5 +93,7 @@
// 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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
