16 files changed, 502 insertions, 34 deletions
diff --git a/clang/test/AST/ast-print-record-decl.c b/clang/test/AST/ast-print-record-decl.c
index d3717a4..fd81588 100644
--- a/clang/test/AST/ast-print-record-decl.c
+++ b/clang/test/AST/ast-print-record-decl.c
@@ -290,9 +290,9 @@ KW DeclGroupInMemberList {
 // A tag decl group in the tag decl's own member list is exercised in
 // defSelfRef above.
 
+#ifdef __cplusplus
 
 // Check out-of-line record definition
-#ifdef __cplusplus
 // PRINT-CXX-NEXT: [[KW]] OutOfLineRecord {
 KW OutOfLineRecord {
   // PRINT-CXX-NEXT: [[KW]] Inner
@@ -304,4 +304,15 @@ KW OutOfLineRecord {
 KW OutOfLineRecord::Inner {
   // PRINT-CXX-NEXT: };
 };
+
+// PRINT-CXX-NEXT: template <typename, typename = int> [[KW]] SmearedTypeDefArgs;
+template <typename, typename = int> KW SmearedTypeDefArgs;
+// PRINT-CXX-NEXT: template <typename = int, typename> [[KW]] SmearedTypeDefArgs;
+template <typename = int, typename> KW SmearedTypeDefArgs;
+
+// PRINT-CXX-NEXT: template <int, int = 0> [[KW]] SmearedNTTPDefArgs;
+template <int, int = 0> KW SmearedNTTPDefArgs;
+// PRINT-CXX-NEXT: template <int = 0, int> [[KW]] SmearedNTTPDefArgs;
+template <int = 0, int> KW SmearedNTTPDefArgs;
+
 #endif
diff --git a/clang/test/CodeGen/X86/avx512vl-builtins.c b/clang/test/CodeGen/X86/avx512vl-builtins.c
index 8800623..6d91870 100644
--- a/clang/test/CodeGen/X86/avx512vl-builtins.c
+++ b/clang/test/CodeGen/X86/avx512vl-builtins.c
@@ -9272,6 +9272,8 @@ __m128i test_mm_cvtepi32_epi8(__m128i __A) {
   return _mm_cvtepi32_epi8(__A); 
 }
 
+TEST_CONSTEXPR(match_v16qi(_mm_cvtepi32_epi8((__m128i)(__v4si){1, 2, 3, 4}), 1 ,2, 3, 4, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+
 __m128i test_mm_mask_cvtepi32_epi8(__m128i __O, __mmask8 __M, __m128i __A) {
   // CHECK-LABEL: test_mm_mask_cvtepi32_epi8
   // CHECK: @llvm.x86.avx512.mask.pmov.db.128
@@ -9297,6 +9299,8 @@ __m128i test_mm256_cvtepi32_epi8(__m256i __A) {
   return _mm256_cvtepi32_epi8(__A); 
 }
 
+TEST_CONSTEXPR(match_v16qi(_mm256_cvtepi32_epi8((__m256i)(__v8si){1, 2, 3, 4, 5, 6, 7, 8}), 1, 2, 3, 4, 5, 6, 7, 8, 0, 0, 0, 0, 0, 0, 0 ,0));
+
 __m128i test_mm256_mask_cvtepi32_epi8(__m128i __O, __mmask8 __M, __m256i __A) {
   // CHECK-LABEL: test_mm256_mask_cvtepi32_epi8
   // CHECK: @llvm.x86.avx512.mask.pmov.db.256
@@ -9322,6 +9326,8 @@ __m128i test_mm_cvtepi32_epi16(__m128i __A) {
   return _mm_cvtepi32_epi16(__A); 
 }
 
+TEST_CONSTEXPR(match_v8hi(_mm_cvtepi32_epi16((__m128i)(__v4si){1, 2, 3, 4}), 1 ,2, 3, 4, 0, 0, 0, 0));
+
 __m128i test_mm_mask_cvtepi32_epi16(__m128i __O, __mmask8 __M, __m128i __A) {
   // CHECK-LABEL: test_mm_mask_cvtepi32_epi16
   // CHECK: @llvm.x86.avx512.mask.pmov.dw.128
@@ -9346,6 +9352,8 @@ __m128i test_mm256_cvtepi32_epi16(__m256i __A) {
   return _mm256_cvtepi32_epi16(__A); 
 }
 
+TEST_CONSTEXPR(match_v8hi(_mm256_cvtepi32_epi16((__m256i)(__v8si){1, 2, 3, 4, 5, 6, 7, 8}), 1, 2, 3, 4, 5, 6, 7, 8));
+
 __m128i test_mm256_mask_cvtepi32_epi16(__m128i __O, __mmask8 __M, __m256i __A) {
   // CHECK-LABEL: test_mm256_mask_cvtepi32_epi16
   // CHECK: @llvm.x86.avx512.mask.pmov.dw.256
@@ -9371,6 +9379,8 @@ __m128i test_mm_cvtepi64_epi8(__m128i __A) {
   return _mm_cvtepi64_epi8(__A); 
 }
 
+TEST_CONSTEXPR(match_v16qi(_mm_cvtepi64_epi8((__m128i)(__v2di){1, 2}), 1, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+
 __m128i test_mm_mask_cvtepi64_epi8(__m128i __O, __mmask8 __M, __m128i __A) {
   // CHECK-LABEL: test_mm_mask_cvtepi64_epi8
   // CHECK: @llvm.x86.avx512.mask.pmov.qb.128
@@ -9396,6 +9406,8 @@ __m128i test_mm256_cvtepi64_epi8(__m256i __A) {
   return _mm256_cvtepi64_epi8(__A); 
 }
 
+TEST_CONSTEXPR(match_v16qi(_mm256_cvtepi64_epi8((__m256i)(__v4di){1, 2, 3, 4}), 1, 2, 3, 4, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+
 __m128i test_mm256_mask_cvtepi64_epi8(__m128i __O, __mmask8 __M, __m256i __A) {
   // CHECK-LABEL: test_mm256_mask_cvtepi64_epi8
   // CHECK: @llvm.x86.avx512.mask.pmov.qb.256
@@ -9421,6 +9433,8 @@ __m128i test_mm_cvtepi64_epi32(__m128i __A) {
   return _mm_cvtepi64_epi32(__A); 
 }
 
+TEST_CONSTEXPR(match_v4si(_mm_cvtepi64_epi32((__m128i)(__v2di){1, 2}),1, 2, 0, 0));
+
 __m128i test_mm_mask_cvtepi64_epi32(__m128i __O, __mmask8 __M, __m128i __A) {
   // CHECK-LABEL: test_mm_mask_cvtepi64_epi32
   // CHECK: @llvm.x86.avx512.mask.pmov.qd.128
@@ -9445,6 +9459,8 @@ __m128i test_mm256_cvtepi64_epi32(__m256i __A) {
   return _mm256_cvtepi64_epi32(__A); 
 }
 
+TEST_CONSTEXPR(match_v4si(_mm256_cvtepi64_epi32((__m256i)(__v4di){1 ,2 ,3 ,4}), 1, 2, 3, 4));
+
 __m128i test_mm256_mask_cvtepi64_epi32(__m128i __O, __mmask8 __M, __m256i __A) {
   // CHECK-LABEL: test_mm256_mask_cvtepi64_epi32
   // CHECK: trunc <4 x i64> %{{.*}} to <4 x i32>
@@ -9452,6 +9468,8 @@ __m128i test_mm256_mask_cvtepi64_epi32(__m128i __O, __mmask8 __M, __m256i __A) {
   return _mm256_mask_cvtepi64_epi32(__O, __M, __A); 
 }
 
+TEST_CONSTEXPR(match_v4si(_mm256_mask_cvtepi64_epi32(_mm_set1_epi32(-777), 0xA,(__m256i)(__v4di){1, -2, 3, -4}), -777, -2, -777, -4));
+
 __m128i test_mm256_maskz_cvtepi64_epi32(__mmask8 __M, __m256i __A) {
   // CHECK-LABEL: test_mm256_maskz_cvtepi64_epi32
   // CHECK: trunc <4 x i64> %{{.*}} to <4 x i32>
@@ -9459,6 +9477,8 @@ __m128i test_mm256_maskz_cvtepi64_epi32(__mmask8 __M, __m256i __A) {
   return _mm256_maskz_cvtepi64_epi32(__M, __A); 
 }
 
+TEST_CONSTEXPR(match_v4si(_mm256_maskz_cvtepi64_epi32( 0xA,(__m256i)(__v4di){1, -2, 3, -4}),0 , -2, 0, -4));
+
 void test_mm256_mask_cvtepi64_storeu_epi32(void * __P, __mmask8 __M, __m256i __A) {
   // CHECK-LABEL: test_mm256_mask_cvtepi64_storeu_epi32
   // CHECK: @llvm.x86.avx512.mask.pmov.qd.mem.256
@@ -9472,6 +9492,8 @@ __m128i test_mm_cvtepi64_epi16(__m128i __A) {
   return _mm_cvtepi64_epi16(__A); 
 }
 
+TEST_CONSTEXPR(match_v8hi(_mm_cvtepi64_epi16((__m128i)(__v2di){1, 2}),1, 2, 0, 0, 0, 0, 0, 0));
+
 __m128i test_mm_mask_cvtepi64_epi16(__m128i __O, __mmask8 __M, __m128i __A) {
   // CHECK-LABEL: test_mm_mask_cvtepi64_epi16
   // CHECK: @llvm.x86.avx512.mask.pmov.qw.128
@@ -9497,6 +9519,8 @@ __m128i test_mm256_cvtepi64_epi16(__m256i __A) {
   return _mm256_cvtepi64_epi16(__A); 
 }
 
+TEST_CONSTEXPR(match_v8hi(_mm256_cvtepi64_epi16((__m256i)(__v4di){1 ,2, 3, 4}),1, 2, 3, 4, 0, 0, 0, 0));
+
 __m128i test_mm256_mask_cvtepi64_epi16(__m128i __O, __mmask8 __M, __m256i __A) {
   // CHECK-LABEL: test_mm256_mask_cvtepi64_epi16
   // CHECK: @llvm.x86.avx512.mask.pmov.qw.256
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index f994adb..e3be262 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -43,6 +43,12 @@
 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_120a -target-feature +ptx86 -DPTX=86 \
 // RUN:            -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX86_SM120a %s
+// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_103a -target-feature +ptx87 -DPTX=87 \
+// RUN:            -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
+// RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX87_SM103a %s
+// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_100a -target-feature +ptx87 -DPTX=87 \
+// RUN:            -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
+// RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX87_SM100a %s
 // ###  The last run to check with the highest SM and PTX version available
 // ###  to make sure target builtins are still accepted.
 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_120a -target-feature +ptx87 -DPTX=87 \
@@ -1203,6 +1209,123 @@ __device__ void nvvm_cvt_sm100a_sm101a_sm120a() {
   // CHECK: ret void
 }
 
+__device__ void nvvm_cvt_sm100a_sm103a() {
+#if (PTX >= 87) && (__CUDA_ARCH_FEAT_SM100_ALL || __CUDA_ARCH_FEAT_SM103_ALL)
+  
+  typedef __fp16 f16x2 __attribute__((ext_vector_type(2)));
+  typedef __bf16 bf16x2 __attribute__((ext_vector_type(2)));
+  typedef char uint8x4 __attribute__((ext_vector_type(4)));
+
+// CHECK_PTX87_SM100a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: store <2 x half> %[[R1]], ptr %r1
+// CHECK_PTX87_SM103a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: store <2 x half> %[[R1]], ptr %r1
+  f16x2 r1 =  __nvvm_ff2f16x2_rs(1.0f, 1.0f, 0);
+  
+// CHECK_PTX87_SM100a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: store <2 x half> %[[R2]], ptr %r2
+// CHECK_PTX87_SM103a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: store <2 x half> %[[R2]], ptr %r2
+  f16x2 r2 =  __nvvm_ff2f16x2_rs_relu(1.0f, 1.0f, 0);
+  
+// CHECK_PTX87_SM100a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: store <2 x half> %[[R3]], ptr %r3
+// CHECK_PTX87_SM103a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: store <2 x half> %[[R3]], ptr %r3
+  f16x2 r3 =  __nvvm_ff2f16x2_rs_satfinite(1.0f, 1.0f, 0);
+
+// CHECK_PTX87_SM100a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: store <2 x half> %[[R4]], ptr %r4
+// CHECK_PTX87_SM103a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: store <2 x half> %[[R4]], ptr %r4
+  f16x2 r4 =  __nvvm_ff2f16x2_rs_relu_satfinite(1.0f, 1.0f, 0);
+
+// CHECK_PTX87_SM100a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R5]], ptr %r5
+// CHECK_PTX87_SM103a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R5]], ptr %r5
+  bf16x2 r5 =  __nvvm_ff2bf16x2_rs(1.0f, 1.0f, 0);
+
+// CHECK_PTX87_SM100a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R6]], ptr %r6
+// CHECK_PTX87_SM103a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R6]], ptr %r6
+  bf16x2 r6 =  __nvvm_ff2bf16x2_rs_relu(1.0f, 1.0f, 0);
+
+// CHECK_PTX87_SM100a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R7]], ptr %r7
+// CHECK_PTX87_SM103a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R7]], ptr %r7
+  bf16x2 r7 =  __nvvm_ff2bf16x2_rs_satfinite(1.0f, 1.0f, 0);
+
+// CHECK_PTX87_SM100a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R8]], ptr %r8
+// CHECK_PTX87_SM103a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R8]], ptr %r8
+  bf16x2 r8 =  __nvvm_ff2bf16x2_rs_relu_satfinite(1.0f, 1.0f, 0);
+
+// CHECK_PTX87_SM100a: %[[R9:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store <4 x i8> %[[R9]], ptr %r9
+// CHECK_PTX87_SM103a: %[[R9:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store <4 x i8> %[[R9]], ptr %r9
+  uint8x4 r9 =  __nvvm_f32x4_to_e4m3x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);
+
+// CHECK_PTX87_SM100a: %[[R10:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store <4 x i8> %[[R10]], ptr %r10
+// CHECK_PTX87_SM103a: %[[R10:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store <4 x i8> %[[R10]], ptr %r10
+  uint8x4 r10 =  __nvvm_f32x4_to_e4m3x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);
+
+// CHECK_PTX87_SM100a: %[[R11:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store <4 x i8> %[[R11]], ptr %r11
+// CHECK_PTX87_SM103a: %[[R11:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store <4 x i8> %[[R11]], ptr %r11
+  uint8x4 r11 =  __nvvm_f32x4_to_e5m2x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);
+
+// CHECK_PTX87_SM100a: %[[R12:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store <4 x i8> %[[R12]], ptr %r12
+// CHECK_PTX87_SM103a: %[[R12:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store <4 x i8> %[[R12]], ptr %r12
+  uint8x4 r12 =  __nvvm_f32x4_to_e5m2x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);
+
+// CHECK_PTX87_SM100a: %[[R13:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store <4 x i8> %[[R13]], ptr %r13
+// CHECK_PTX87_SM103a: %[[R13:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store <4 x i8> %[[R13]], ptr %r13
+  uint8x4 r13 = __nvvm_f32x4_to_e2m3x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);  
+
+// CHECK_PTX87_SM100a: %[[R14:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store <4 x i8> %[[R14]], ptr %r14
+// CHECK_PTX87_SM103a: %[[R14:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store <4 x i8> %[[R14]], ptr %r14
+  uint8x4 r14 = __nvvm_f32x4_to_e2m3x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);
+
+// CHECK_PTX87_SM100a: %[[R15:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store <4 x i8> %[[R15]], ptr %r15
+// CHECK_PTX87_SM103a: %[[R15:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store <4 x i8> %[[R15]], ptr %r15
+  uint8x4 r15 = __nvvm_f32x4_to_e3m2x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);
+
+// CHECK_PTX87_SM100a: %[[R16:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store <4 x i8> %[[R16]], ptr %r16
+// CHECK_PTX87_SM103a: %[[R16:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store <4 x i8> %[[R16]], ptr %r16
+  uint8x4 r16 = __nvvm_f32x4_to_e3m2x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);
+
+// CHECK_PTX87_SM100a: %[[R17:.*]] = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store i16 %[[R17]], ptr %r17
+// CHECK_PTX87_SM103a: %[[R17:.*]] = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store i16 %[[R17]], ptr %r17
+  short r17 = __nvvm_f32x4_to_e2m1x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);
+
+// CHECK_PTX87_SM100a: %[[R18:.*]] = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM100a: store i16 %[[R18]], ptr %r18
+// CHECK_PTX87_SM103a: %[[R18:.*]] = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)
+// CHECK_PTX87_SM103a: store i16 %[[R18]], ptr %r18
+  short r18 = __nvvm_f32x4_to_e2m1x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);
+#endif
+}
+
 #define NAN32 0x7FBFFFFF
 #define NAN16 (__bf16)0x7FBF
 #define BF16 (__bf16)0.1f
diff --git a/clang/test/CodeGenOpenCL/amdgpu-features.cl b/clang/test/CodeGenOpenCL/amdgpu-features.cl
index af1ef64..c0c22bc 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-features.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-features.cl
@@ -109,8 +109,8 @@
 // GFX1153: "target-features"="+16-bit-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f32,+ci-insts,+dl-insts,+dot10-insts,+dot12-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
 // GFX1200: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
 // GFX1201: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
-// GFX1250: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+wavefrontsize32"
-// GFX1251: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+wavefrontsize32"
+// GFX1250: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bitop3-insts,+ci-insts,+cluster,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+wavefrontsize32"
+// GFX1251: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bitop3-insts,+ci-insts,+cluster,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+wavefrontsize32"
 
 // GFX1103-W64: "target-features"="+16-bit-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f32,+ci-insts,+dl-insts,+dot10-insts,+dot12-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize64"
 
diff --git a/clang/test/Headers/arm-acle-header.c b/clang/test/Headers/arm-acle-header.c
index fea8472..58fcc66 100644
--- a/clang/test/Headers/arm-acle-header.c
+++ b/clang/test/Headers/arm-acle-header.c
@@ -10,6 +10,8 @@
 // RUN: %clang_cc1 -x c++ -triple arm64ec-windows -target-cpu cortex-a53 -fsyntax-only -ffreestanding -fms-extensions -fms-compatibility -fms-compatibility-version=19.11 %s
 // expected-no-diagnostics
 
+#include "system_reserved_names.h"
+
 #include <arm_acle.h>
 #ifdef _MSC_VER
 #include <intrin.h>
diff --git a/clang/test/Headers/arm-cde-header.c b/clang/test/Headers/arm-cde-header.c
index 1f60368..526202a 100644
--- a/clang/test/Headers/arm-cde-header.c
+++ b/clang/test/Headers/arm-cde-header.c
@@ -9,5 +9,7 @@
 
 // Check that the headers don't conflict with each other
 
+#include "system_reserved_names.h"
+
 #include <arm_cde.h>
 #include <arm_mve.h>
diff --git a/clang/test/Headers/arm-cmse-header.c b/clang/test/Headers/arm-cmse-header.c
index 862572d..c21c1ff 100644
--- a/clang/test/Headers/arm-cmse-header.c
+++ b/clang/test/Headers/arm-cmse-header.c
@@ -2,6 +2,8 @@
 // RUN: %clang_cc1 -triple thumbv8m.base-eabi  -fsyntax-only -ffreestanding -x c++ %s -verify -mcmse
 // expected-no-diagnostics
 
+#include "system_reserved_names.h"
+
 #include <arm_cmse.h>
 
 typedef void (*callback_t)(void);
diff --git a/clang/test/Headers/arm-fp16-header.c b/clang/test/Headers/arm-fp16-header.c
index b1a87fa..e472654 100644
--- a/clang/test/Headers/arm-fp16-header.c
+++ b/clang/test/Headers/arm-fp16-header.c
@@ -18,4 +18,6 @@
 
 // REQUIRES: aarch64-registered-target || arm-registered-target
 
+#include "system_reserved_names.h"
+
 #include <arm_fp16.h>
diff --git a/clang/test/Headers/arm-neon-header.c b/clang/test/Headers/arm-neon-header.c
index 89bd5aa..43b1b35 100644
--- a/clang/test/Headers/arm-neon-header.c
+++ b/clang/test/Headers/arm-neon-header.c
@@ -26,4 +26,6 @@
 
 // REQUIRES: aarch64-registered-target || arm-registered-target
 
+#include "system_reserved_names.h"
+
 #include <arm_neon.h>
diff --git a/clang/test/Headers/system_reserved_names.h b/clang/test/Headers/system_reserved_names.h
new file mode 100644
index 0000000..1a53f4f
--- /dev/null
+++ b/clang/test/Headers/system_reserved_names.h
@@ -0,0 +1,165 @@
+// Test that headers are not tripped up by the surrounding code defining various
+// alphabetic macros. Also ensure that we don't swallow the definition of user
+// provided macros (in other words, ensure that we push/pop correctly everywhere).
+//
+// The contents of this header is a lightly trimmed version of
+// libcxx/test/libcxx/system_reserved_names.gen.py; additions to that testcase
+// can be synced into this header as well.
+
+#define SYSTEM_RESERVED_NAME This name should not be used in Clang headers
+
+// libc++ does not use single-letter names as a matter of principle.
+// But Windows' own <wchar.h>, <math.h>, and <exception> use many of these
+// (at least C,E,F,I,M,N,P,S,X,Y,Z) as uglified function parameter names,
+// so don't define these on Windows.
+//
+#ifndef _WIN32
+#define _A SYSTEM_RESERVED_NAME
+#define _B SYSTEM_RESERVED_NAME
+#define _C SYSTEM_RESERVED_NAME
+#define _D SYSTEM_RESERVED_NAME
+#define _E SYSTEM_RESERVED_NAME
+#define _F SYSTEM_RESERVED_NAME
+#define _G SYSTEM_RESERVED_NAME
+#define _H SYSTEM_RESERVED_NAME
+#define _I SYSTEM_RESERVED_NAME
+#define _J SYSTEM_RESERVED_NAME
+#define _K SYSTEM_RESERVED_NAME
+#define _L SYSTEM_RESERVED_NAME
+#define _M SYSTEM_RESERVED_NAME
+#define _N SYSTEM_RESERVED_NAME
+#define _O SYSTEM_RESERVED_NAME
+#define _P SYSTEM_RESERVED_NAME
+#define _Q SYSTEM_RESERVED_NAME
+#define _R SYSTEM_RESERVED_NAME
+#define _S SYSTEM_RESERVED_NAME
+#define _T SYSTEM_RESERVED_NAME
+#define _U SYSTEM_RESERVED_NAME
+#define _V SYSTEM_RESERVED_NAME
+#define _W SYSTEM_RESERVED_NAME
+#define _X SYSTEM_RESERVED_NAME
+#define _Y SYSTEM_RESERVED_NAME
+#define _Z SYSTEM_RESERVED_NAME
+#endif
+
+// FreeBSD's <sys/types.h> uses _M
+//
+#ifdef __FreeBSD__
+# undef _M
+#endif
+
+// Test that libc++ doesn't use names that collide with FreeBSD system macros.
+// newlib and picolibc also define these macros
+#if !defined(__FreeBSD__) && !defined(_NEWLIB_VERSION)
+#  define __null_sentinel SYSTEM_RESERVED_NAME
+#  define __generic SYSTEM_RESERVED_NAME
+#endif
+
+// tchar.h defines these macros on Windows
+#ifndef _WIN32
+# define _UI   SYSTEM_RESERVED_NAME
+# define _PUC  SYSTEM_RESERVED_NAME
+# define _CPUC SYSTEM_RESERVED_NAME
+# define _PC   SYSTEM_RESERVED_NAME
+# define _CRPC SYSTEM_RESERVED_NAME
+# define _CPC  SYSTEM_RESERVED_NAME
+#endif
+
+// yvals.h on MINGW defines this macro
+#ifndef _WIN32
+# define _C2 SYSTEM_RESERVED_NAME
+#endif
+
+// Test that libc++ doesn't use names that collide with Win32 API macros.
+// Obviously we can only define these on non-Windows platforms.
+#ifndef _WIN32
+# define __allocator SYSTEM_RESERVED_NAME
+# define __bound SYSTEM_RESERVED_NAME
+# define __deallocate SYSTEM_RESERVED_NAME
+# define __deref SYSTEM_RESERVED_NAME
+# define __format_string SYSTEM_RESERVED_NAME
+# define __full SYSTEM_RESERVED_NAME
+# define __in SYSTEM_RESERVED_NAME
+# define __inout SYSTEM_RESERVED_NAME
+# define __nz SYSTEM_RESERVED_NAME
+# define __out SYSTEM_RESERVED_NAME
+# define __part SYSTEM_RESERVED_NAME
+# define __post SYSTEM_RESERVED_NAME
+# define __pre SYSTEM_RESERVED_NAME
+#endif
+
+// Newlib & picolibc use __input as a parameter name of a64l & l64a
+#ifndef _NEWLIB_VERSION
+# define __input SYSTEM_RESERVED_NAME
+#endif
+#define __output SYSTEM_RESERVED_NAME
+
+#define __acquire SYSTEM_RESERVED_NAME
+#define __release SYSTEM_RESERVED_NAME
+
+// Android and FreeBSD use this for __attribute__((__unused__))
+#if !defined(__FreeBSD__)  && !defined(__ANDROID__)
+#define __unused SYSTEM_RESERVED_NAME
+#endif
+
+// These names are not reserved, so the user can macro-define them.
+// These are intended to find improperly _Uglified template parameters.
+#define A SYSTEM_RESERVED_NAME
+#define Arg SYSTEM_RESERVED_NAME
+#define Args SYSTEM_RESERVED_NAME
+#define As SYSTEM_RESERVED_NAME
+#define B SYSTEM_RESERVED_NAME
+#define Bs SYSTEM_RESERVED_NAME
+#define C SYSTEM_RESERVED_NAME
+#define Cp SYSTEM_RESERVED_NAME
+#define Cs SYSTEM_RESERVED_NAME
+// Windows setjmp.h contains a struct member named 'D' on ARM/AArch64.
+#ifndef _WIN32
+# define D SYSTEM_RESERVED_NAME
+#endif
+#define Dp SYSTEM_RESERVED_NAME
+#define Ds SYSTEM_RESERVED_NAME
+#define E SYSTEM_RESERVED_NAME
+#define Ep SYSTEM_RESERVED_NAME
+#define Es SYSTEM_RESERVED_NAME
+#define N SYSTEM_RESERVED_NAME
+#define Np SYSTEM_RESERVED_NAME
+#define Ns SYSTEM_RESERVED_NAME
+#define R SYSTEM_RESERVED_NAME
+#define Rp SYSTEM_RESERVED_NAME
+#define Rs SYSTEM_RESERVED_NAME
+#define T SYSTEM_RESERVED_NAME
+#define Tp SYSTEM_RESERVED_NAME
+#define Ts SYSTEM_RESERVED_NAME
+#define Type SYSTEM_RESERVED_NAME
+#define Types SYSTEM_RESERVED_NAME
+#define U SYSTEM_RESERVED_NAME
+#define Up SYSTEM_RESERVED_NAME
+#define Us SYSTEM_RESERVED_NAME
+#define V SYSTEM_RESERVED_NAME
+#define Vp SYSTEM_RESERVED_NAME
+#define Vs SYSTEM_RESERVED_NAME
+#define X SYSTEM_RESERVED_NAME
+#define Xp SYSTEM_RESERVED_NAME
+#define Xs SYSTEM_RESERVED_NAME
+
+// The classic Windows min/max macros
+#define min SYSTEM_RESERVED_NAME
+#define max SYSTEM_RESERVED_NAME
+
+// Test to make sure curses has no conflicting macros with the standard library
+#define move SYSTEM_RESERVED_NAME
+#define erase SYSTEM_RESERVED_NAME
+#define refresh SYSTEM_RESERVED_NAME
+
+// Dinkumware libc ctype.h uses these definitions
+#define _XA SYSTEM_RESERVED_NAME
+#define _XS SYSTEM_RESERVED_NAME
+#define _BB SYSTEM_RESERVED_NAME
+#define _CN SYSTEM_RESERVED_NAME
+#define _DI SYSTEM_RESERVED_NAME
+#define _LO SYSTEM_RESERVED_NAME
+#define _PU SYSTEM_RESERVED_NAME
+#define _SP SYSTEM_RESERVED_NAME
+#define _UP SYSTEM_RESERVED_NAME
+#define _XD SYSTEM_RESERVED_NAME
diff --git a/clang/test/Headers/x86-intrinsics-headers-clean.cpp b/clang/test/Headers/x86-intrinsics-headers-clean.cpp
index a19207f..0a04bce 100644
--- a/clang/test/Headers/x86-intrinsics-headers-clean.cpp
+++ b/clang/test/Headers/x86-intrinsics-headers-clean.cpp
@@ -10,4 +10,6 @@
 
 // expected-no-diagnostics
 
+#include "system_reserved_names.h"
+
 #include <x86intrin.h>
diff --git a/clang/test/Headers/x86-intrinsics-headers.c b/clang/test/Headers/x86-intrinsics-headers.c
index dc06cbd..89a7d4d 100644
--- a/clang/test/Headers/x86-intrinsics-headers.c
+++ b/clang/test/Headers/x86-intrinsics-headers.c
@@ -5,6 +5,8 @@
 // XFAIL: target=arm64ec-pc-windows-msvc
 // These intrinsics are not yet implemented for Arm64EC.
 
+#include "system_reserved_names.h"
+
 #if defined(i386) || defined(__x86_64__)
 
 #ifdef __SSE4_2__
diff --git a/clang/test/Modules/pr133057.cpp b/clang/test/Modules/pr133057.cpp
new file mode 100644
index 0000000..b273fc3
--- /dev/null
+++ b/clang/test/Modules/pr133057.cpp
@@ -0,0 +1,143 @@
+// RUN: rm -rf %t
+// RUN: mkdir -p %t
+// RUN: split-file %s %t
+//
+// RUN: %clang_cc1 -xc++ -std=c++20 -emit-module -fmodule-name=hf -fno-cxx-modules -fmodules -fno-implicit-modules %t/CMO.cppmap -o %t/WI9.pcm
+// RUN: %clang_cc1 -xc++ -std=c++20 -emit-module -fmodule-name=g -fno-cxx-modules -fmodules -fno-implicit-modules -fmodule-file=%t/WI9.pcm %t/E6H.cppmap -o %t/4BK.pcm
+// RUN: %clang_cc1 -xc++ -std=c++20 -emit-module -fmodule-name=r -fno-cxx-modules -fmodules -fno-implicit-modules -fmodule-file=%t/WI9.pcm %t/HMT.cppmap -o %t/LUM.pcm
+// RUN: %clang_cc1 -xc++ -std=c++20 -emit-module -fmodule-name=q -fno-cxx-modules -fmodules -fno-implicit-modules -fmodule-file=%t/LUM.pcm -fmodule-file=%t/4BK.pcm %t/JOV.cppmap -o %t/9VX.pcm
+// RUN: %clang_cc1 -xc++ -std=c++20 -verify -fsyntax-only -fno-cxx-modules -fmodules -fno-implicit-modules -fmodule-file=%t/9VX.pcm %t/XFD.cc
+
+//--- 2OT.h
+#include "LQ1.h"
+
+namespace ciy {
+namespace xqk {
+template <typename>
+class vum {
+ public:
+  using sc = std::C::wmd;
+  friend bool operator==(vum, vum);
+};
+template <typename>
+class me {
+ public:
+  using vbh = vum<me>;
+  using sc = std::C::vy<vbh>::sc;
+  template <typename db>
+  operator db() { return {}; }
+};
+}  // namespace xqk
+template <typename vus>
+xqk::me<vus> uvo(std::C::wmd, vus);
+}  // namespace ciy
+
+class ua {
+  std::C::wmd kij() {
+    ciy::uvo(kij(), '-');
+    return {};
+  }
+};
+
+//--- 9KF.h
+#include "LQ1.h"
+#include "2OT.h"
+namespace {
+void al(std::C::wmd lou) { std::C::jv<std::C::wmd> yt = ciy::uvo(lou, '/'); }
+}  // namespace
+
+//--- CMO.cppmap
+module "hf" {
+header "LQ1.h"
+}
+
+
+//--- E6H.cppmap
+module "g" {
+export *
+header "2OT.h"
+}
+
+
+//--- HMT.cppmap
+module "r" {
+header "2OT.h"
+}
+
+
+//--- JOV.cppmap
+module "q" {
+header "9KF.h"
+}
+
+
+//--- LQ1.h
+namespace std {
+namespace C {
+template <class zd>
+struct vy : zd {};
+template <class ub>
+struct vy<ub*> {
+  typedef ub jz;
+};
+struct wmd {};
+template <class uo, class zt>
+void sk(uo k, zt gf) {
+  (void)(k != gf);
+}
+template <class uo>
+class fm {
+ public:
+  fm(uo);
+};
+template <class kj, class kju>
+bool operator==(kj, kju);
+template <class epn>
+void afm(epn) {
+  using yp = vy<epn>;
+  if (__is_trivially_copyable(yp)) {
+    sk(fm(epn()), nullptr);
+  }
+}
+template <class ub>
+class jv {
+ public:
+  constexpr void gq();
+  ub *nef;
+};
+template <class ub>
+constexpr void jv<ub>::gq() {
+    afm(nef);
+}
+}  // namespace C
+}  // namespace std
+namespace ciy {
+}  // namespace ciy
+
+//--- XFD.cc
+// expected-no-diagnostics
+#include "LQ1.h"
+#include "2OT.h"
+class wiy {
+ public:
+  std::C::wmd eyb();
+};
+template <typename wpa>
+void i(wpa fg) {
+  std::C::jv<std::C::wmd> zs;
+  zs = ciy::uvo(fg.eyb(), '\n');
+}
+namespace ciy {
+namespace xqk {
+struct sbv;
+std::C::jv<sbv> ns() {
+  std::C::jv<sbv> ubs;
+  ubs.gq();
+  return ubs;
+}
+}  // namespace xqk
+}  // namespace ciy
+void s() {
+  wiy fg;
+  i(fg);
+}
diff --git a/clang/test/SemaCUDA/vararg.cu b/clang/test/SemaCUDA/vararg.cu
index 0238f42..62693e1 100644
--- a/clang/test/SemaCUDA/vararg.cu
+++ b/clang/test/SemaCUDA/vararg.cu
@@ -1,11 +1,9 @@
 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \
-// RUN:   -verify -DEXPECT_VA_ARG_ERR -DEXPECT_VARARG_ERR %s
+// RUN:   -verify -DEXPECT_VA_ARG_ERR %s
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \
 // RUN:   -fcuda-allow-variadic-functions -verify -DEXPECT_VA_ARG_ERR %s
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify \
-// RUN:   -DEXPECT_VARARG_ERR %s
 
 #include <stdarg.h>
 #include "Inputs/cuda.h"
@@ -30,28 +28,15 @@ __device__ void baz() {
 #endif
 }
 
-__device__ void vararg(const char* x, ...) {}
-#ifdef EXPECT_VARARG_ERR
-// expected-error@-2 {{CUDA device code does not support variadic functions}}
-#endif
+__device__ void vararg(const char* x, ...) {} // OK
 
 template <typename T>
-__device__ void vararg(T t, ...) {}
-#ifdef EXPECT_VARARG_ERR
-// expected-error@-2 {{CUDA device code does not support variadic functions}}
-#endif
+__device__ void vararg(T t, ...) {} // OK
 
 extern "C" __device__ int printf(const char* fmt, ...);  // OK, special case.
 
-// Definition of printf not allowed.
-extern "C" __device__ int printf(const char* fmt, ...) { return 0; }
-#ifdef EXPECT_VARARG_ERR
-// expected-error@-2 {{CUDA device code does not support variadic functions}}
-#endif
+extern "C" __device__ int printf(const char* fmt, ...) { return 0; } // OK
 
 namespace ns {
-__device__ int printf(const char* fmt, ...);
-#ifdef EXPECT_VARARG_ERR
-// expected-error@-2 {{CUDA device code does not support variadic functions}}
-#endif
+__device__ int printf(const char* fmt, ...); // OK
 }
diff --git a/clang/test/SemaTemplate/alias-template-deprecated.cpp b/clang/test/SemaTemplate/alias-template-deprecated.cpp
index 7418e222..6dffd37 100644
--- a/clang/test/SemaTemplate/alias-template-deprecated.cpp
+++ b/clang/test/SemaTemplate/alias-template-deprecated.cpp
@@ -46,23 +46,19 @@ using UsingInstWithCPPAttr [[deprecated("Do not use this")]] = NoAttr<int>;
 void bar() {
   NoAttr<int> obj; // Okay
 
-  // expected-warning@+2 {{'UsingWithAttr' is deprecated}}
-  // expected-note@+1 {{in instantiation of template type alias 'UsingWithAttr' requested here}}
+  // expected-warning@+1 {{'UsingWithAttr' is deprecated}}
   UsingWithAttr<int> objUsingWA;
 
-  // expected-warning@+2 {{'UsingWithAttr' is deprecated}}
-  // expected-note@+1 {{in instantiation of template type alias 'UsingWithAttr' requested here}}
+  // expected-warning@+1 {{'UsingWithAttr' is deprecated}}
   NoAttr<UsingWithAttr<int>> s;
 
   // expected-note@+1 {{'DepInt' has been explicitly marked deprecated here}}
   using DepInt [[deprecated]] = int;
-  // expected-warning@+3 {{'UsingWithAttr' is deprecated}}
-  // expected-warning@+2 {{'DepInt' is deprecated}}
-  // expected-note@+1 {{in instantiation of template type alias 'UsingWithAttr' requested here}}
+  // expected-warning@+2 {{'UsingWithAttr' is deprecated}}
+  // expected-warning@+1 {{'DepInt' is deprecated}}
   using X = UsingWithAttr<DepInt>;
 
-  // expected-warning@+2 {{'UsingWithAttr' is deprecated}}
-  // expected-note@+1 {{in instantiation of template type alias 'UsingWithAttr' requested here}}
+  // expected-warning@+1 {{'UsingWithAttr' is deprecated}}
   UsingWithAttr<int>().foo();
 
   // expected-warning@+1 {{'UsingInstWithAttr' is deprecated}}
@@ -74,8 +70,7 @@ void bar() {
   // expected-warning@+1 {{'UsingTDWithAttr' is deprecated}}
   UsingTDWithAttr objUTDWA;
 
-  // expected-warning@+2 {{'UsingWithCPPAttr' is deprecated}}
-  // expected-note@+1 {{in instantiation of template type alias 'UsingWithCPPAttr' requested here}}
+  // expected-warning@+1 {{'UsingWithCPPAttr' is deprecated}}
   UsingWithCPPAttr<int> objUsingWCPPA;
 
   // expected-warning@+1 {{'UsingInstWithCPPAttr' is deprecated: Do not use this}}
diff --git a/clang/test/SemaTemplate/alias-templates.cpp b/clang/test/SemaTemplate/alias-templates.cpp
index ab5cad7..09fe72e 100644
--- a/clang/test/SemaTemplate/alias-templates.cpp
+++ b/clang/test/SemaTemplate/alias-templates.cpp
@@ -312,3 +312,11 @@ namespace resolved_nttp {
 
   using TC2 = decltype(C<bool, 2, 3>::p); // expected-note {{instantiation of}}
 }
+
+namespace OuterSubstFailure {
+  template <class T> struct A {
+      template <class> using B = T&;
+      // expected-error@-1 {{cannot form a reference to 'void'}}
+  };
+  template struct A<void>; // expected-note {{requested here}}
+} // namespace OuterSubstFailure