Skip to content

[OpenACC][CIR] link clause lowering for global declare #169524

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Nov 25, 2025
Merged

[OpenACC][CIR] link clause lowering for global declare #169524

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 9 additions & 5 deletions clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -231,16 +231,12 @@ namespace {
class OpenACCGlobalDeclareClauseEmitter final
: public OpenACCClauseVisitor<OpenACCGlobalDeclareClauseEmitter> {
CIRGenModule &cgm;
void clauseNotImplemented(const OpenACCClause &c) {
cgm.errorNYI(c.getSourceRange(), "OpenACC Global Declare Clause",
c.getClauseKind());
}

public:
OpenACCGlobalDeclareClauseEmitter(CIRGenModule &cgm) : cgm(cgm) {}

void VisitClause(const OpenACCClause &clause) {
clauseNotImplemented(clause);
llvm_unreachable("Invalid OpenACC clause on global Declare");
}

void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
Expand Down Expand Up @@ -271,6 +267,14 @@ class OpenACCGlobalDeclareClauseEmitter final
/*structured=*/true,
/*implicit=*/false, /*requiresDtor=*/true);
}

void VisitLinkClause(const OpenACCLinkClause &clause) {
for (const Expr *var : clause.getVarList())
cgm.emitGlobalOpenACCDeclareDataOperands<mlir::acc::DeclareLinkOp>(
var, mlir::acc::DataClause::acc_declare_link, {},
/*structured=*/true,
/*implicit=*/false, /*requiresDtor=*/false);
}
};
} // namespace

Expand Down
2 changes: 0 additions & 2 deletions clang/test/CIR/CodeGenOpenACC/combined-copy.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,8 +73,6 @@ void acc_compute(int parmVar) {
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc

// TODO: OpenACC: Represent alwaysin/alwaysout/always correctly. For now,
// these do nothing to the IR.
#pragma acc parallel loop copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3)
for(int i = 0; i < 5; ++i);
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc
Expand Down
2 changes: 0 additions & 2 deletions clang/test/CIR/CodeGenOpenACC/compute-copy.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,8 +65,6 @@ void acc_compute(int parmVar) {
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc

// TODO: OpenACC: Represent alwaysin/alwaysout/always correctly. For now,
// these do nothing to the IR.
#pragma acc parallel copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3)
;
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc
Expand Down
4 changes: 0 additions & 4 deletions clang/test/CIR/CodeGenOpenACC/declare-copy.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,15 +5,11 @@ struct HasSideEffects {
~HasSideEffects();
};

// TODO: OpenACC: Implement 'global', NS lowering.

struct Struct {
static const HasSideEffects StaticMemHSE;
static const HasSideEffects StaticMemHSEArr[5];
static const int StaticMemInt;

// TODO: OpenACC: Implement static-local lowering.

void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) {
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}})
// CHECK-NEXT: cir.alloca{{.*}}["this"
Expand Down
4 changes: 0 additions & 4 deletions clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,15 +5,11 @@ struct HasSideEffects {
~HasSideEffects();
};

// TODO: OpenACC: Implement 'global', NS lowering.

struct Struct {
static const HasSideEffects StaticMemHSE;
static const HasSideEffects StaticMemHSEArr[5];
static const int StaticMemInt;

// TODO: OpenACC: Implement static-local lowering.

void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) {
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}})
// CHECK-NEXT: cir.alloca{{.*}}["this"
Expand Down
4 changes: 0 additions & 4 deletions clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,15 +5,11 @@ struct HasSideEffects {
~HasSideEffects();
};

// TODO: OpenACC: Implement 'global', NS lowering.

struct Struct {
static const HasSideEffects StaticMemHSE;
static const HasSideEffects StaticMemHSEArr[5];
static const int StaticMemInt;

// TODO: OpenACC: Implement static-local lowering.

void MemFunc1(HasSideEffects *ArgHSE, int *ArgInt) {
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})
// CHECK-NEXT: cir.alloca{{.*}}["this"
Expand Down
114 changes: 108 additions & 6 deletions clang/test/CIR/CodeGenOpenACC/declare-link.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,14 +5,116 @@ struct HasSideEffects {
~HasSideEffects();
};

// TODO: OpenACC: Implement 'global', NS lowering.
HasSideEffects GlobalHSE1;
HasSideEffects GlobalHSEArr[5];
int GlobalInt1;

struct Struct {
static const HasSideEffects StaticMemHSE;
static const HasSideEffects StaticMemHSEArr[5];
static const int StaticMemInt;
#pragma acc declare link(GlobalHSE1, GlobalInt1, GlobalHSEArr[1:1])
// CHECK: acc.global_ctor @GlobalHSE1_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSE1 : !cir.ptr<!rec_HasSideEffects>
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "GlobalHSE1"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
// CHECK-NOT: acc.global_dtor
//
// CHECK: acc.global_ctor @GlobalInt1_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalInt1 : !cir.ptr<!s32i>
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "GlobalInt1"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
// CHECK: acc.global_ctor @GlobalHSEArr_acc_ctor {
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "GlobalHSEArr[1:1]"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }

namespace NS {

HasSideEffects NSHSE1;
HasSideEffects NSHSEArr[5];
int NSInt1;

#pragma acc declare link(NSHSE1, NSInt1, NSHSEArr[1:1])
// CHECK: acc.global_ctor @{{.*}}NSHSE1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "NSHSE1"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
// CHECK: acc.global_ctor @{{.*}}NSInt1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "NSInt1"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
// CHECK: acc.global_ctor @{{.*}}NSHSEArr{{.*}}_acc_ctor {
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "NSHSEArr[1:1]"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }

} // namespace NS

namespace {

// TODO: OpenACC: Implement static-local lowering.
HasSideEffects AnonNSHSE1;
HasSideEffects AnonNSHSEArr[5];
int AnonNSInt1;

#pragma acc declare link(AnonNSHSE1, AnonNSInt1, AnonNSHSEArr[1:1])
// CHECK: acc.global_ctor @{{.*}}AnonNSHSE1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "AnonNSHSE1"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
// CHECK: acc.global_ctor @{{.*}}AnonNSInt1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSInt1{{.*}} : !cir.ptr<!s32i>
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "AnonNSInt1"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
// CHECK: acc.global_ctor @{{.*}}AnonNSHSEArr{{.*}}_acc_ctor {
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "AnonNSHSEArr[1:1]"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }

} // namespace NS


struct Struct {

void MemFunc1() {
// CHECK: cir.func {{.*}}MemFunc1{{.*}}({{.*}}) {
Expand Down
4 changes: 0 additions & 4 deletions clang/test/CIR/CodeGenOpenACC/declare-present.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,15 +5,11 @@ struct HasSideEffects {
~HasSideEffects();
};

// TODO: OpenACC: Implement 'global', NS lowering.

struct Struct {
static const HasSideEffects StaticMemHSE;
static const HasSideEffects StaticMemHSEArr[5];
static const int StaticMemInt;

// TODO: OpenACC: Implement static-local lowering.

void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) {
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}})
// CHECK-NEXT: cir.alloca{{.*}}["this"
Expand Down
Loading