Skip to content

Commit ff97ec5

Browse files
sperytsperyt
committed
Merge from 'main' to 'sycl-web' (#1)
2 parents cf9e670 + 55a9615 commit ff97ec5

File tree

80 files changed

+4292
-676
lines changed

Some content is hidden

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

80 files changed

+4292
-676
lines changed

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

+13
Original file line numberDiff line numberDiff line change
@@ -196,6 +196,19 @@ TARGET_BUILTIN(__builtin_amdgcn_perm, "UiUiUiUi", "nc", "gfx8-insts")
196196

197197
TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
198198

199+
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f64, "dd*1d", "t", "gfx90a-insts")
200+
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f32, "ff*1f", "t", "gfx90a-insts")
201+
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2f16, "V2hV2h*1V2h", "t", "gfx90a-insts")
202+
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmin_f64, "dd*1d", "t", "gfx90a-insts")
203+
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmax_f64, "dd*1d", "t", "gfx90a-insts")
204+
205+
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f64, "dd*0d", "t", "gfx90a-insts")
206+
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmin_f64, "dd*0d", "t", "gfx90a-insts")
207+
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmax_f64, "dd*0d", "t", "gfx90a-insts")
208+
209+
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d", "t", "gfx90a-insts")
210+
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f", "t", "gfx8-insts")
211+
199212
//===----------------------------------------------------------------------===//
200213
// Deep learning builtins.
201214
//===----------------------------------------------------------------------===//

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

+38-28
Original file line numberDiff line numberDiff line change
@@ -274,7 +274,7 @@ class MigratorOpts<string base>
274274
// Args.hasArg(OPT_ffoo) can be used to check that the flag is enabled.
275275
// This is useful if the option is usually disabled.
276276
// Use this only when the option cannot be declared via BoolFOption.
277-
multiclass OptInFFlag<string name, string pos_prefix, string neg_prefix="",
277+
multiclass OptInCC1FFlag<string name, string pos_prefix, string neg_prefix="",
278278
string help="", list<OptionFlag> flags=[]> {
279279
def f#NAME : Flag<["-"], "f"#name>, Flags<[CC1Option] # flags>,
280280
Group<f_Group>, HelpText<pos_prefix # help>;
@@ -285,14 +285,35 @@ multiclass OptInFFlag<string name, string pos_prefix, string neg_prefix="",
285285
// A boolean option which is opt-out in CC1. The negative option exists in CC1 and
286286
// Args.hasArg(OPT_fno_foo) can be used to check that the flag is disabled.
287287
// Use this only when the option cannot be declared via BoolFOption.
288-
multiclass OptOutFFlag<string name, string pos_prefix, string neg_prefix,
288+
multiclass OptOutCC1FFlag<string name, string pos_prefix, string neg_prefix,
289289
string help="", list<OptionFlag> flags=[]> {
290290
def f#NAME : Flag<["-"], "f"#name>, Flags<flags>,
291291
Group<f_Group>, HelpText<pos_prefix # help>;
292292
def fno_#NAME : Flag<["-"], "fno-"#name>, Flags<[CC1Option] # flags>,
293293
Group<f_Group>, HelpText<neg_prefix # help>;
294294
}
295295

296+
// A boolean option which is opt-in in FC1. The positive option exists in FC1 and
297+
// Args.hasArg(OPT_ffoo) can be used to check that the flag is enabled.
298+
// This is useful if the option is usually disabled.
299+
multiclass OptInFC1FFlag<string name, string pos_prefix, string neg_prefix="",
300+
string help="", list<OptionFlag> flags=[]> {
301+
def f#NAME : Flag<["-"], "f"#name>, Flags<[FC1Option] # flags>,
302+
Group<f_Group>, HelpText<pos_prefix # help>;
303+
def fno_#NAME : Flag<["-"], "fno-"#name>, Flags<flags>,
304+
Group<f_Group>, HelpText<neg_prefix # help>;
305+
}
306+
307+
// A boolean option which is opt-out in FC1. The negative option exists in FC1 and
308+
// Args.hasArg(OPT_fno_foo) can be used to check that the flag is disabled.
309+
multiclass OptOutFC1FFlag<string name, string pos_prefix, string neg_prefix,
310+
string help="", list<OptionFlag> flags=[]> {
311+
def f#NAME : Flag<["-"], "f"#name>, Flags<flags>,
312+
Group<f_Group>, HelpText<pos_prefix # help>;
313+
def fno_#NAME : Flag<["-"], "fno-"#name>, Flags<[FC1Option] # flags>,
314+
Group<f_Group>, HelpText<neg_prefix # help>;
315+
}
316+
296317
// Creates a positive and negative flags where both of them are prefixed with
297318
// "m", have help text specified for positive and negative option, and a Group
298319
// optionally specified by the opt_group argument, otherwise Group<m_Group>.
@@ -1278,7 +1299,7 @@ defm addrsig : BoolFOption<"addrsig",
12781299
CodeGenOpts<"Addrsig">, DefaultFalse,
12791300
PosFlag<SetTrue, [CC1Option], "Emit">, NegFlag<SetFalse, [], "Don't emit">,
12801301
BothFlags<[CoreOption], " an address-significance table">>;
1281-
defm blocks : OptInFFlag<"blocks", "Enable the 'blocks' language feature", "", "", [CoreOption]>;
1302+
defm blocks : OptInCC1FFlag<"blocks", "Enable the 'blocks' language feature", "", "", [CoreOption]>;
12821303
def fbootclasspath_EQ : Joined<["-"], "fbootclasspath=">, Group<f_Group>;
12831304
defm borland_extensions : BoolFOption<"borland-extensions",
12841305
LangOpts<"Borland">, DefaultFalse,
@@ -1294,7 +1315,7 @@ def fclang_abi_compat_EQ : Joined<["-"], "fclang-abi-compat=">, Group<f_clang_Gr
12941315
Flags<[CC1Option]>, MetaVarName<"<version>">, Values<"<major>.<minor>,latest">,
12951316
HelpText<"Attempt to match the ABI of Clang <version>">;
12961317
def fclasspath_EQ : Joined<["-"], "fclasspath=">, Group<f_Group>;
1297-
defm color_diagnostics : OptInFFlag<"color-diagnostics", "Enable", "Disable", " colors in diagnostics",
1318+
defm color_diagnostics : OptInCC1FFlag<"color-diagnostics", "Enable", "Disable", " colors in diagnostics",
12981319
[CoreOption, FlangOption]>;
12991320
def fdiagnostics_color : Flag<["-"], "fdiagnostics-color">, Group<f_Group>,
13001321
Flags<[CoreOption, NoXarchOption]>;
@@ -1415,7 +1436,7 @@ def fno_elide_type : Flag<["-"], "fno-elide-type">, Group<f_Group>,
14151436
HelpText<"Do not elide types when printing diagnostics">,
14161437
MarshallingInfoNegativeFlag<DiagnosticOpts<"ElideType">>;
14171438
def feliminate_unused_debug_symbols : Flag<["-"], "feliminate-unused-debug-symbols">, Group<f_Group>;
1418-
defm eliminate_unused_debug_types : OptOutFFlag<"eliminate-unused-debug-types",
1439+
defm eliminate_unused_debug_types : OptOutCC1FFlag<"eliminate-unused-debug-types",
14191440
"Do not emit ", "Emit ", " debug info for defined but unused types">;
14201441
def femit_all_decls : Flag<["-"], "femit-all-decls">, Group<f_Group>, Flags<[CC1Option]>,
14211442
HelpText<"Emit all declarations, even if unused">,
@@ -1520,7 +1541,7 @@ defm cxx_static_destructors : BoolFOption<"c++-static-destructors",
15201541
def fsymbol_partition_EQ : Joined<["-"], "fsymbol-partition=">, Group<f_Group>,
15211542
Flags<[CC1Option]>, MarshallingInfoString<CodeGenOpts<"SymbolPartition">>;
15221543

1523-
defm memory_profile : OptInFFlag<"memory-profile", "Enable", "Disable", " heap memory profiling">;
1544+
defm memory_profile : OptInCC1FFlag<"memory-profile", "Enable", "Disable", " heap memory profiling">;
15241545
def fmemory_profile_EQ : Joined<["-"], "fmemory-profile=">,
15251546
Group<f_Group>, Flags<[CC1Option]>, MetaVarName<"<directory>">,
15261547
HelpText<"Enable heap memory profiling and dump results into <directory>">;
@@ -2164,9 +2185,9 @@ defm pch_instantiate_templates : BoolFOption<"pch-instantiate-templates",
21642185
LangOpts<"PCHInstantiateTemplates">, DefaultFalse,
21652186
PosFlag<SetTrue, [], "Instantiate templates already while building a PCH">,
21662187
NegFlag<SetFalse>, BothFlags<[CC1Option, CoreOption]>>;
2167-
defm pch_codegen: OptInFFlag<"pch-codegen", "Generate ", "Do not generate ",
2188+
defm pch_codegen: OptInCC1FFlag<"pch-codegen", "Generate ", "Do not generate ",
21682189
"code for uses of this PCH that assumes an explicit object file will be built for the PCH">;
2169-
defm pch_debuginfo: OptInFFlag<"pch-debuginfo", "Generate ", "Do not generate ",
2190+
defm pch_debuginfo: OptInCC1FFlag<"pch-debuginfo", "Generate ", "Do not generate ",
21702191
"debug info for types in an object file built from this PCH and do not generate them elsewhere">;
21712192

21722193
def fimplicit_module_maps : Flag <["-"], "fimplicit-module-maps">, Group<f_Group>,
@@ -4574,7 +4595,7 @@ def : Flag<["-"], "fsycl-explicit-simd">, Flags<[CoreOption]>, Group<clang_ignor
45744595
HelpText<"Enable SYCL explicit SIMD extension. (deprecated)">;
45754596
def : Flag<["-"], "fno-sycl-explicit-simd">, Flags<[CoreOption]>, Group<clang_ignored_legacy_options_Group>,
45764597
HelpText<"Disable SYCL explicit SIMD extension. (deprecated)">;
4577-
defm sycl_early_optimizations : OptOutFFlag<"sycl-early-optimizations", "Enable", "Disable", " standard optimization pipeline for SYCL device compiler", [CoreOption]>;
4598+
defm sycl_early_optimizations : OptOutCC1FFlag<"sycl-early-optimizations", "Enable", "Disable", " standard optimization pipeline for SYCL device compiler", [CoreOption]>;
45784599
def fsycl_dead_args_optimization : Flag<["-"], "fsycl-dead-args-optimization">,
45794600
Group<sycl_Group>, Flags<[NoArgumentUnused, CoreOption]>, HelpText<"Enables "
45804601
"elimination of DPC++ dead kernel arguments">;
@@ -4633,26 +4654,18 @@ def fdefault_real_8 : Flag<["-"],"fdefault-real-8">, Group<f_Group>,
46334654
HelpText<"Set the default real kind to an 8 byte wide type">;
46344655
def flarge_sizes : Flag<["-"],"flarge-sizes">, Group<f_Group>,
46354656
HelpText<"Use INTEGER(KIND=8) for the result type in size-related intrinsics">;
4636-
def fbackslash : Flag<["-"], "fbackslash">, Group<f_Group>,
4637-
HelpText<"Specify that backslash in string introduces an escape character">,
4638-
DocBrief<[{Change the interpretation of backslashes in string literals from
4639-
a single backslash character to "C-style" escape characters.}]>;
4640-
def fno_backslash : Flag<["-"], "fno-backslash">, Group<f_Group>;
4641-
def fxor_operator : Flag<["-"], "fxor-operator">, Group<f_Group>,
4642-
HelpText<"Enable .XOR. as a synonym of .NEQV.">;
4643-
def fno_xor_operator : Flag<["-"], "fno-xor-operator">, Group<f_Group>;
4644-
def flogical_abbreviations : Flag<["-"], "flogical-abbreviations">, Group<f_Group>,
4645-
HelpText<"Enable logical abbreviations">;
4646-
def fno_logical_abbreviations : Flag<["-"], "fno-logical-abbreviations">, Group<f_Group>;
4647-
def fimplicit_none : Flag<["-"], "fimplicit-none">, Group<f_Group>,
4648-
HelpText<"No implicit typing allowed unless overridden by IMPLICIT statements">;
4649-
def fno_implicit_none : Flag<["-"], "fno-implicit-none">, Group<f_Group>;
4657+
46504658
def falternative_parameter_statement : Flag<["-"], "falternative-parameter-statement">, Group<f_Group>,
46514659
HelpText<"Enable the old style PARAMETER statement">;
46524660
def fintrinsic_modules_path : Separate<["-"], "fintrinsic-modules-path">, Group<f_Group>, MetaVarName<"<dir>">,
46534661
HelpText<"Specify where to find the compiled intrinsic modules">,
46544662
DocBrief<[{This option specifies the location of pre-compiled intrinsic modules,
46554663
if they are not in the default location expected by the compiler.}]>;
4664+
4665+
defm backslash : OptInFC1FFlag<"backslash", "Specify that backslash in string introduces an escape character">;
4666+
defm xor_operator : OptInFC1FFlag<"xor-operator", "Enable .XOR. as a synonym of .NEQV.">;
4667+
defm logical_abbreviations : OptInFC1FFlag<"logical-abbreviations", "Enable logical abbreviations">;
4668+
defm implicit_none : OptInFC1FFlag<"implicit-none", "No implicit typing allowed unless overridden by IMPLICIT statements">;
46564669
}
46574670

46584671
def J : JoinedOrSeparate<["-"], "J">,
@@ -4707,13 +4720,10 @@ def fget_symbols_sources : Flag<["-"], "fget-symbols-sources">, Group<Action_Gro
47074720

47084721
def module_suffix : Separate<["-"], "module-suffix">, Group<f_Group>, MetaVarName<"<suffix>">,
47094722
HelpText<"Use <suffix> as the suffix for module files (the default value is `.mod`)">;
4710-
def fanalyzed_objects_for_unparse : Flag<["-"],
4711-
"fanalyzed-objects-for-unparse">, Group<f_Group>;
4712-
def fno_analyzed_objects_for_unparse : Flag<["-"],
4713-
"fno-analyzed-objects-for-unparse">, Group<f_Group>,
4714-
HelpText<"Do not use the analyzed objects when unparsing">;
47154723
def fno_reformat : Flag<["-"], "fno-reformat">, Group<Preprocessor_Group>,
47164724
HelpText<"Dump the cooked character stream in -E mode">;
4725+
defm analyzed_objects_for_unparse : OptOutFC1FFlag<"analyzed-objects-for-unparse", "", "Do not use the analyzed objects when unparsing">;
4726+
47174727
}
47184728

47194729
//===----------------------------------------------------------------------===//

‎clang/lib/CodeGen/CGBuiltin.cpp

+68
Original file line numberDiff line numberDiff line change
@@ -16203,6 +16203,74 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1620316203
Src0 = Builder.CreatePointerBitCastOrAddrSpaceCast(Src0, PTy);
1620416204
return Builder.CreateCall(F, { Src0, Src1, Src2, Src3, Src4 });
1620516205
}
16206+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:
16207+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:
16208+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
16209+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
16210+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
16211+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
16212+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
16213+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
16214+
Intrinsic::ID IID;
16215+
llvm::Type *ArgTy = llvm::Type::getDoubleTy(getLLVMContext());
16216+
switch (BuiltinID) {
16217+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:
16218+
ArgTy = llvm::Type::getFloatTy(getLLVMContext());
16219+
IID = Intrinsic::amdgcn_global_atomic_fadd;
16220+
break;
16221+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
16222+
ArgTy = llvm::FixedVectorType::get(
16223+
llvm::Type::getHalfTy(getLLVMContext()), 2);
16224+
IID = Intrinsic::amdgcn_global_atomic_fadd;
16225+
break;
16226+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:
16227+
IID = Intrinsic::amdgcn_global_atomic_fadd;
16228+
break;
16229+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
16230+
IID = Intrinsic::amdgcn_global_atomic_fmin;
16231+
break;
16232+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
16233+
IID = Intrinsic::amdgcn_global_atomic_fmax;
16234+
break;
16235+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
16236+
IID = Intrinsic::amdgcn_flat_atomic_fadd;
16237+
break;
16238+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
16239+
IID = Intrinsic::amdgcn_flat_atomic_fmin;
16240+
break;
16241+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64:
16242+
IID = Intrinsic::amdgcn_flat_atomic_fmax;
16243+
break;
16244+
}
16245+
llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
16246+
llvm::Value *Val = EmitScalarExpr(E->getArg(1));
16247+
llvm::Function *F =
16248+
CGM.getIntrinsic(IID, {ArgTy, Addr->getType(), Val->getType()});
16249+
return Builder.CreateCall(F, {Addr, Val});
16250+
}
16251+
case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
16252+
case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32: {
16253+
Intrinsic::ID IID;
16254+
llvm::Type *ArgTy;
16255+
switch (BuiltinID) {
16256+
case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:
16257+
ArgTy = llvm::Type::getFloatTy(getLLVMContext());
16258+
IID = Intrinsic::amdgcn_ds_fadd;
16259+
break;
16260+
case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
16261+
ArgTy = llvm::Type::getDoubleTy(getLLVMContext());
16262+
IID = Intrinsic::amdgcn_ds_fadd;
16263+
break;
16264+
}
16265+
llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
16266+
llvm::Value *Val = EmitScalarExpr(E->getArg(1));
16267+
llvm::Constant *ZeroI32 = llvm::ConstantInt::getIntegerValue(
16268+
llvm::Type::getInt32Ty(getLLVMContext()), APInt(32, 0, true));
16269+
llvm::Constant *ZeroI1 = llvm::ConstantInt::getIntegerValue(
16270+
llvm::Type::getInt1Ty(getLLVMContext()), APInt(1, 0));
16271+
llvm::Function *F = CGM.getIntrinsic(IID, {ArgTy});
16272+
return Builder.CreateCall(F, {Addr, Val, ZeroI32, ZeroI32, ZeroI1});
16273+
}
1620616274
case AMDGPU::BI__builtin_amdgcn_read_exec: {
1620716275
CallInst *CI = cast<CallInst>(
1620816276
EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, NormalRead, "exec"));

0 commit comments

Comments
 (0)