Skip to content

Commit f51a5ba

Browse files
zhaohuiw42zhaohuiw42
authored
Merge branch 'main' into fix-self-reference-checker
2 parents e52e53f + 9fefc01 commit f51a5ba

File tree

52 files changed

+947
-304
lines changed

Some content is hidden

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

52 files changed

+947
-304
lines changed

‎clang/docs/ClangFormatStyleOptions.rst‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2582,9 +2582,9 @@ the configuration (without a prefix: ``Auto``).
25822582

25832583
.. _BracedInitializerIndentWidth:
25842584

2585-
**BracedInitializerIndentWidth** (``Unsigned``) :versionbadge:`clang-format 17` :ref:`<BracedInitializerIndentWidth>`
2585+
**BracedInitializerIndentWidth** (``Integer``) :versionbadge:`clang-format 17` :ref:`<BracedInitializerIndentWidth>`
25862586
The number of columns to use to indent the contents of braced init lists.
2587-
If unset, ``ContinuationIndentWidth`` is used.
2587+
If unset or negative, ``ContinuationIndentWidth`` is used.
25882588

25892589
.. code-block:: c++
25902590

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

Lines changed: 23 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -544,29 +544,29 @@ TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12, "V4fiiV4f",
544544
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12, "V4fiiV4f", "nc", "gfx12-insts,wavefrontsize64")
545545
TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12, "V4iIbiIbiV4iIb", "nc", "gfx12-insts,wavefrontsize64")
546546

547-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32, "V8fV8hV16hV8fs", "nc", "gfx12-insts,wavefrontsize32")
548-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32, "V8fV8sV16sV8fs", "nc", "gfx12-insts,wavefrontsize32")
549-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32, "V8hV8hV16hV8hs", "nc", "gfx12-insts,wavefrontsize32")
550-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32, "V8sV8sV16sV8ss", "nc", "gfx12-insts,wavefrontsize32")
551-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32, "V8iIbV2iIbV4iV8isIb", "nc", "gfx12-insts,wavefrontsize32")
552-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32, "V8iIbiIbV2iV8isIb", "nc", "gfx12-insts,wavefrontsize32")
553-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32, "V8iIbV2iIbV4iV8isIb", "nc", "gfx12-insts,wavefrontsize32")
554-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32, "V8fV2iV4iV8fs", "nc", "gfx12-insts,wavefrontsize32")
555-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32, "V8fV2iV4iV8fs", "nc", "gfx12-insts,wavefrontsize32")
556-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32, "V8fV2iV4iV8fs", "nc", "gfx12-insts,wavefrontsize32")
557-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32, "V8fV2iV4iV8fs", "nc", "gfx12-insts,wavefrontsize32")
558-
559-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64, "V4fV4hV8hV4fs", "nc", "gfx12-insts,wavefrontsize64")
560-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64, "V4fV4sV8sV4fs", "nc", "gfx12-insts,wavefrontsize64")
561-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64, "V4hV4hV8hV4hs", "nc", "gfx12-insts,wavefrontsize64")
562-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64, "V4sV4sV8sV4ss", "nc", "gfx12-insts,wavefrontsize64")
563-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64, "V4iIbiIbV2iV4isIb", "nc", "gfx12-insts,wavefrontsize64")
564-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64, "V4iIbiIbiV4isIb", "nc", "gfx12-insts,wavefrontsize64")
565-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64, "V4iIbiIbV2iV4isIb", "nc", "gfx12-insts,wavefrontsize64")
566-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64, "V4fiV2iV4fs", "nc", "gfx12-insts,wavefrontsize64")
567-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64, "V4fiV2iV4fs", "nc", "gfx12-insts,wavefrontsize64")
568-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64, "V4fiV2iV4fs", "nc", "gfx12-insts,wavefrontsize64")
569-
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64, "V4fiV2iV4fs", "nc", "gfx12-insts,wavefrontsize64")
547+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32, "V8fV8hV16hV8fi", "nc", "gfx12-insts,wavefrontsize32")
548+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32, "V8fV8sV16sV8fi", "nc", "gfx12-insts,wavefrontsize32")
549+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32, "V8hV8hV16hV8hi", "nc", "gfx12-insts,wavefrontsize32")
550+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32, "V8sV8sV16sV8si", "nc", "gfx12-insts,wavefrontsize32")
551+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32, "V8iIbV2iIbV4iV8iiIb", "nc", "gfx12-insts,wavefrontsize32")
552+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32, "V8iIbiIbV2iV8iiIb", "nc", "gfx12-insts,wavefrontsize32")
553+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32, "V8iIbV2iIbV4iV8iiIb", "nc", "gfx12-insts,wavefrontsize32")
554+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32, "V8fV2iV4iV8fi", "nc", "gfx12-insts,wavefrontsize32")
555+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32, "V8fV2iV4iV8fi", "nc", "gfx12-insts,wavefrontsize32")
556+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32, "V8fV2iV4iV8fi", "nc", "gfx12-insts,wavefrontsize32")
557+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32, "V8fV2iV4iV8fi", "nc", "gfx12-insts,wavefrontsize32")
558+
559+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64, "V4fV4hV8hV4fi", "nc", "gfx12-insts,wavefrontsize64")
560+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64, "V4fV4sV8sV4fi", "nc", "gfx12-insts,wavefrontsize64")
561+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64, "V4hV4hV8hV4hi", "nc", "gfx12-insts,wavefrontsize64")
562+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64, "V4sV4sV8sV4si", "nc", "gfx12-insts,wavefrontsize64")
563+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64, "V4iIbiIbV2iV4iiIb", "nc", "gfx12-insts,wavefrontsize64")
564+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64, "V4iIbiIbiV4iiIb", "nc", "gfx12-insts,wavefrontsize64")
565+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64, "V4iIbiIbV2iV4iiIb", "nc", "gfx12-insts,wavefrontsize64")
566+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64, "V4fiV2iV4fi", "nc", "gfx12-insts,wavefrontsize64")
567+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64, "V4fiV2iV4fi", "nc", "gfx12-insts,wavefrontsize64")
568+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64, "V4fiV2iV4fi", "nc", "gfx12-insts,wavefrontsize64")
569+
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64, "V4fiV2iV4fi", "nc", "gfx12-insts,wavefrontsize64")
570570

571571
TARGET_BUILTIN(__builtin_amdgcn_prng_b32, "UiUi", "nc", "prng-inst")
572572
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_fp6_f16, "V6UiV32hf", "nc", "f16bf16-to-fp6bf6-cvt-scale-insts")

‎clang/include/clang/Basic/DiagnosticDriverKinds.td‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -563,6 +563,11 @@ def err_test_module_file_extension_format : Error<
563563
def err_drv_module_output_with_multiple_arch : Error<
564564
"option '-fmodule-output' cannot be used with multiple arch options">;
565565

566+
def err_drv_reduced_module_output_overrided : Warning<
567+
"the implicit output of reduced BMI may be overrided by the output file specified by '--precompile'. "
568+
"please consider use '-fmodule-output=' to specify the output file for reduced BMI explicitly">,
569+
InGroup<DiagGroup<"reduced-bmi-output-overrided">>;
570+
566571
def warn_drv_delayed_template_parsing_after_cxx20 : Warning<
567572
"-fdelayed-template-parsing is deprecated after C++20">,
568573
InGroup<DiagGroup<"delayed-template-parsing-in-cxx20">>;

‎clang/include/clang/Format/Format.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1289,7 +1289,7 @@ struct FormatStyle {
12891289
BitFieldColonSpacingStyle BitFieldColonSpacing;
12901290

12911291
/// The number of columns to use to indent the contents of braced init lists.
1292-
/// If unset, ``ContinuationIndentWidth`` is used.
1292+
/// If unset or negative, ``ContinuationIndentWidth`` is used.
12931293
/// \code
12941294
/// AlignAfterOpenBracket: AlwaysBreak
12951295
/// BracedInitializerIndentWidth: 2
@@ -1319,7 +1319,7 @@ struct FormatStyle {
13191319
/// }
13201320
/// \endcode
13211321
/// \version 17
1322-
std::optional<unsigned> BracedInitializerIndentWidth;
1322+
int BracedInitializerIndentWidth;
13231323

13241324
/// Different ways to wrap braces after control statements.
13251325
enum BraceWrappingAfterControlStatementStyle : int8_t {

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

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4250,10 +4250,18 @@ static bool RenderModulesOptions(Compilation &C, const Driver &D,
42504250

42514251
if (Args.hasArg(options::OPT_fmodule_output_EQ))
42524252
Args.AddLastArg(CmdArgs, options::OPT_fmodule_output_EQ);
4253-
else
4253+
else {
4254+
if (Args.hasArg(options::OPT__precompile) &&
4255+
(!Args.hasArg(options::OPT_o) ||
4256+
Args.getLastArg(options::OPT_o)->getValue() ==
4257+
getCXX20NamedModuleOutputPath(Args, Input.getBaseInput()))) {
4258+
D.Diag(diag::err_drv_reduced_module_output_overrided);
4259+
}
4260+
42544261
CmdArgs.push_back(Args.MakeArgString(
42554262
"-fmodule-output=" +
42564263
getCXX20NamedModuleOutputPath(Args, Input.getBaseInput())));
4264+
}
42574265
}
42584266

42594267
// Noop if we see '-fmodules-reduced-bmi' with other translation

‎clang/lib/Format/ContinuationIndenter.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1921,9 +1921,9 @@ void ContinuationIndenter::moveStatePastScopeOpener(LineState &State,
19211921
NewIndent = Style.IndentWidth +
19221922
std::min(State.Column, CurrentState.NestedBlockIndent);
19231923
} else if (Current.is(tok::l_brace)) {
1924-
NewIndent =
1925-
CurrentState.LastSpace + Style.BracedInitializerIndentWidth.value_or(
1926-
Style.ContinuationIndentWidth);
1924+
const auto Width = Style.BracedInitializerIndentWidth;
1925+
NewIndent = CurrentState.LastSpace +
1926+
(Width < 0 ? Style.ContinuationIndentWidth : Width);
19271927
} else {
19281928
NewIndent = CurrentState.LastSpace + Style.ContinuationIndentWidth;
19291929
}

‎clang/lib/Format/Format.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1512,7 +1512,7 @@ FormatStyle getLLVMStyle(FormatStyle::LanguageKind Language) {
15121512
LLVMStyle.BinPackLongBracedList = true;
15131513
LLVMStyle.BinPackParameters = FormatStyle::BPPS_BinPack;
15141514
LLVMStyle.BitFieldColonSpacing = FormatStyle::BFCS_Both;
1515-
LLVMStyle.BracedInitializerIndentWidth = std::nullopt;
1515+
LLVMStyle.BracedInitializerIndentWidth = -1;
15161516
LLVMStyle.BraceWrapping = {/*AfterCaseLabel=*/false,
15171517
/*AfterClass=*/false,
15181518
/*AfterControlStatement=*/FormatStyle::BWACS_Never,

‎clang/lib/Sema/SemaHLSL.cpp‎

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1446,18 +1446,20 @@ static bool DiagnoseLocalRegisterBinding(Sema &S, SourceLocation &ArgLoc,
14461446
Ty = Ty->getArrayElementTypeNoTypeQual();
14471447

14481448
// Basic types
1449-
if (Ty->isArithmeticType()) {
1449+
if (Ty->isArithmeticType() || Ty->isVectorType()) {
14501450
bool DeclaredInCOrTBuffer = isa<HLSLBufferDecl>(D->getDeclContext());
14511451
if (SpecifiedSpace && !DeclaredInCOrTBuffer)
14521452
S.Diag(ArgLoc, diag::err_hlsl_space_on_global_constant);
14531453

1454-
if (!DeclaredInCOrTBuffer &&
1455-
(Ty->isIntegralType(S.getASTContext()) || Ty->isFloatingType())) {
1456-
// Default Globals
1454+
if (!DeclaredInCOrTBuffer && (Ty->isIntegralType(S.getASTContext()) ||
1455+
Ty->isFloatingType() || Ty->isVectorType())) {
1456+
// Register annotation on default constant buffer declaration ($Globals)
14571457
if (RegType == RegisterType::CBuffer)
14581458
S.Diag(ArgLoc, diag::warn_hlsl_deprecated_register_type_b);
14591459
else if (RegType != RegisterType::C)
14601460
S.Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
1461+
else
1462+
return true;
14611463
} else {
14621464
if (RegType == RegisterType::C)
14631465
S.Diag(ArgLoc, diag::warn_hlsl_register_type_c_packoffset);

‎clang/test/AST/HLSL/resource_binding_attr.hlsl‎

Lines changed: 41 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -1,41 +1,56 @@
1-
// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -ast-dump -o - %s | FileCheck %s
1+
// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -finclude-default-header -ast-dump -o - %s | FileCheck %s
22

3-
// CHECK:HLSLBufferDecl 0x[[CB:[0-9a-f]+]] {{.*}} line:8:9 cbuffer CB
4-
// CHECK-NEXT:HLSLResourceClassAttr 0x[[CB:[0-9a-f]+]] {{.*}} Implicit CBuffer
5-
// CHECK-NEXT:HLSLResourceAttr 0x[[CB:[0-9a-f]+]] {{.*}} Implicit CBuffer
6-
// CHECK-NEXT:HLSLResourceBindingAttr 0x{{[0-9a-f]+}} <col:14> "b3" "space2"
7-
// CHECK-NEXT:VarDecl 0x[[A:[0-9a-f]+]] {{.*}} col:9 used a 'hlsl_constant float'
3+
// CHECK: HLSLBufferDecl {{.*}} line:[[# @LINE + 5]]:9 cbuffer CB
4+
// CHECK-NEXT: HLSLResourceClassAttr {{.*}} Implicit CBuffer
5+
// CHECK-NEXT: HLSLResourceAttr {{.*}} Implicit CBuffer
6+
// CHECK-NEXT: HLSLResourceBindingAttr {{.*}} "b3" "space2"
7+
// CHECK-NEXT: VarDecl {{.*}} used a 'hlsl_constant float'
88
cbuffer CB : register(b3, space2) {
99
float a;
1010
}
1111

12-
// CHECK:HLSLBufferDecl 0x[[TB:[0-9a-f]+]] {{.*}} line:17:9 tbuffer TB
13-
// CHECK-NEXT:HLSLResourceClassAttr 0x[[CB:[0-9a-f]+]] {{.*}} Implicit SRV
14-
// CHECK-NEXT:HLSLResourceAttr 0x[[CB:[0-9a-f]+]] {{.*}} Implicit TBuffer
15-
// CHECK-NEXT:HLSLResourceBindingAttr 0x{{[0-9a-f]+}} <col:14> "t2" "space1"
16-
// CHECK-NEXT:VarDecl 0x[[B:[0-9a-f]+]] {{.*}} col:9 used b 'hlsl_constant float'
12+
// CHECK: HLSLBufferDecl {{.*}} line:[[# @LINE + 5]]:9 tbuffer TB
13+
// CHECK-NEXT: HLSLResourceClassAttr {{.*}} Implicit SRV
14+
// CHECK-NEXT: HLSLResourceAttr {{.*}} Implicit TBuffer
15+
// CHECK-NEXT: HLSLResourceBindingAttr {{.*}} "t2" "space1"
16+
// CHECK-NEXT: VarDecl {{.*}} used b 'hlsl_constant float'
1717
tbuffer TB : register(t2, space1) {
1818
float b;
1919
}
2020

21-
float foo() {
22-
// CHECK: BinaryOperator 0x{{[0-9a-f]+}} <col:10, col:14> 'float' '+'
23-
// CHECK-NEXT: ImplicitCastExpr 0x{{[0-9a-f]+}} <col:10> 'float' <LValueToRValue>
24-
// CHECK-NEXT: DeclRefExpr 0x{{[0-9a-f]+}} <col:10> 'hlsl_constant float' lvalue Var 0x[[A]] 'a' 'hlsl_constant float'
25-
// CHECK-NEXT: ImplicitCastExpr 0x{{[0-9a-f]+}} <col:14> 'float' <LValueToRValue>
26-
// CHECK-NEXT: DeclRefExpr 0x{{[0-9a-f]+}} <col:14> 'hlsl_constant float' lvalue Var 0x[[B]] 'b' 'hlsl_constant float'
21+
export float foo() {
2722
return a + b;
2823
}
2924

30-
// CHECK: VarDecl 0x{{[0-9a-f]+}} <{{.*}}> col:17 UAV 'RWBuffer<float>':'hlsl::RWBuffer<float>' callinit
31-
// CHECK-NEXT:-CXXConstructExpr 0x{{[0-9a-f]+}} <col:17> 'RWBuffer<float>':'hlsl::RWBuffer<float>' 'void ()'
32-
// CHECK-NEXT:-HLSLResourceBindingAttr 0x{{[0-9a-f]+}} <col:23> "u3" "space0"
25+
// CHECK: VarDecl {{.*}} UAV 'RWBuffer<float>':'hlsl::RWBuffer<float>'
26+
// CHECK: HLSLResourceBindingAttr {{.*}} "u3" "space0"
3327
RWBuffer<float> UAV : register(u3);
3428

35-
// CHECK: -VarDecl 0x{{[0-9a-f]+}} <{{.*}}> col:17 UAV1 'RWBuffer<float>':'hlsl::RWBuffer<float>' callinit
36-
// CHECK-NEXT:-CXXConstructExpr 0x{{[0-9a-f]+}} <col:17> 'RWBuffer<float>':'hlsl::RWBuffer<float>' 'void ()'
37-
// CHECK-NEXT:-HLSLResourceBindingAttr 0x{{[0-9a-f]+}} <col:24> "u2" "space0"
38-
// CHECK-NEXT:-VarDecl 0x{{[0-9a-f]+}} <col:1, col:38> col:38 UAV2 'RWBuffer<float>':'hlsl::RWBuffer<float>' callinit
39-
// CHECK-NEXT:-CXXConstructExpr 0x{{[0-9a-f]+}} <col:38> 'RWBuffer<float>':'hlsl::RWBuffer<float>' 'void ()'
40-
// CHECK-NEXT:-HLSLResourceBindingAttr 0x{{[0-9a-f]+}} <col:45> "u4" "space0"
29+
// CHECK: VarDecl {{.*}} UAV1 'RWBuffer<float>':'hlsl::RWBuffer<float>'
30+
// CHECK: HLSLResourceBindingAttr {{.*}} "u2" "space0"
31+
// CHECK: VarDecl {{.*}} UAV2 'RWBuffer<float>':'hlsl::RWBuffer<float>'
32+
// CHECK: HLSLResourceBindingAttr {{.*}} "u4" "space0"
4133
RWBuffer<float> UAV1 : register(u2), UAV2 : register(u4);
34+
35+
//
36+
// Default constants ($Globals) layout annotations
37+
38+
// CHECK: VarDecl {{.*}} f 'hlsl_constant float'
39+
// CHECK: HLSLResourceBindingAttr {{.*}} "c5" "space0"
40+
float f : register(c5);
41+
42+
// CHECK: VarDecl {{.*}} intv 'hlsl_constant int4':'vector<int hlsl_constant, 4>'
43+
// CHECK: HLSLResourceBindingAttr {{.*}} "c2" "space0"
44+
int4 intv : register(c2);
45+
46+
// CHECK: VarDecl {{.*}} dar 'hlsl_constant double[5]'
47+
// CHECK: HLSLResourceBindingAttr {{.*}} "c3" "space0"
48+
double dar[5] : register(c3);
49+
50+
struct S {
51+
int a;
52+
};
53+
54+
// CHECK: VarDecl {{.*}} s 'hlsl_constant S'
55+
// CHECK: HLSLResourceBindingAttr {{.*}} "c10" "space0
56+
S s : register(c10);

0 commit comments

Comments
 (0)