commit 77db154cf9383fa20685fca56633601ce0ea47fa
author erichkeane <ekeane@nvidia.com> Thu Apr 10 10:13:09 2025 -0700
committer erichkeane <ekeane@nvidia.com> Thu Apr 10 17:45:23 2025 -0700
tree 6894b2ae657967b417fc01845f6b4649691cf405
parent f4203ca2b7b047fc556535661f406e7fef758835

[OpenACC][CIR] implement basic 'set' lowering with device_type clause

The 'set' lowering is pretty trivial.  'device_type' is a little more
restricted since both the MLIR-Dialect and language limit it to only 1
value (as confirmed by standards-discussion).

This patch implements 'set', with 'device_type', since 'set' requires at
least 1 clause, and  this is the least difficult to implement at the
moment.

clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp

diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 6fc5a7e..5e3f826 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -92,6 +92,7 @@
 
     switch (dirKind) {
     case OpenACCDirectiveKind::Init:
+    case OpenACCDirectiveKind::Set:
     case OpenACCDirectiveKind::Shutdown: {
       // Device type has a list that is either a 'star' (emitted as 'star'),
       // or an identifer list, all of which get added for attributes.
@@ -133,6 +134,11 @@
 
         op.setDeviceTypesAttr(
             mlir::ArrayAttr::get(builder.getContext(), deviceTypes));
+      } else if constexpr (isOneOfTypes<Op, SetOp>) {
+        assert(attrData.deviceTypeArchs.size() <= 1 &&
+               "Set can only have a single architecture");
+        if (!attrData.deviceTypeArchs.empty())
+          op.setDeviceType(attrData.deviceTypeArchs[0]);
       } else {
         cgm.errorNYI(dirLoc, "OpenACC 'device_type' clause lowering for ",
                      dirKind);
@@ -232,6 +238,13 @@
                                s.clauses());
 }
 
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCSetConstruct(const OpenACCSetConstruct &s) {
+  mlir::Location start = getLoc(s.getSourceRange().getEnd());
+  return emitOpenACCOp<SetOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
+                              s.clauses());
+}
+
 mlir::LogicalResult CIRGenFunction::emitOpenACCShutdownConstruct(
     const OpenACCShutdownConstruct &s) {
   mlir::Location start = getLoc(s.getSourceRange().getEnd());
@@ -270,11 +283,6 @@
   return mlir::failure();
 }
 mlir::LogicalResult
-CIRGenFunction::emitOpenACCSetConstruct(const OpenACCSetConstruct &s) {
-  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Set Construct");
-  return mlir::failure();
-}
-mlir::LogicalResult
 CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
   getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Update Construct");
   return mlir::failure();

clang/test/CIR/CodeGenOpenACC/set.c

diff --git a/clang/test/CIR/CodeGenOpenACC/set.c b/clang/test/CIR/CodeGenOpenACC/set.c
new file mode 100644
index 0000000..0e85dd1
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/set.c
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_set(void) {
+  // CHECK: cir.func @acc_set() {
+
+#pragma acc set device_type(*)
+  // CHECK-NEXT: acc.set attributes {device_type = #acc.device_type<star>}
+
+  // Set doesn't allow multiple device_type clauses, so no need to test them.
+#pragma acc set device_type(radeon)
+  // CHECK-NEXT: acc.set attributes {device_type = #acc.device_type<radeon>}
+
+  // CHECK-NEXT: cir.return
+}