Author: erichkeane Date: 2025-04-22T10:06:28-07:00 New Revision: c05da6e8611e4da7b16729cc2b8d2dc6c77ba4be
URL: https://github.com/llvm/llvm-project/commit/c05da6e8611e4da7b16729cc2b8d2dc6c77ba4be DIFF: https://github.com/llvm/llvm-project/commit/c05da6e8611e4da7b16729cc2b8d2dc6c77ba4be.diff LOG: [OpenACC][CIR] Implement lowering for 'set' clauses The 'set' clauses are default_async, device_num, and if. The latter two are implemented identically to other constructs by that name. default_async works exactly like device_num, (and others) that take an int-expression. Added: Modified: clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp clang/test/CIR/CodeGenOpenACC/set.c Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp index 621d5fa335db3..3696bfbd8e4db 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -170,6 +170,7 @@ class OpenACCClauseCIREmitter final break; } } else { + // Combined Constructs left. return clauseNotImplemented(clause); } } @@ -208,6 +209,7 @@ class OpenACCClauseCIREmitter final // they just modify the other clauses IR. So setting of `lastDeviceType` // (done above) is all we need. } else { + // update, data, loop, routine, combined remain. return clauseNotImplemented(clause); } } @@ -221,6 +223,7 @@ class OpenACCClauseCIREmitter final } else if constexpr (isOneOfTypes<OpTy, SerialOp>) { llvm_unreachable("num_workers not valid on serial"); } else { + // Combined Remain. return clauseNotImplemented(clause); } } @@ -234,6 +237,7 @@ class OpenACCClauseCIREmitter final } else if constexpr (isOneOfTypes<OpTy, SerialOp>) { llvm_unreachable("vector_length not valid on serial"); } else { + // Combined remain. return clauseNotImplemented(clause); } } @@ -250,6 +254,7 @@ class OpenACCClauseCIREmitter final createIntExpr(clause.getIntExpr()), &range)); } } else { + // Data, enter data, exit data, update, wait, combined remain. return clauseNotImplemented(clause); } } @@ -266,19 +271,21 @@ class OpenACCClauseCIREmitter final llvm_unreachable("var-list version of self shouldn't get here"); } } else { + // update and combined remain. return clauseNotImplemented(clause); } } void VisitIfClause(const OpenACCIfClause &clause) { if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, InitOp, - ShutdownOp>) { + ShutdownOp, SetOp>) { operation.getIfCondMutable().append( createCondition(clause.getConditionExpr())); } else { // 'if' applies to most of the constructs, but hold off on lowering them // until we can write tests/know what we're doing with codegen to make // sure we get it right. + // Enter data, exit data, host_data, update, wait, combined remain. return clauseNotImplemented(clause); } } @@ -287,8 +294,23 @@ class OpenACCClauseCIREmitter final if constexpr (isOneOfTypes<OpTy, InitOp, ShutdownOp>) { operation.getDeviceNumOperandMutable().append( createIntExpr(clause.getIntExpr())); + } else if constexpr (isOneOfTypes<OpTy, SetOp>) { + // This is only a separate case because the getter name is diff erent in + // 'set' for some reason. + operation.getDeviceNumMutable().append( + createIntExpr(clause.getIntExpr())); } else { - return clauseNotImplemented(clause); + llvm_unreachable( + "init, shutdown, set, are only valid device_num constructs"); + } + } + + void VisitDefaultAsyncClause(const OpenACCDefaultAsyncClause &clause) { + if constexpr (isOneOfTypes<OpTy, SetOp>) { + operation.getDefaultAsyncMutable().append( + createIntExpr(clause.getIntExpr())); + } else { + llvm_unreachable("set, is only valid device_num constructs"); } } }; diff --git a/clang/test/CIR/CodeGenOpenACC/set.c b/clang/test/CIR/CodeGenOpenACC/set.c index 0e85dd1b5f78e..3a8a9c24569e5 100644 --- a/clang/test/CIR/CodeGenOpenACC/set.c +++ b/clang/test/CIR/CodeGenOpenACC/set.c @@ -1,7 +1,9 @@ // RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s -void acc_set(void) { - // CHECK: cir.func @acc_set() { +void acc_set(int cond) { + // CHECK: cir.func @acc_set(%[[ARG:.*]]: !s32i{{.*}}) { + // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init] + // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i> #pragma acc set device_type(*) // CHECK-NEXT: acc.set attributes {device_type = #acc.device_type<star>} @@ -10,5 +12,33 @@ void acc_set(void) { #pragma acc set device_type(radeon) // CHECK-NEXT: acc.set attributes {device_type = #acc.device_type<radeon>} +#pragma acc set default_async(cond) + // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i + // CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32 + // CHECK-NEXT: acc.set default_async(%[[COND_CONV]] : si32) + +#pragma acc set default_async(1) + // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32 + // CHECK-NEXT: acc.set default_async(%[[ONE_CONV]] : si32) + +#pragma acc set device_num(cond) if (cond) + // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i + // CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32 + // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i + // CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool + // CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.set device_num(%[[COND_CONV]] : si32) if(%[[BOOL_CONV]]) + +#pragma acc set device_type(radeon) default_async(1) device_num(cond) if (cond) + // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32 + // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i + // CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32 + // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i + // CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool + // CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.set default_async(%[[ONE_CONV]] : si32) device_num(%[[COND_CONV]] : si32) if(%[[BOOL_CONV]]) attributes {device_type = #acc.device_type<radeon>} + // CHECK-NEXT: cir.return } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
