[OpenMP][Clang] Allow ancestor device modifier only with reverse offloading

[saiislam]

Summary

OpenMP Spec 5.0 [2.12.5, Restrictions]: If a device clause in which the
ancestor device-modifier appears is present on the target construct,
then a requires directive with the reverse_offload clause must be
specified.

[ABataev]

clang/test/OpenMP/target_device_codegen.cpp

43–47

Do we have a codegen test for the ancestor modifier?

[saiislam]

clang/test/OpenMP/target_device_codegen.cpp

43–47

We shouldn't have a test for ancestor modifier because it can only be used when requires reverse_offload is specified and Spec 5.2 says that if an implementation is not supporting a requirement (reverse offload in this case) then it should give compile-time error termination [1].

I am going to propose this change in a different phab review.

[1] OpenMP API Specification 5.2, Section 8.2.1, lines 12-13, page 212.

[ABataev]

clang/test/OpenMP/target_device_codegen.cpp

43–47

I'm not asking about runtime test but the unit test. We should check at least that using this modifier does not crash the compiler. Could you restore this test here and add #pragma omp requires reverse_offload instead?

[saiislam]

Restored device ancestor codegen unit test with requires reverese_offload.

[ABataev]

LG

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_requires_clause_redeclaration : Error <
   "Only one %0 clause can appear on a requires directive in a single translation unit">;
 def note_omp_requires_previous_clause : Note <
   "%0 clause previously used here">;
 def err_omp_directive_before_requires : Error <
   "'%0' region encountered before requires directive with '%1' clause">;
 def note_omp_requires_encountered_directive : Note <
   "'%0' previously encountered here">;
+def err_omp_device_ancestor_without_requires_reverse_offload : Error <
+  "Device clause with ancestor device-modifier used without specifying 'requires reverse_offload'">;
 def err_omp_invalid_scope : Error <
   "'#pragma omp %0' directive must appear only in file scope">;
 def note_omp_invalid_length_on_this_ptr_mapping : Note <
   "expected length on mapping of 'this' array section expression to be '1'">;
 def note_omp_invalid_lower_bound_on_this_ptr_mapping : Note <
   "expected lower bound on mapping of 'this' array section expression to be '0' or not specified">;
 def note_omp_invalid_subscript_on_this_ptr_map : Note <
   "expected 'this' subscript expression on map clause to be 'this[0]'">;

clang/lib/Sema/SemaOpenMP.cpp

OMPClause *Sema::ActOnOpenMPDeviceClause(OpenMPDeviceClauseModifier Modifier,
   // OpenMP [2.9.1, Restrictions]
   // The device expression must evaluate to a non-negative integer value.
   ErrorFound = !isNonNegativeIntegerValue(ValExpr, *this, OMPC_device,
                                           /*StrictlyPositive=*/false) ||
                ErrorFound;
   if (ErrorFound)
     return nullptr;
 
+  // OpenMP 5.0 [2.12.5, Restrictions]
+  // In case of ancestor device-modifier, a requires directive with
+  // the reverse_offload clause must be specified.
+  if (Modifier == OMPC_DEVICE_ancestor) {
+    if (!DSAStack->hasRequiresDeclWithClause<OMPReverseOffloadClause>()) {
+      targetDiag(
+          StartLoc,
+          diag::err_omp_device_ancestor_without_requires_reverse_offload);
+      ErrorFound = true;
+    }
+  }
+
   OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective();
   OpenMPDirectiveKind CaptureRegion =
       getOpenMPCaptureRegionForClause(DKind, OMPC_device, LangOpts.OpenMP);
   if (CaptureRegion != OMPD_unknown && !CurContext->isDependentContext()) {
     ValExpr = MakeFullExpr(ValExpr).get();
     llvm::MapVector<const Expr *, DeclRefExpr *> Captures;
     ValExpr = tryBuildCapture(*this, ValExpr, Captures).get();
     HelperValStmt = buildPreInits(Context, Captures);

clang/test/OpenMP/target_ast_print.cpp

 ///==========================================================================///
 // RUN: %clang_cc1 -DOMP5 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix OMP5
 // RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix OMP5
 
 // RUN: %clang_cc1 -DOMP5 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix OMP5
 // RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix OMP5
-
+#pragma omp requires reverse_offload
 typedef void **omp_allocator_handle_t;
 extern const omp_allocator_handle_t omp_null_allocator;
 extern const omp_allocator_handle_t omp_default_mem_alloc;
 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
 extern const omp_allocator_handle_t omp_const_mem_alloc;
 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;

clang/test/OpenMP/target_device_ancestor_messages.cpp

+// RUN: %clang_cc1 -triple=x86_64 -verify -fopenmp -fopenmp-targets=x86_64 -x c++ -fexceptions -fcxx-exceptions %s
+// RUN: %clang_cc1 -triple=x86_64 -verify -fopenmp-simd -fopenmp-targets=x86_64 -x c++ -fexceptions -fcxx-exceptions %s
+
+void bar() {
+#pragma omp target device(ancestor : 1) // expected-error {{Device clause with ancestor device-modifier used without specifying 'requires reverse_offload'}}
+  ;
+}

clang/test/OpenMP/target_device_codegen.cpp

 // Test host codegen.
 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
-
+#pragma omp requires reverse_offload
 void foo(int n) {
 
   // CHECK:       [[N:%.+]] = load i32, i32* [[N_ADDR:%.+]],
   // CHECK:       store i32 [[N]], i32* [[DEVICE_CAP:%.+]],
   // CHECK:       [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
   // CHECK:       [[DEVICE:%.+]] = sext i32 [[DEV]] to i64
   // CHECK:       [[RET:%.+]] = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 [[DEVICE]], i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i64* null, i64* null, i8** null, i8** null)
   // CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
void foo(int n) {
   // CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
   // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
   // CHECK:       [[FAIL]]
   // CHECK:       call void [[HVT0:@.+]]()
   // CHECK-NEXT:  br label %[[END]]
   // CHECK:       [[END]]
   #pragma omp target device(device_num: n)
   ;
   // CHECK-NOT:   call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}},
   // CHECK:       call void @__omp_offloading_{{.+}}_l46()
   // CHECK-NOT:   call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}},
   #pragma omp target device(ancestor: n)
   ;
 }
 
 #endif