Skip to content

Commit 299849f

Browse files
author
z1_cciauto
authored
merge main into amd-staging (llvm#3435)
2 parents 1e855f6 + aca63a2 commit 299849f

31 files changed

+3132
-159
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -795,6 +795,8 @@ TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8, "V8fV16iV16iIsV8fIbI
795795
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8, "V8fV16iV16iIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
796796
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8, "V8fV16iV16iIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
797797
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8, "V8fV16iV16iIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
798+
TARGET_BUILTIN(__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4, "V8fIiV16iIiV16iIsV8fIiIiiIiIiiIbIb", "nc", "gfx1250-insts,wavefrontsize32")
799+
TARGET_BUILTIN(__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4, "V8fIiV16iIiV16iIsV8fIiIiLiIiIiLiIbIb", "nc", "gfx1250-insts,wavefrontsize32")
798800
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x32_f16, "V8fIbV16hIbV16hIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
799801
TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x32_f16, "V8hIbV16hIbV16hIsV8hIbIb", "nc", "gfx1250-insts,wavefrontsize32")
800802
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_32x16x128_f4, "V16fV16iV8iIsV16f", "nc", "gfx1250-insts,wavefrontsize32")

‎clang/lib/AST/ByteCode/EvalEmitter.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -233,6 +233,14 @@ template <> bool EvalEmitter::emitRet<PT_Ptr>(const SourceInfo &Info) {
233233
return false;
234234
}
235235
} else {
236+
// If this is pointing to a local variable, just return
237+
// the result, even if the pointer is dead.
238+
// This will later be diagnosed by CheckLValueConstantExpression.
239+
if (Ptr.isBlockPointer() && !Ptr.block()->isStatic()) {
240+
EvalResult.setValue(Ptr.toAPValue(Ctx.getASTContext()));
241+
return true;
242+
}
243+
236244
if (!Ptr.isLive() && !Ptr.isTemporary())
237245
return false;
238246

‎clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -921,6 +921,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
921921
case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:
922922
case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4:
923923
case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:
924+
case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
925+
case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:
924926
case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:
925927
case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16:
926928
case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16:
@@ -1187,6 +1189,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
11871189
ArgsForMatchingMatrixTypes = {5, 1, 3};
11881190
BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4;
11891191
break;
1192+
case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
1193+
ArgsForMatchingMatrixTypes = {5, 1, 3};
1194+
BuiltinWMMAOp = Intrinsic::amdgcn_wmma_scale_f32_16x16x128_f8f6f4;
1195+
break;
1196+
case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:
1197+
ArgsForMatchingMatrixTypes = {5, 1, 3};
1198+
BuiltinWMMAOp = Intrinsic::amdgcn_wmma_scale16_f32_16x16x128_f8f6f4;
1199+
break;
11901200
case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:
11911201
ArgsForMatchingMatrixTypes = {3, 0, 1};
11921202
BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_32x16x128_f4;

‎clang/lib/Driver/ToolChains/Clang.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1359,12 +1359,10 @@ void RenderARMABI(const Driver &D, const llvm::Triple &Triple,
13591359
// FIXME: Support -meabi.
13601360
// FIXME: Parts of this are duplicated in the backend, unify this somehow.
13611361
const char *ABIName = nullptr;
1362-
if (Arg *A = Args.getLastArg(options::OPT_mabi_EQ)) {
1362+
if (Arg *A = Args.getLastArg(options::OPT_mabi_EQ))
13631363
ABIName = A->getValue();
1364-
} else {
1365-
std::string CPU = getCPUName(D, Args, Triple, /*FromAs*/ false);
1364+
else
13661365
ABIName = llvm::ARM::computeDefaultTargetABI(Triple).data();
1367-
}
13681366

13691367
CmdArgs.push_back("-target-abi");
13701368
CmdArgs.push_back(ABIName);

‎clang/test/AST/ByteCode/cxx11.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -301,4 +301,11 @@ namespace NonConstLocal {
301301
}
302302
}
303303

304+
#define ATTR __attribute__((require_constant_initialization))
305+
int somefunc() {
306+
const int non_global = 42; // both-note {{declared here}}
307+
ATTR static const int &local_init = non_global; // both-error {{variable does not have a constant initializer}} \
308+
// both-note {{required by}} \
309+
// both-note {{reference to 'non_global' is not a constant expression}}
310+
}
304311

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-wmma-w32.cl

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -169,6 +169,30 @@ void test_amdgcn_wmma_f32_16x16x128_f8f6f4(global v8f* out, v16i a, v16i b, v8f
169169
*out = __builtin_amdgcn_wmma_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c);
170170
}
171171

172+
// CHECK-GFX1250-LABEL: @test_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(
173+
// CHECK-GFX1250-NEXT: entry:
174+
// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = shufflevector <16 x i32> [[B:%.*]], <16 x i32> poison, <12 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11>
175+
// CHECK-GFX1250-NEXT: [[TMP1:%.*]] = tail call <8 x float> @llvm.amdgcn.wmma.scale.f32.16x16x128.f8f6f4.v8f32.v16i32.v12i32(i32 1, <16 x i32> [[A:%.*]], i32 2, <12 x i32> [[TMP0]], i16 0, <8 x float> [[C:%.*]], i32 1, i32 2, i32 [[SCALE_SRC0:%.*]], i32 2, i32 1, i32 [[SCALE_SRC1:%.*]], i1 false, i1 true)
176+
// CHECK-GFX1250-NEXT: store <8 x float> [[TMP1]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
177+
// CHECK-GFX1250-NEXT: ret void
178+
//
179+
void test_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(global v8f* out, v16i a, v16i b, v8f c, int scale_src0, int scale_src1)
180+
{
181+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 2, scale_src0, 2, 1, scale_src1, 0, 1);
182+
}
183+
184+
// CHECK-GFX1250-LABEL: @test_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(
185+
// CHECK-GFX1250-NEXT: entry:
186+
// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = shufflevector <16 x i32> [[B:%.*]], <16 x i32> poison, <12 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11>
187+
// CHECK-GFX1250-NEXT: [[TMP1:%.*]] = tail call <8 x float> @llvm.amdgcn.wmma.scale16.f32.16x16x128.f8f6f4.v8f32.v16i32.v12i32(i32 1, <16 x i32> [[A:%.*]], i32 2, <12 x i32> [[TMP0]], i16 0, <8 x float> [[C:%.*]], i32 1, i32 2, i64 [[SCALE_SRC0:%.*]], i32 2, i32 1, i64 [[SCALE_SRC1:%.*]], i1 false, i1 true)
188+
// CHECK-GFX1250-NEXT: store <8 x float> [[TMP1]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
189+
// CHECK-GFX1250-NEXT: ret void
190+
//
191+
void test_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(global v8f* out, v16i a, v16i b, v8f c, long scale_src0, long scale_src1)
192+
{
193+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 2, scale_src0, 2, 1, scale_src1, 0, 1);
194+
}
195+
172196
// CHECK-GFX1250-LABEL: @test_amdgcn_wmma_f32_16x16x32_f16(
173197
// CHECK-GFX1250-NEXT: entry:
174198
// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> [[A:%.*]], i1 false, <16 x half> [[B:%.*]], i16 0, <8 x float> [[C:%.*]], i1 false, i1 true)

‎clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-wmma-w32-param.cl

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -121,6 +121,36 @@ void test_amdgcn_wmma_f32_16x16x128_f8f6f4(global v8f* out, v16i a, v16i b, v8f
121121
*out = __builtin_amdgcn_wmma_f32_16x16x128_f8f6f4(1, a, 2, b, mod, c); // expected-error {{'__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4' must be a constant integer}}
122122
}
123123

124+
void test_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(global v8f* out, v16i a, v16i b, v8f c, int mod, int scale_src0, int scale_src1, bool reuse)
125+
{
126+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(mod, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
127+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, mod, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, 0, 1); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
128+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, mod, c, 1, 0, scale_src0, 2, 0, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
129+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, mod, 0, scale_src0, 2, 0, scale_src1, 0, 0); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
130+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, mod, 0, scale_src1, 0, 0); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
131+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, reuse, 0); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
132+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, 0, reuse); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
133+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, mod, scale_src0, 2, 0, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
134+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, mod, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
135+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, mod, 0); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
136+
*out = __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, 1, mod); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
137+
}
138+
139+
void test_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(global v8f* out, v16i a, v16i b, v8f c, int mod, long scale_src0, long scale_src1, bool reuse)
140+
{
141+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(mod, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
142+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, mod, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, 0, 1); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
143+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, mod, c, 1, 0, scale_src0, 2, 0, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
144+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, mod, 0, scale_src0, 2, 0, scale_src1, 0, 0); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
145+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, mod, 0, scale_src1, 0, 0); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
146+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, reuse, 0); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
147+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, 0, reuse); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
148+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, mod, scale_src0, 2, 0, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
149+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, mod, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
150+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, mod, 0); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
151+
*out = __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4(1, a, 2, b, 0, c, 1, 0, scale_src0, 2, 0, scale_src1, 1, mod); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4' must be a constant integer}}
152+
}
153+
124154
void test_amdgcn_wmma_f32_16x16x32_f16(global v8f* out, v16h a, v16h b, v8f c, int mod)
125155
{
126156
*out = __builtin_amdgcn_wmma_f32_16x16x32_f16(mod, a, 0, b, 0, c, false, false); // expected-error {{'__builtin_amdgcn_wmma_f32_16x16x32_f16' must be a constant integer}}

‎flang/lib/Lower/Bridge.cpp

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1320,7 +1320,7 @@ class FirConverter : public Fortran::lower::AbstractConverter {
13201320
auto copyData = [&](hlfir::Entity l, hlfir::Entity r) {
13211321
// Dereference RHS and load it if trivial scalar.
13221322
r = hlfir::loadTrivialScalar(loc, *builder, r);
1323-
builder->create<hlfir::AssignOp>(loc, r, l, isAllocatable);
1323+
hlfir::AssignOp::create(*builder, loc, r, l, isAllocatable);
13241324
};
13251325

13261326
if (isPointer) {
@@ -3061,11 +3061,11 @@ class FirConverter : public Fortran::lower::AbstractConverter {
30613061
exprVal = builder->createConvert(loc, builder->getI1Type(), exprVal);
30623062
if (innerInsertionPoint.isSet())
30633063
builder->restoreInsertionPoint(innerInsertionPoint);
3064-
builder->create<hlfir::YieldOp>(loc, exprVal);
3064+
hlfir::YieldOp::create(*builder, loc, exprVal);
30653065
};
30663066
for (const Fortran::parser::ConcurrentControl &control :
30673067
std::get<std::list<Fortran::parser::ConcurrentControl>>(header.t)) {
3068-
auto forallOp = builder->create<hlfir::ForallOp>(loc);
3068+
auto forallOp = hlfir::ForallOp::create(*builder, loc);
30693069
if (isOutterForall && !outerForall)
30703070
outerForall = forallOp;
30713071
evaluateControl(std::get<1>(control.t), forallOp.getLbRegion());
@@ -3082,8 +3082,8 @@ class FirConverter : public Fortran::lower::AbstractConverter {
30823082
mlir::Type controlVarType = genType(*controlVar);
30833083
mlir::Block *forallBody = builder->createBlock(&forallOp.getBody(), {},
30843084
{controlVarType}, {loc});
3085-
auto forallIndex = builder->create<hlfir::ForallIndexOp>(
3086-
loc, fir::ReferenceType::get(controlVarType),
3085+
auto forallIndex = hlfir::ForallIndexOp::create(
3086+
*builder, loc, fir::ReferenceType::get(controlVarType),
30873087
forallBody->getArguments()[0],
30883088
builder->getStringAttr(controlVar->name().ToString()));
30893089
localSymbols.addVariableDefinition(*controlVar, forallIndex,
@@ -3096,7 +3096,7 @@ class FirConverter : public Fortran::lower::AbstractConverter {
30963096
std::get<std::optional<Fortran::parser::ScalarLogicalExpr>>(
30973097
header.t)) {
30983098
// Create hlfir.forall_mask and set insertion point in its body.
3099-
auto forallMaskOp = builder->create<hlfir::ForallMaskOp>(loc);
3099+
auto forallMaskOp = hlfir::ForallMaskOp::create(*builder, loc);
31003100
evaluateControl(*maskExpr, forallMaskOp.getMaskRegion(), /*isMask=*/true);
31013101
builder->createBlock(&forallMaskOp.getBody());
31023102
auto end = fir::FirEndOp::create(*builder, loc);
@@ -4577,14 +4577,14 @@ class FirConverter : public Fortran::lower::AbstractConverter {
45774577
// descriptor address/value and later implemented with a store.
45784578
// The RHS is fully prepared in lowering, so that all that is left
45794579
// in hlfir.region_assign code generation is the store.
4580-
auto regionAssignOp = builder->create<hlfir::RegionAssignOp>(loc);
4580+
auto regionAssignOp = hlfir::RegionAssignOp::create(*builder, loc);
45814581

45824582
// Lower LHS in its own region.
45834583
builder->createBlock(&regionAssignOp.getLhsRegion());
45844584
Fortran::lower::StatementContext lhsContext;
45854585
hlfir::Entity lhs = Fortran::lower::convertExprToHLFIR(
45864586
loc, *this, assign.lhs, localSymbols, lhsContext);
4587-
auto lhsYieldOp = builder->create<hlfir::YieldOp>(loc, lhs);
4587+
auto lhsYieldOp = hlfir::YieldOp::create(*builder, loc, lhs);
45884588
Fortran::lower::genCleanUpInRegionIfAny(
45894589
loc, *builder, lhsYieldOp.getCleanup(), lhsContext);
45904590

@@ -4593,7 +4593,7 @@ class FirConverter : public Fortran::lower::AbstractConverter {
45934593
Fortran::lower::StatementContext rhsContext;
45944594
mlir::Value rhs =
45954595
genForallPointerAssignmentRhs(loc, lhs, assign, rhsContext);
4596-
auto rhsYieldOp = builder->create<hlfir::YieldOp>(loc, rhs);
4596+
auto rhsYieldOp = hlfir::YieldOp::create(*builder, loc, rhs);
45974597
Fortran::lower::genCleanUpInRegionIfAny(
45984598
loc, *builder, rhsYieldOp.getCleanup(), rhsContext);
45994599

@@ -5364,7 +5364,7 @@ class FirConverter : public Fortran::lower::AbstractConverter {
53645364
if (!lowerToHighLevelFIR()) {
53655365
implicitIterSpace.growStack();
53665366
} else {
5367-
whereOp = builder->create<hlfir::WhereOp>(loc);
5367+
whereOp = hlfir::WhereOp::create(*builder, loc);
53685368
builder->createBlock(&whereOp.getMaskRegion());
53695369
}
53705370

@@ -5426,7 +5426,7 @@ class FirConverter : public Fortran::lower::AbstractConverter {
54265426
hlfir::Entity mask = Fortran::lower::convertExprToHLFIR(
54275427
loc, *this, *maskExpr, localSymbols, maskContext);
54285428
mask = hlfir::loadTrivialScalar(loc, *builder, mask);
5429-
auto yieldOp = builder->create<hlfir::YieldOp>(loc, mask);
5429+
auto yieldOp = hlfir::YieldOp::create(*builder, loc, mask);
54305430
Fortran::lower::genCleanUpInRegionIfAny(loc, *builder, yieldOp.getCleanup(),
54315431
maskContext);
54325432
}
@@ -5442,7 +5442,7 @@ class FirConverter : public Fortran::lower::AbstractConverter {
54425442
mlir::Location loc = getCurrentLocation();
54435443
hlfir::ElseWhereOp elsewhereOp;
54445444
if (lowerToHighLevelFIR()) {
5445-
elsewhereOp = builder->create<hlfir::ElseWhereOp>(loc);
5445+
elsewhereOp = hlfir::ElseWhereOp::create(*builder, loc);
54465446
// Lower mask in the mask region.
54475447
builder->createBlock(&elsewhereOp.getMaskRegion());
54485448
}
@@ -5470,7 +5470,7 @@ class FirConverter : public Fortran::lower::AbstractConverter {
54705470
void genFIR(const Fortran::parser::WhereConstruct::Elsewhere &ew) {
54715471
if (lowerToHighLevelFIR()) {
54725472
auto elsewhereOp =
5473-
builder->create<hlfir::ElseWhereOp>(getCurrentLocation());
5473+
hlfir::ElseWhereOp::create(*builder, getCurrentLocation());
54745474
builder->createBlock(&elsewhereOp.getBody());
54755475
}
54765476
genNestedStatement(
@@ -5496,7 +5496,7 @@ class FirConverter : public Fortran::lower::AbstractConverter {
54965496
std::get<Fortran::parser::LogicalExpr>(stmt.t));
54975497
if (lowerToHighLevelFIR()) {
54985498
mlir::Location loc = getCurrentLocation();
5499-
auto whereOp = builder->create<hlfir::WhereOp>(loc);
5499+
auto whereOp = hlfir::WhereOp::create(*builder, loc);
55005500
builder->createBlock(&whereOp.getMaskRegion());
55015501
lowerWhereMaskToHlfir(loc, mask);
55025502
builder->createBlock(&whereOp.getBody());

‎flang/lib/Lower/ConvertArrayConstructor.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -795,7 +795,7 @@ hlfir::EntityWithAttributes Fortran::lower::ArrayConstructorBuilder<T>::gen(
795795
// Insert the clean-up for the created hlfir.expr.
796796
fir::FirOpBuilder *bldr = &builder;
797797
stmtCtx.attachCleanup(
798-
[=]() { bldr->create<hlfir::DestroyOp>(loc, hlfirExpr); });
798+
[=]() { hlfir::DestroyOp::create(*bldr, loc, hlfirExpr); });
799799
return hlfir::EntityWithAttributes{hlfirExpr};
800800
}
801801

0 commit comments

Comments
 (0)