Skip to content

Commit 25ad296

Browse files
s-barannikovs-barannikov
committed
Merge remote-tracking branch 'upstream/main' into tti/constify-this
2 parents 3c67155 + c873ca2 commit 25ad296

File tree

84 files changed

+1415
-552
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

84 files changed

+1415
-552
lines changed

‎clang/cmake/caches/Fuchsia-stage2.cmake

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,7 @@ set(LLVM_ENABLE_LIBEDIT OFF CACHE BOOL "")
1818
set(LLVM_ENABLE_LLD ON CACHE BOOL "")
1919
set(LLVM_ENABLE_LTO ON CACHE BOOL "")
2020
set(LLVM_ENABLE_PER_TARGET_RUNTIME_DIR ON CACHE BOOL "")
21+
set(LLVM_ENABLE_PIC OFF CACHE BOOL "")
2122
set(LLVM_ENABLE_PLUGINS OFF CACHE BOOL "")
2223
set(LLVM_ENABLE_UNWIND_TABLES OFF CACHE BOOL "")
2324
set(LLVM_ENABLE_Z3_SOLVER OFF CACHE BOOL "")

‎clang/docs/ReleaseNotes.rst

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -301,6 +301,9 @@ related warnings within the method body.
301301
particularly relevant for AMDGPU targets, where they map to corresponding IR
302302
metadata.
303303

304+
- Clang now disallows the use of attributes applied before an
305+
``extern template`` declaration (#GH79893).
306+
304307
Improvements to Clang's diagnostics
305308
-----------------------------------
306309

‎clang/include/clang/CIR/Dialect/IR/CIRTypesDetails.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -54,7 +54,7 @@ struct RecordTypeStorage : public mlir::TypeStorage {
5454
RecordType::RecordKind kind)
5555
: members(members), name(name), incomplete(incomplete), packed(packed),
5656
padded(padded), kind(kind) {
57-
assert(name || !incomplete && "Incomplete records must have a name");
57+
assert((name || !incomplete) && "Incomplete records must have a name");
5858
}
5959

6060
KeyTy getAsKey() const {

‎clang/include/clang/Driver/Options.td

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -4702,15 +4702,17 @@ def EB : Flag<["-"], "EB">, Alias<mbig_endian>;
47024702
def m16 : Flag<["-"], "m16">, Group<m_Group>, Flags<[NoXarchOption]>,
47034703
Visibility<[ClangOption, CLOption, DXCOption]>;
47044704
def m32 : Flag<["-"], "m32">, Group<m_Group>, Flags<[NoXarchOption]>,
4705-
Visibility<[ClangOption, CLOption, DXCOption]>;
4706-
def maix32 : Flag<["-"], "maix32">, Group<m_Group>, Flags<[NoXarchOption]>;
4705+
Visibility<[ClangOption, CLOption, DXCOption, FlangOption]>;
4706+
def maix32 : Flag<["-"], "maix32">, Group<m_Group>, Flags<[NoXarchOption]>,
4707+
Visibility<[FlangOption]>;
47074708
def mqdsp6_compat : Flag<["-"], "mqdsp6-compat">, Group<m_Group>,
47084709
Flags<[NoXarchOption]>, Visibility<[ClangOption, CC1Option]>,
47094710
HelpText<"Enable hexagon-qdsp6 backward compatibility">,
47104711
MarshallingInfoFlag<LangOpts<"HexagonQdsp6Compat">>;
47114712
def m64 : Flag<["-"], "m64">, Group<m_Group>, Flags<[NoXarchOption]>,
47124713
Visibility<[ClangOption, CLOption, DXCOption, FlangOption]>;
4713-
def maix64 : Flag<["-"], "maix64">, Group<m_Group>, Flags<[NoXarchOption]>;
4714+
def maix64 : Flag<["-"], "maix64">, Group<m_Group>, Flags<[NoXarchOption]>,
4715+
Visibility<[FlangOption]>;
47144716
def mx32 : Flag<["-"], "mx32">, Group<m_Group>, Flags<[NoXarchOption]>,
47154717
Visibility<[ClangOption, CLOption, DXCOption]>;
47164718
def miamcu : Flag<["-"], "miamcu">, Group<m_Group>, Flags<[NoXarchOption]>,

‎clang/include/clang/Frontend/CompilerInstance.h

Lines changed: 11 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -845,18 +845,25 @@ class CompilerInstance : public ModuleLoader {
845845
/// Creates a \c CompilerInstance for compiling a module.
846846
///
847847
/// This expects a properly initialized \c FrontendInputFile.
848+
///
849+
/// Explicitly-specified \c VFS takes precedence over the VFS of this instance
850+
/// when creating the clone and also prevents \c FileManager sharing.
848851
std::unique_ptr<CompilerInstance> cloneForModuleCompileImpl(
849852
SourceLocation ImportLoc, StringRef ModuleName, FrontendInputFile Input,
850-
StringRef OriginalModuleMapFile, StringRef ModuleFileName);
853+
StringRef OriginalModuleMapFile, StringRef ModuleFileName,
854+
IntrusiveRefCntPtr<llvm::vfs::FileSystem> VFS = nullptr);
851855

852856
public:
853857
/// Creates a new \c CompilerInstance for compiling a module.
854858
///
855859
/// This takes care of creating appropriate \c FrontendInputFile for
856860
/// public/private frameworks, inferred modules and such.
857-
std::unique_ptr<CompilerInstance>
858-
cloneForModuleCompile(SourceLocation ImportLoc, Module *Module,
859-
StringRef ModuleFileName);
861+
///
862+
/// Explicitly-specified \c VFS takes precedence over the VFS of this instance
863+
/// when creating the clone and also prevents \c FileManager sharing.
864+
std::unique_ptr<CompilerInstance> cloneForModuleCompile(
865+
SourceLocation ImportLoc, Module *Module, StringRef ModuleFileName,
866+
IntrusiveRefCntPtr<llvm::vfs::FileSystem> VFS = nullptr);
860867

861868
/// Compile a module file for the given module, using the options
862869
/// provided by the importing compiler instance. Returns true if the module

‎clang/lib/CIR/CodeGen/CIRGenBuilder.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -63,6 +63,7 @@ class CIRGenBuilderTy : public cir::CIRBaseBuilderTy {
6363
case clang::TagTypeKind::Enum:
6464
llvm_unreachable("enums are not records");
6565
}
66+
llvm_unreachable("Unsupported record kind");
6667
}
6768

6869
/// Get an incomplete CIR struct type. If we have a complete record

‎clang/lib/CIR/CodeGen/CIRGenTypes.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -153,7 +153,7 @@ isSafeToConvert(const RecordDecl *rd, CIRGenTypes &cgt,
153153
// out, don't do it. This includes virtual base classes which get laid out
154154
// when a class is translated, even though they aren't embedded by-value into
155155
// the class.
156-
if (const CXXRecordDecl *crd = dyn_cast<CXXRecordDecl>(rd)) {
156+
if (isa<CXXRecordDecl>(rd)) {
157157
assert(!cir::MissingFeatures::cxxSupport());
158158
cgt.getCGModule().errorNYI(rd->getSourceRange(),
159159
"isSafeToConvert: CXXRecordDecl");
@@ -237,7 +237,7 @@ mlir::Type CIRGenTypes::convertRecordDeclType(const clang::RecordDecl *rd) {
237237
assert(insertResult && "isSafeToCovert() should have caught this.");
238238

239239
// Force conversion of non-virtual base classes recursively.
240-
if (const auto *cxxRecordDecl = dyn_cast<CXXRecordDecl>(rd)) {
240+
if (isa<CXXRecordDecl>(rd)) {
241241
cgm.errorNYI(rd->getSourceRange(), "CXXRecordDecl");
242242
}
243243

‎clang/lib/Driver/Driver.cpp

Lines changed: 10 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -733,11 +733,16 @@ static llvm::Triple computeTargetTriple(const Driver &D,
733733
Target.setEnvironment(llvm::Triple::GNUX32);
734734
} else if (A->getOption().matches(options::OPT_m32) ||
735735
A->getOption().matches(options::OPT_maix32)) {
736-
AT = Target.get32BitArchVariant().getArch();
737-
if (Target.getEnvironment() == llvm::Triple::GNUX32)
738-
Target.setEnvironment(llvm::Triple::GNU);
739-
else if (Target.getEnvironment() == llvm::Triple::MuslX32)
740-
Target.setEnvironment(llvm::Triple::Musl);
736+
if (D.IsFlangMode() && !Target.isOSAIX()) {
737+
D.Diag(diag::err_drv_unsupported_opt_for_target)
738+
<< A->getAsString(Args) << Target.str();
739+
} else {
740+
AT = Target.get32BitArchVariant().getArch();
741+
if (Target.getEnvironment() == llvm::Triple::GNUX32)
742+
Target.setEnvironment(llvm::Triple::GNU);
743+
else if (Target.getEnvironment() == llvm::Triple::MuslX32)
744+
Target.setEnvironment(llvm::Triple::Musl);
745+
}
741746
} else if (A->getOption().matches(options::OPT_m16) &&
742747
Target.get32BitArchVariant().getArch() == llvm::Triple::x86) {
743748
AT = llvm::Triple::x86;

‎clang/lib/Frontend/CompilerInstance.cpp

Lines changed: 17 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -1152,7 +1152,8 @@ static Language getLanguageFromOptions(const LangOptions &LangOpts) {
11521152

11531153
std::unique_ptr<CompilerInstance> CompilerInstance::cloneForModuleCompileImpl(
11541154
SourceLocation ImportLoc, StringRef ModuleName, FrontendInputFile Input,
1155-
StringRef OriginalModuleMapFile, StringRef ModuleFileName) {
1155+
StringRef OriginalModuleMapFile, StringRef ModuleFileName,
1156+
IntrusiveRefCntPtr<llvm::vfs::FileSystem> VFS) {
11561157
// Construct a compiler invocation for creating this module.
11571158
auto Invocation = std::make_shared<CompilerInvocation>(getInvocation());
11581159

@@ -1212,19 +1213,21 @@ std::unique_ptr<CompilerInstance> CompilerInstance::cloneForModuleCompileImpl(
12121213
auto &Inv = *Invocation;
12131214
Instance.setInvocation(std::move(Invocation));
12141215

1216+
if (VFS) {
1217+
Instance.createFileManager(std::move(VFS));
1218+
} else if (FrontendOpts.ModulesShareFileManager) {
1219+
Instance.setFileManager(&getFileManager());
1220+
} else {
1221+
Instance.createFileManager(&getVirtualFileSystem());
1222+
}
1223+
12151224
Instance.createDiagnostics(
1216-
getVirtualFileSystem(),
1225+
Instance.getVirtualFileSystem(),
12171226
new ForwardingDiagnosticConsumer(getDiagnosticClient()),
12181227
/*ShouldOwnClient=*/true);
1219-
12201228
if (llvm::is_contained(DiagOpts.SystemHeaderWarningsModules, ModuleName))
12211229
Instance.getDiagnostics().setSuppressSystemWarnings(false);
12221230

1223-
if (FrontendOpts.ModulesShareFileManager) {
1224-
Instance.setFileManager(&getFileManager());
1225-
} else {
1226-
Instance.createFileManager(&getVirtualFileSystem());
1227-
}
12281231
Instance.createSourceManager(Instance.getFileManager());
12291232
SourceManager &SourceMgr = Instance.getSourceManager();
12301233

@@ -1318,7 +1321,8 @@ static OptionalFileEntryRef getPublicModuleMap(FileEntryRef File,
13181321
}
13191322

13201323
std::unique_ptr<CompilerInstance> CompilerInstance::cloneForModuleCompile(
1321-
SourceLocation ImportLoc, Module *Module, StringRef ModuleFileName) {
1324+
SourceLocation ImportLoc, Module *Module, StringRef ModuleFileName,
1325+
IntrusiveRefCntPtr<llvm::vfs::FileSystem> VFS) {
13221326
StringRef ModuleName = Module->getTopLevelModuleName();
13231327

13241328
InputKind IK(getLanguageFromOptions(getLangOpts()), InputKind::ModuleMap);
@@ -1363,7 +1367,8 @@ std::unique_ptr<CompilerInstance> CompilerInstance::cloneForModuleCompile(
13631367
return cloneForModuleCompileImpl(
13641368
ImportLoc, ModuleName,
13651369
FrontendInputFile(ModuleMapFilePath, IK, IsSystem),
1366-
ModMap.getModuleMapFileForUniquing(Module)->getName(), ModuleFileName);
1370+
ModMap.getModuleMapFileForUniquing(Module)->getName(), ModuleFileName,
1371+
std::move(VFS));
13671372
}
13681373

13691374
// FIXME: We only need to fake up an input file here as a way of
@@ -1380,7 +1385,8 @@ std::unique_ptr<CompilerInstance> CompilerInstance::cloneForModuleCompile(
13801385
auto Instance = cloneForModuleCompileImpl(
13811386
ImportLoc, ModuleName,
13821387
FrontendInputFile(FakeModuleMapFile, IK, +Module->IsSystem),
1383-
ModMap.getModuleMapFileForUniquing(Module)->getName(), ModuleFileName);
1388+
ModMap.getModuleMapFileForUniquing(Module)->getName(), ModuleFileName,
1389+
std::move(VFS));
13841390

13851391
std::unique_ptr<llvm::MemoryBuffer> ModuleMapBuffer =
13861392
llvm::MemoryBuffer::getMemBufferCopy(InferredModuleMapContent);

‎clang/lib/Parse/Parser.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1049,6 +1049,8 @@ Parser::ParseExternalDeclaration(ParsedAttributes &Attrs,
10491049

10501050
case tok::kw_extern:
10511051
if (getLangOpts().CPlusPlus && NextToken().is(tok::kw_template)) {
1052+
ProhibitAttributes(Attrs);
1053+
ProhibitAttributes(DeclSpecAttrs);
10521054
// Extern templates
10531055
SourceLocation ExternLoc = ConsumeToken();
10541056
SourceLocation TemplateLoc = ConsumeToken();

‎clang/lib/Sema/SemaExprCXX.cpp

Lines changed: 4 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -4745,19 +4745,13 @@ Sema::PerformImplicitConversion(Expr *From, QualType ToType,
47454745
case ICK_HLSL_Array_RValue:
47464746
if (ToType->isArrayParameterType()) {
47474747
FromType = Context.getArrayParameterType(FromType);
4748-
From = ImpCastExprToType(From, FromType, CK_HLSLArrayRValue, VK_PRValue,
4749-
/*BasePath=*/nullptr, CCK)
4750-
.get();
4751-
} else { // FromType must be ArrayParameterType
4752-
assert(FromType->isArrayParameterType() &&
4753-
"FromType must be ArrayParameterType in ICK_HLSL_Array_RValue \
4754-
if it is not ToType");
4748+
} else if (FromType->isArrayParameterType()) {
47554749
const ArrayParameterType *APT = cast<ArrayParameterType>(FromType);
47564750
FromType = APT->getConstantArrayType(Context);
4757-
From = ImpCastExprToType(From, FromType, CK_HLSLArrayRValue, VK_PRValue,
4758-
/*BasePath=*/nullptr, CCK)
4759-
.get();
47604751
}
4752+
From = ImpCastExprToType(From, FromType, CK_HLSLArrayRValue, VK_PRValue,
4753+
/*BasePath=*/nullptr, CCK)
4754+
.get();
47614755
break;
47624756

47634757
case ICK_Function_To_Pointer:

‎clang/lib/Sema/SemaOverload.cpp

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -2275,17 +2275,16 @@ static bool IsStandardConversion(Sema &S, Expr* From, QualType ToType,
22752275
// handling here.
22762276
if (ToType->isArrayParameterType()) {
22772277
FromType = S.Context.getArrayParameterType(FromType);
2278-
SCS.First = ICK_HLSL_Array_RValue;
22792278
} else if (FromType->isArrayParameterType()) {
22802279
const ArrayParameterType *APT = cast<ArrayParameterType>(FromType);
22812280
FromType = APT->getConstantArrayType(S.Context);
2282-
SCS.First = ICK_HLSL_Array_RValue;
2283-
} else {
2284-
SCS.First = ICK_Identity;
22852281
}
22862282

2287-
if (S.Context.getCanonicalType(FromType) !=
2288-
S.Context.getCanonicalType(ToType))
2283+
SCS.First = ICK_HLSL_Array_RValue;
2284+
2285+
// Don't consider qualifiers, which include things like address spaces
2286+
if (FromType.getCanonicalType().getUnqualifiedType() !=
2287+
ToType.getCanonicalType().getUnqualifiedType())
22892288
return false;
22902289

22912290
SCS.setAllToTypes(ToType);

‎clang/test/CodeGenDirectX/unsupported_intrinsic.hlsl

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,8 @@
1+
// REQUIRES: directx-registered-target
2+
// RUN: not %clang_dxc -T lib_6_3 %s 2>&1 | FileCheck %s
3+
4+
// CHECK: error: Unsupported intrinsic llvm.vector.reduce.and.v4i32 for DXIL lowering
5+
6+
export int vecReduceAndTest(int4 vec) {
7+
return __builtin_reduce_and(vec);
8+
}

‎clang/test/CodeGenHLSL/ArrayAssignable.hlsl

Lines changed: 62 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,23 @@
1-
// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s --enable-var-scope
1+
// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -finclude-default-header -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s
2+
3+
struct S {
4+
int x;
5+
float f;
6+
};
7+
8+
// CHECK: [[CBLayout:%.*]] = type <{ [2 x float], [2 x <4 x i32>], [2 x [2 x i32]], [1 x target("dx.Layout", %S, 8, 0, 4)] }>
9+
// CHECK: @CBArrays.cb = global target("dx.CBuffer", target("dx.Layout", [[CBLayout]], 136, 0, 32, 64, 128))
10+
// CHECK: @c1 = external addrspace(2) global [2 x float], align 4
11+
// CHECK: @c2 = external addrspace(2) global [2 x <4 x i32>], align 16
12+
// CHECK: @c3 = external addrspace(2) global [2 x [2 x i32]], align 4
13+
// CHECK: @c4 = external addrspace(2) global [1 x target("dx.Layout", %S, 8, 0, 4)], align 4
14+
15+
cbuffer CBArrays : register(b0) {
16+
float c1[2];
17+
int4 c2[2];
18+
int c3[2][2];
19+
S c4[1];
20+
}
221

322
// CHECK-LABEL: define void {{.*}}arr_assign1
423
// CHECK: [[Arr:%.*]] = alloca [2 x i32], align 4
@@ -116,3 +135,45 @@ void arr_assign7() {
116135
int Arr2[2][2] = {{0, 0}, {1, 1}};
117136
(Arr = Arr2)[0] = {6, 6};
118137
}
138+
139+
// Verify you can assign from a cbuffer array
140+
141+
// CHECK-LABEL: define void {{.*}}arr_assign8
142+
// CHECK: [[C:%.*]] = alloca [2 x float], align 4
143+
// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[C]], ptr align 4 {{.*}}, i32 8, i1 false)
144+
// CHECK-NEXT: call void @llvm.memcpy.p0.p2.i32(ptr align 4 [[C]], ptr addrspace(2) align 4 @c1, i32 8, i1 false)
145+
// CHECK-NEXT: ret void
146+
void arr_assign8() {
147+
float C[2] = {1.0, 2.0};
148+
C = c1;
149+
}
150+
151+
// CHECK-LABEL: define void {{.*}}arr_assign9
152+
// CHECK: [[C:%.*]] = alloca [2 x <4 x i32>], align 16
153+
// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 16 [[C]], ptr align 16 {{.*}}, i32 32, i1 false)
154+
// CHECK-NEXT: call void @llvm.memcpy.p0.p2.i32(ptr align 16 [[C]], ptr addrspace(2) align 16 @c2, i32 32, i1 false)
155+
// CHECK-NEXT: ret void
156+
void arr_assign9() {
157+
int4 C[2] = {1,2,3,4,5,6,7,8};
158+
C = c2;
159+
}
160+
161+
// CHECK-LABEL: define void {{.*}}arr_assign10
162+
// CHECK: [[C:%.*]] = alloca [2 x [2 x i32]], align 4
163+
// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[C]], ptr align 4 {{.*}}, i32 16, i1 false)
164+
// CHECK-NEXT: call void @llvm.memcpy.p0.p2.i32(ptr align 4 [[C]], ptr addrspace(2) align 4 @c3, i32 16, i1 false)
165+
// CHECK-NEXT: ret void
166+
void arr_assign10() {
167+
int C[2][2] = {1,2,3,4};
168+
C = c3;
169+
}
170+
171+
// CHECK-LABEL: define void {{.*}}arr_assign11
172+
// CHECK: [[C:%.*]] = alloca [1 x %struct.S], align 4
173+
// CHECK: call void @llvm.memcpy.p0.p2.i32(ptr align 4 [[C]], ptr addrspace(2) align 4 @c4, i32 8, i1 false)
174+
// CHECK-NEXT: ret void
175+
void arr_assign11() {
176+
S s = {1, 2.0};
177+
S C[1] = {s};
178+
C = c4;
179+
}

‎clang/test/Parser/extern-template-attributes.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,8 @@
1+
// RUN: %clang_cc1 -std=c++17 -fms-extensions -verify %s
2+
3+
template <class>
4+
struct S {};
5+
6+
[[deprecated]] extern template struct S<int>; // expected-error {{an attribute list cannot appear here}}
7+
__attribute__((deprecated)) extern template struct S<int>; // expected-error {{an attribute list cannot appear here}}
8+
__declspec(deprecated) extern template struct S<int>; // expected-error {{expected unqualified-id}}

‎flang/lib/Optimizer/Analysis/AliasAnalysis.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -542,6 +542,14 @@ AliasAnalysis::Source AliasAnalysis::getSource(mlir::Value v,
542542
v = op->getOperand(0);
543543
defOp = v.getDefiningOp();
544544
})
545+
.Case<fir::PackArrayOp>([&](auto op) {
546+
// The packed array is not distinguishable from the original
547+
// array, so skip PackArrayOp and track further through
548+
// the array operand.
549+
v = op.getArray();
550+
defOp = v.getDefiningOp();
551+
approximateSource = true;
552+
})
545553
.Case<fir::BoxAddrOp>([&](auto op) {
546554
v = op->getOperand(0);
547555
defOp = v.getDefiningOp();

‎flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp

Lines changed: 14 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -82,6 +82,11 @@ struct GPULaunchKernelConversion
8282
mlir::LogicalResult
8383
matchAndRewrite(mlir::gpu::LaunchFuncOp op, OpAdaptor adaptor,
8484
mlir::ConversionPatternRewriter &rewriter) const override {
85+
// Only convert gpu.launch_func for CUDA Fortran.
86+
if (!op.getOperation()->getAttrOfType<cuf::ProcAttributeAttr>(
87+
cuf::getProcAttrName()))
88+
return mlir::failure();
89+
8590
mlir::Location loc = op.getLoc();
8691
auto *ctx = rewriter.getContext();
8792
mlir::ModuleOp mod = op->getParentOfType<mlir::ModuleOp>();
@@ -293,7 +298,15 @@ class CUFGPUToLLVMConversion
293298
fir::LLVMTypeConverter typeConverter(module, /*applyTBAA=*/false,
294299
/*forceUnifiedTBAATree=*/false, *dl);
295300
cuf::populateCUFGPUToLLVMConversionPatterns(typeConverter, patterns);
296-
target.addIllegalOp<mlir::gpu::LaunchFuncOp>();
301+
302+
target.addDynamicallyLegalOp<mlir::gpu::LaunchFuncOp>(
303+
[&](mlir::gpu::LaunchFuncOp op) {
304+
if (op.getOperation()->getAttrOfType<cuf::ProcAttributeAttr>(
305+
cuf::getProcAttrName()))
306+
return false;
307+
return true;
308+
});
309+
297310
target.addIllegalOp<cuf::SharedMemoryOp>();
298311
target.addLegalDialect<mlir::LLVM::LLVMDialect>();
299312
if (mlir::failed(mlir::applyPartialConversion(getOperation(), target,

‎flang/lib/Optimizer/Transforms/CUFOpConversion.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -888,6 +888,11 @@ struct CUFLaunchOpConversion
888888
}
889889
if (procAttr)
890890
gpuLaunchOp->setAttr(cuf::getProcAttrName(), procAttr);
891+
else
892+
// Set default global attribute of the original was not found.
893+
gpuLaunchOp->setAttr(cuf::getProcAttrName(),
894+
cuf::ProcAttributeAttr::get(
895+
op.getContext(), cuf::ProcAttribute::Global));
891896
rewriter.replaceOp(op, gpuLaunchOp);
892897
return mlir::success();
893898
}

0 commit comments

Comments
 (0)