llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-llvm-ir Author: Yaxun (Sam) Liu (yxsamliu) <details> <summary>Changes</summary> AMDGPU has FMA mix instructions that let callers select the F32, F16, or BF16 interpretation of each source. Clang does not provide direct access to these forms. Add builtins for the F32 and packed F16 and BF16 forms. Require constant source selectors and diagnose use on targets that do not support the corresponding instructions. --- Patch is 45.32 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/214814.diff 16 Files Affected: - (modified) clang/docs/ReleaseNotes.md (+3) - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+29) - (modified) clang/include/clang/Basic/BuiltinsAMDGPUDocs.td (+50) - (modified) clang/lib/Sema/SemaAMDGPU.cpp (+12) - (modified) clang/test/CodeGen/amdgpu-builtin-is-invocable.c (+1-1) - (modified) clang/test/CodeGen/amdgpu-builtin-processor-is.c (+1-1) - (modified) clang/test/CodeGen/link-builtin-bitcode.c (+3-3) - (modified) clang/test/CodeGenCXX/dynamic-cast-address-space.cpp (+2-2) - (added) clang/test/CodeGenOpenCL/builtins-amdgcn-fma-mix-target-features.cl (+23) - (added) clang/test/CodeGenOpenCL/builtins-amdgcn-fma-mix.cl (+41) - (added) clang/test/SemaOpenCL/builtins-amdgcn-fma-mix-err.cl (+50) - (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+38) - (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+1) - (modified) llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp (+13) - (modified) llvm/lib/Target/AMDGPU/VOP3PInstructions.td (+38) - (added) llvm/test/CodeGen/AMDGPU/amdgcn-fma-mix-intrinsics.ll (+57) ``````````diff diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index 7976b82b63f6e..0636808e50f01 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -467,6 +467,9 @@ features cannot lower the translation-unit ABI level; #### AMDGPU Support +- Added builtins for the `v_fma_mix_f32`, `v_fma_mixlo`, and `v_fma_mixhi` + instruction families. + - Deprecated the following builtins in favor of `__builtin_amdgcn_ballot_w32` or `__builtin_amdgcn_ballot_w64`: - `__builtin_amdgcn_uicmp` diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 77b402a8a63b2..8aa340c67ff26 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -1051,6 +1051,35 @@ def __builtin_amdgcn_exp2_bf16 : AMDGPUBuiltin<"__bf16(__bf16)", [Const], "bf16- def __builtin_amdgcn_sin_bf16 : AMDGPUBuiltin<"__bf16(__bf16)", [Const], "bf16-trans-insts">; def __builtin_amdgcn_cos_bf16 : AMDGPUBuiltin<"__bf16(__bf16)", [Const], "bf16-trans-insts">; +class AMDGPUFmaMixBuiltin<string feature, Documentation doc> + : AMDGPUBuiltin<"float(unsigned int, unsigned int, unsigned int, _Constant unsigned int, _Constant unsigned int, _Constant unsigned int)", + [Const], feature> { + let Documentation = [doc]; + let ArgNames = ["src0", "src1", "src2", "src0_sel", "src1_sel", + "src2_sel"]; +} + +class AMDGPUFmaMixPackedBuiltin<string feature, Documentation doc> + : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int, unsigned int, unsigned int, _Constant unsigned int, _Constant unsigned int, _Constant unsigned int)", + [Const], feature> { + let Documentation = [doc]; + let ArgNames = ["src0", "src1", "src2", "dst", "src0_sel", "src1_sel", + "src2_sel"]; +} + +def __builtin_amdgcn_fma_mix_f32 + : AMDGPUFmaMixBuiltin<"fma-mix-insts", DocFmaMixF32>; +def __builtin_amdgcn_fma_mix_f32_bf16 + : AMDGPUFmaMixBuiltin<"fma-mix-bf16-insts", DocFmaMixF32BF16>; +def __builtin_amdgcn_fma_mixlo_f16 + : AMDGPUFmaMixPackedBuiltin<"fma-mix-insts", DocFmaMixLoF16>; +def __builtin_amdgcn_fma_mixhi_f16 + : AMDGPUFmaMixPackedBuiltin<"fma-mix-insts", DocFmaMixHiF16>; +def __builtin_amdgcn_fma_mixlo_bf16 + : AMDGPUFmaMixPackedBuiltin<"fma-mix-bf16-insts", DocFmaMixLoBF16>; +def __builtin_amdgcn_fma_mixhi_bf16 + : AMDGPUFmaMixPackedBuiltin<"fma-mix-bf16-insts", DocFmaMixHiBF16>; + def __builtin_amdgcn_cvt_sr_pk_bf16_f32 : AMDGPUBuiltin<"_ExtVector<2, __bf16>(float, float, int)", [Const], "cvt-sr-pk-bf16-f32-inst">; def __builtin_amdgcn_cvt_sr_pk_f16_f32 : AMDGPUBuiltin<"_ExtVector<2, _Float16>(float, float, int)", [Const], "gfx1250-insts">; def __builtin_amdgcn_cvt_f16_fp8 : AMDGPUBuiltin<"_Float16(int, _Constant int)", [Const], "gfx1250-insts">; diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td index 8dafa45cbbd99..3be4e15ea05b3 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td @@ -78,6 +78,23 @@ between standard floating-point types, integer types and packed low-precision fo }]; } +def DocCatAMDGPUFmaMix : DocumentationCategory<"FMA Mix Builtins"> { + let Content = [{ +These builtins expose the AMDGPU FMA mix instructions. Each 32-bit source is +paired with a compile-time source selector in the range [0, 3]. Selector bit 0 +chooses the low or high half when bit 1 selects 16-bit interpretation. Bit 0 +is ignored for 32-bit interpretation. Selectors 0 and 1 select F32, selector 2 +selects the low F16 or BF16 value, and selector 3 selects the high F16 or BF16 +value. + +The unsigned source arguments carry raw F32, F16, or BF16 bit patterns. They +are not numerically converted from unsigned integers. + +The ``mixlo`` and ``mixhi`` builtins replace one 16-bit half of ``dst`` and +preserve the other half. +}]; +} + def DocCatAMDGPUPrefetch : DocumentationCategory<"Prefetch Builtins"> { let Content = [{ These builtins provide access to AMDGPU prefetch instructions, including instructions @@ -85,6 +102,39 @@ for prefetching data and instructions. }]; } +//===----------------------------------------------------------------------===// +// FMA mix builtins — Documentation records +//===----------------------------------------------------------------------===// + +class DocFmaMix<string instruction, string result, string feature> + : Documentation { + let Category = DocCatAMDGPUFmaMix; + let Content = !strconcat("Computes ``src0 * src1 + src2`` and returns ", + result, ". Requires target feature ``", feature, + "`` and maps to the ``", instruction, + "`` instruction."); +} + +def DocFmaMixF32 : + DocFmaMix<"v_fma_mix_f32", "an F32 result", "fma-mix-insts">; +def DocFmaMixF32BF16 : + DocFmaMix<"v_fma_mix_f32_bf16", "an F32 result", + "fma-mix-bf16-insts">; +def DocFmaMixLoF16 : + DocFmaMix<"v_fma_mixlo_f16", "the packed result with its low half replaced", + "fma-mix-insts">; +def DocFmaMixHiF16 : + DocFmaMix<"v_fma_mixhi_f16", "the packed result with its high half replaced", + "fma-mix-insts">; +def DocFmaMixLoBF16 : + DocFmaMix<"v_fma_mixlo_bf16", + "the packed result with its low half replaced", + "fma-mix-bf16-insts">; +def DocFmaMixHiBF16 : + DocFmaMix<"v_fma_mixhi_bf16", + "the packed result with its high half replaced", + "fma-mix-bf16-insts">; + //===----------------------------------------------------------------------===// // ABI / Special Register Builtins — Documentation records //===----------------------------------------------------------------------===// diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 2208865e3237b..2913e63e4b2ce 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -444,6 +444,18 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, /*High=*/0) || SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/2, /*Low=*/0, /*High=*/0); + case AMDGPU::BI__builtin_amdgcn_fma_mix_f32: + case AMDGPU::BI__builtin_amdgcn_fma_mix_f32_bf16: + return SemaRef.BuiltinConstantArgRange(TheCall, 3, 0, 3) || + SemaRef.BuiltinConstantArgRange(TheCall, 4, 0, 3) || + SemaRef.BuiltinConstantArgRange(TheCall, 5, 0, 3); + case AMDGPU::BI__builtin_amdgcn_fma_mixlo_f16: + case AMDGPU::BI__builtin_amdgcn_fma_mixhi_f16: + case AMDGPU::BI__builtin_amdgcn_fma_mixlo_bf16: + case AMDGPU::BI__builtin_amdgcn_fma_mixhi_bf16: + return SemaRef.BuiltinConstantArgRange(TheCall, 4, 0, 3) || + SemaRef.BuiltinConstantArgRange(TheCall, 5, 0, 3) || + SemaRef.BuiltinConstantArgRange(TheCall, 6, 0, 3); default: return false; } diff --git a/clang/test/CodeGen/amdgpu-builtin-is-invocable.c b/clang/test/CodeGen/amdgpu-builtin-is-invocable.c index 06c8bea3aadfc..ff7eec3492850 100644 --- a/clang/test/CodeGen/amdgpu-builtin-is-invocable.c +++ b/clang/test/CodeGen/amdgpu-builtin-is-invocable.c @@ -59,7 +59,7 @@ void foo() { // AMDGCN-GFX1010: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } // AMDGCN-GFX1010: attributes #[[ATTR2]] = { noreturn nounwind } //. -// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+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,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-sr-pk-bf16-f32-inst,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+extended-image-insts,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+smem-prefetch-insts,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } +// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+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,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-sr-pk-bf16-f32-inst,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+extended-image-insts,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fma-mix-bf16-insts,+fma-mix-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+smem-prefetch-insts,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } // AMDGCNSPIRV: attributes #[[ATTR1:[0-9]+]] = { nounwind } // AMDGCNSPIRV: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } // AMDGCNSPIRV: attributes #[[ATTR3]] = { noreturn nounwind } diff --git a/clang/test/CodeGen/amdgpu-builtin-processor-is.c b/clang/test/CodeGen/amdgpu-builtin-processor-is.c index 0056bf81be202..04f05989b95db 100644 --- a/clang/test/CodeGen/amdgpu-builtin-processor-is.c +++ b/clang/test/CodeGen/amdgpu-builtin-processor-is.c @@ -68,7 +68,7 @@ void foo() { //. // AMDGCN-GFX1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } //. -// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+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,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-sr-pk-bf16-f32-inst,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+extended-image-insts,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+smem-prefetch-insts,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } +// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+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,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-sr-pk-bf16-f32-inst,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+extended-image-insts,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fma-mix-bf16-insts,+fma-mix-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+smem-prefetch-insts,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } // AMDGCNSPIRV: attributes #[[ATTR1:[0-9]+]] = { nounwind } // AMDGCNSPIRV: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } // AMDGCNSPIRV: attributes #[[ATTR3]] = { noreturn nounwind } diff --git a/clang/test/CodeGen/link-builtin-bitcode.c b/clang/test/CodeGen/link-builtin-bitcode.c index 8a6aebdbd66dd..cf3cb98c289fe 100644 --- a/clang/test/CodeGen/link-builtin-bitcode.c +++ b/clang/test/CodeGen/link-builtin-bitcode.c @@ -44,6 +44,6 @@ int bar() { return no_attr() + attr_in_target() + attr_not_in_target() + attr_in // CHECK-SAME: () #[[ATTR_INCOMPATIBLE:[0-9]+]] { // CHECK: attributes #[[ATTR_BAR]] = { {{.*}} "no-trapping-math"="true" {{.*}} } -// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } -// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } -// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" } +// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+fma-mix-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } +// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-global-insts,+fma-mix-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } +// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+fma-mix-insts,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" } diff --git a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp index d0e069804e5c2..b9673d397ec82 100644 --- a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp +++ b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp @@ -107,9 +107,9 @@ const B& f(A *a) { // CHECK: attributes #[[ATTR3]] = { nounwind } // CHECK: attributes #[[ATTR4]] = { noreturn } //. -// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-fl... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/214814 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
