saiislam updated this revision to Diff 397912.
saiislam added a comment.
Herald added a subscriber: jvesely.
Added target specific tests for ISA traits, for CPU as well as GPU.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D116549/new/
https://reviews.llvm.org/D116549
Files:
clang/lib/Parse/ParseOpenMP.cpp
clang/test/OpenMP/amdgcn_target_codegen.cpp
clang/test/OpenMP/metadirective_implementation_codegen.c
clang/test/OpenMP/metadirective_implementation_codegen.cpp
Index: clang/test/OpenMP/metadirective_implementation_codegen.cpp
===================================================================
--- clang/test/OpenMP/metadirective_implementation_codegen.cpp
+++ clang/test/OpenMP/metadirective_implementation_codegen.cpp
@@ -1,6 +1,7 @@
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -target-cpu x86-64| FileCheck %s -check-prefixes=SUPPORTEDISA,UNSUPPORTEDISA
// expected-no-diagnostics
#ifndef HEADER
@@ -35,6 +36,12 @@
: parallel) default(parallel for)
for (int i = 0; i < 100; i++)
;
+#pragma omp metadirective when(device = {isa("sse2")} \
+ : parallel) default(single)
+ bar();
+#pragma omp metadirective when(device = {isa("some-unsupported-feature")} \
+ : parallel) default(single)
+ bar();
}
// CHECK-LABEL: void @_Z3foov()
@@ -44,6 +51,10 @@
// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_5:@.+]] to void
// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_6:@.+]] to void
// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_7:@.+]] to void
+// SUPPORTEDISA: ...) @__kmpc_fork_call(
+// UNSUPPORTEDISA: call i32 @__kmpc_single
+// UNSUPPORTEDISA: @_Z3barv
+// UNSUPPORTEDISA: call void @__kmpc_end_single
// CHECK: ret void
// CHECK: define internal void [[OUTLINED_2]](
@@ -73,4 +84,7 @@
// NO-CHECK: call void @__kmpc_for_static_fini
// CHECK: ret void
+// SUPPORTEDISA: define internal void @.omp_outlined..6(
+// SUPPORTEDISA: @_Z3barv
+// SUPPORTEDISA: ret void
#endif
Index: clang/test/OpenMP/metadirective_implementation_codegen.c
===================================================================
--- clang/test/OpenMP/metadirective_implementation_codegen.c
+++ clang/test/OpenMP/metadirective_implementation_codegen.c
@@ -1,6 +1,7 @@
// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c -triple aarch64-unknown-linux -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c -triple ppc64le-unknown-linux -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm -target-cpu x86-64 %s -o - | FileCheck %s -check-prefixes=SUPPORTEDISA,UNSUPPORTEDISA
// expected-no-diagnostics
#ifndef HEADER
@@ -35,10 +36,20 @@
: parallel) default(parallel for)
for (int i = 0; i < 100; i++)
;
+#pragma omp metadirective when(device = {isa("sse2")} \
+ : parallel) default(single)
+ bar();
+#pragma omp metadirective when(device = {isa("some-unsupported-feature")} \
+ : parallel) default(single)
+ bar();
}
// CHECK: void @foo()
// CHECK-COUNT-6: ...) @__kmpc_fork_call(
+// SUPPORTEDISA: ...) @__kmpc_fork_call(
+// UNSUPPORTEDISA: call i32 @__kmpc_single
+// UNSUPPORTEDISA: @bar
+// UNSUPPORTEDISA: call void @__kmpc_end_single
// CHECK: ret void
// CHECK: define internal void @.omp_outlined.(
@@ -68,4 +79,7 @@
// NO-CHECK: call void @__kmpc_for_static_fini
// CHECK: ret void
+// SUPPORTEDISA: define internal void @.omp_outlined..6(
+// SUPPORTEDISA: @bar
+// SUPPORTEDISA: ret void
#endif
Index: clang/test/OpenMP/amdgcn_target_codegen.cpp
===================================================================
--- clang/test/OpenMP/amdgcn_target_codegen.cpp
+++ clang/test/OpenMP/amdgcn_target_codegen.cpp
@@ -2,6 +2,7 @@
// RUN: %clang_cc1 -fopenmp -x c++ -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++ -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 -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -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 -check-prefix=AMDGPUTARGET
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
@@ -35,4 +36,46 @@
return arr[0];
}
+// AMDGPUTARGET: define weak amdgpu_kernel void @__omp_offloading_{{.*}}test_amdgcn_device_isa_selected
+// AMDGPUTARGET: user_code.entry:
+// AMDGPUTARGET: call void @__kmpc_parallel_51
+// AMDGPUTARGET-NOT: call i32 @__kmpc_single
+// AMDGPUTARGET: ret void
+int test_amdgcn_device_isa_selected() {
+ int threadCount = 0;
+
+#pragma omp target map(tofrom \
+ : threadCount)
+ {
+#pragma omp metadirective \
+ when(device = {isa("flat-address-space")} \
+ : parallel) default(single)
+ threadCount++;
+ }
+
+ return threadCount;
+}
+
+// AMDGPUTARGET: define weak amdgpu_kernel void @__omp_offloading_{{.*}}test_amdgcn_device_isa_not_selected
+// AMDGPUTARGET: user_code.entry:
+// AMDGPUTARGET: call i32 @__kmpc_single
+// AMDGPUTARGET-NOT: call void @__kmpc_parallel_51
+// AMDGPUTARGET: ret void
+int test_amdgcn_device_isa_not_selected() {
+ int threadCount = 0;
+
+#pragma omp target map(tofrom \
+ : threadCount)
+ {
+#pragma omp metadirective \
+ when(device = {isa("sse")} \
+ : parallel) \
+ when(device = {isa("another-unsupported-gpu-feature")} \
+ : parallel) default(single)
+ threadCount++;
+ }
+
+ return threadCount;
+}
+
#endif
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ clang/lib/Parse/ParseOpenMP.cpp
@@ -2529,7 +2529,9 @@
TPA.Revert();
// End of the first iteration. Parser is reset to the start of metadirective
- TargetOMPContext OMPCtx(ASTContext, /* DiagUnknownTrait */ nullptr,
+ std::function<void(StringRef)> DiagUnknownTrait =
+ [this, Loc](StringRef ISATrait) {};
+ TargetOMPContext OMPCtx(ASTContext, std::move(DiagUnknownTrait),
/* CurrentFunctionDecl */ nullptr,
ArrayRef<llvm::omp::TraitProperty>());
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
