https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/213270
>From 69fb339399ba292f23dfc17d4ebff943f937a242 Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Fri, 31 Jul 2026 07:12:57 +0200 Subject: [PATCH 1/2] AMDGPU: Do not give gfx12.5 bvh-ray-tracing-insts bvh-ray-tracing-insts was listed in the FeatureGFX12 generation, so gfx1250/gfx1251/gfx12-5-generic inherited it even though they have no BVH. Move the feature out of the common base and into FeatureISAVersion12, which This stops clang from wrongly accepting __builtin_amdgcn_image_bvh_intersect_ray* on gfx1250. Co-authored-by: Claude (Claude-Opus-4.8) --- ...builtins-amdgcn-image-bvh-intersect-ray-err.cl | 15 +++++++++++++++ llvm/lib/Target/AMDGPU/AMDGPU.td | 3 ++- .../GlobalISel/llvm.amdgcn.intersect_ray.ll | 1 + .../CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll | 2 ++ 4 files changed, 20 insertions(+), 1 deletion(-) create mode 100644 clang/test/SemaOpenCL/builtins-amdgcn-image-bvh-intersect-ray-err.cl diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-image-bvh-intersect-ray-err.cl b/clang/test/SemaOpenCL/builtins-amdgcn-image-bvh-intersect-ray-err.cl new file mode 100644 index 0000000000000..d81f223da0f13 --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-amdgcn-image-bvh-intersect-ray-err.cl @@ -0,0 +1,15 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu12.5 -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgpu12.50 -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgpu12.51 -verify -S -o - %s + +typedef unsigned int uint4 __attribute__((ext_vector_type(4))); +typedef float float4 __attribute__((ext_vector_type(4))); + +// gfx12.5 does not have the BVH ray tracing instructions; unlike gfx1200 it does +// not inherit bvh-ray-tracing-insts from the GFX12 generation. +void test_image_bvh_intersect_ray(global uint4 *out, unsigned node, float ext, + float4 origin, float4 dir, float4 invdir, + uint4 desc) { + *out = __builtin_amdgcn_image_bvh_intersect_ray(node, ext, origin, dir, invdir, desc); // expected-error{{'__builtin_amdgcn_image_bvh_intersect_ray' needs target feature bvh-ray-tracing-insts}} +} diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index 01c39232ff479..45224250d8e5a 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -1679,7 +1679,7 @@ def FeatureGFX12 : GCNSubtargetFeatureGeneration<"GFX12", FeatureFlatAddressSpace, Feature16BitInsts, FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts, - FeatureBVHRayTracingInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, + FeatureGFX10_BEncoding, FeatureGFX10_3Insts, FeatureGFX11Insts, FeatureGFX12Insts, FeatureVOP3PInsts, FeatureVOPDInsts, FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp, FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts, @@ -2205,6 +2205,7 @@ def FeatureISAVersion12 : FeatureSet< FeatureFlatAtomicFaddF32Inst, FeatureImageInsts, FeatureExtendedImageInsts, + FeatureBVHRayTracingInsts, FeatureFP8ConversionInsts, FeatureCvtFP8ByteSel, FeatureWMMA128bInsts, diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll index 1d0ae809d3707..a2518b90bb0f2 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll @@ -4,6 +4,7 @@ ; RUN: llc -global-isel -mtriple=amdgpu11.00 -mattr=+real-true16 < %s | FileCheck -check-prefixes=GCN,GFX11,GFX11-TRUE16 %s ; RUN: llc -global-isel -mtriple=amdgpu11.00 -mattr=-real-true16 < %s | FileCheck -check-prefixes=GCN,GFX11,GFX11-FAKE16 %s ; RUN: not llc -global-isel -mtriple=amdgpu10.12 < %s -filetype=null 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel -mtriple=amdgpu12.50 < %s -filetype=null 2>&1 | FileCheck -check-prefix=ERR %s ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr) ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr) diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll index d8117b76fb576..e8ac7c139edc9 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll @@ -15,6 +15,8 @@ ; RUN: not llc -global-isel=0 -mtriple=amdgpu10.12 < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=1 -mtriple=amdgpu10.12 < %s 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel=0 -mtriple=amdgpu12.50 < %s 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel=1 -mtriple=amdgpu12.50 < %s 2>&1 | FileCheck -check-prefix=ERR %s ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr) ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr) >From becbf176825b292175338a4d60acc5c7e13bbbf6 Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Fri, 31 Jul 2026 08:23:08 +0200 Subject: [PATCH 2/2] AMDGPU: Separate image_msaa_load from bvh-ray-tracing-insts Add a separate msaa-load-insts feature so image_msaa_load is available on gfx13, which has it but not the BVH intersect-ray instructions. These were assumed to be the same feature previously, but gfx13 does not have the bvh instructions, but does have image_msaa_load. Co-authored-by: Claude (Claude-Opus-4.8) --- llvm/lib/Target/AMDGPU/AMDGPU.td | 17 ++++++++++++++--- llvm/lib/Target/AMDGPU/MIMGInstructions.td | 8 +++++--- .../GlobalISel/llvm.amdgcn.intersect_ray.ll | 1 + .../CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll | 2 ++ llvm/test/MC/AMDGPU/gfx13_asm_vimage_err.s | 13 +++++++++++++ 5 files changed, 35 insertions(+), 6 deletions(-) create mode 100644 llvm/test/MC/AMDGPU/gfx13_asm_vimage_err.s diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index 45224250d8e5a..6a7299482bdc7 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -796,6 +796,13 @@ defm BVHRayTracingInsts : AMDGPUSubtargetFeature<"bvh-ray-tracing-insts", "Has BVH ray tracing instructions" >; +defm MSAALoadInsts : AMDGPUSubtargetFeature<"msaa-load-insts", + "Has image_msaa_load instructions", + /*GenPredicate=*/1, + /*GenAssemblerPredicate=*/1, + /*Deps=*/[FeatureImageInsts] +>; + defm GFX10_BEncoding : AMDGPUSubtargetFeature<"gfx10_b-encoding", "Encoding format GFX10_B", /*GenPredicate=*/0 @@ -1650,7 +1657,8 @@ def FeatureGFX11 : GCNSubtargetFeatureGeneration<"GFX11", FeatureFlatAddressSpace, Feature16BitInsts, FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts, - FeatureBVHRayTracingInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, + FeatureBVHRayTracingInsts, FeatureMSAALoadInsts, + FeatureGFX10_BEncoding, FeatureGFX10_3Insts, FeatureGFX11Insts, FeatureVOP3PInsts, FeatureVOPDInsts, FeatureTrue16BitInsts, FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp, FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts, @@ -1707,7 +1715,7 @@ def FeatureGFX13 : GCNSubtargetFeatureGeneration<"GFX13", FeatureFlatAddressSpace, Feature16BitInsts, FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts, - FeatureBVHRayTracingInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, + FeatureMSAALoadInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, FeatureGFX11Insts, FeatureGFX12Insts, FeatureGFX13Insts, FeatureVOP3PInsts, FeatureVOPDInsts, FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp, FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts, @@ -2046,11 +2054,13 @@ def FeatureISAVersion10_1_2 : FeatureSet< def FeatureISAVersion10_1_3 : FeatureSet< !listconcat(FeatureISAVersion10_1_Common.Features, - [FeatureBVHRayTracingInsts])>; + [FeatureBVHRayTracingInsts, + FeatureMSAALoadInsts])>; def FeatureISAVersion10_3_0 : FeatureSet< !listconcat(FeatureISAVersion10_Common.Features, [FeatureBVHRayTracingInsts, + FeatureMSAALoadInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, FeatureDot1Insts, @@ -2206,6 +2216,7 @@ def FeatureISAVersion12 : FeatureSet< FeatureImageInsts, FeatureExtendedImageInsts, FeatureBVHRayTracingInsts, + FeatureMSAALoadInsts, FeatureFP8ConversionInsts, FeatureCvtFP8ByteSel, FeatureWMMA128bInsts, diff --git a/llvm/lib/Target/AMDGPU/MIMGInstructions.td b/llvm/lib/Target/AMDGPU/MIMGInstructions.td index 004d4647d0797..cb554073f6398 100644 --- a/llvm/lib/Target/AMDGPU/MIMGInstructions.td +++ b/llvm/lib/Target/AMDGPU/MIMGInstructions.td @@ -1926,17 +1926,19 @@ defm IMAGE_SAMPLE_C_CD_CL_O_G16 : MIMG_Sampler <mimgopc<MIMG.NOP, MIMG.NOP, MIMG //def IMAGE_RSRC256 : MIMG_NoPattern_RSRC256 <"image_rsrc256", mimgopc<0x7e>>; //def IMAGE_SAMPLER : MIMG_NoPattern_ <"image_sampler", mimgopc<0x7f>>; -let OtherPredicates = [HasImageInsts, HasBVHRayTracingInsts, isGFX10Only] in +let SubtargetPredicate = HasMSAALoadInsts, OtherPredicates = [isGFX10Only] in defm IMAGE_MSAA_LOAD_X : MIMG_NoSampler <mimgopc<MIMG.NOP, MIMG.NOP, MIMG.NOP, 0x80>, "image_msaa_load", 1, 0, 0, 1>; -let OtherPredicates = [HasImageInsts, HasBVHRayTracingInsts] in { +let SubtargetPredicate = HasMSAALoadInsts in { defm IMAGE_MSAA_LOAD : MIMG_MSAA_Load <mimgopc<0x80, 0x18, 0x18, MIMG.NOP>, "image_msaa_load">; +} // End SubtargetPredicate = HasMSAALoadInsts +let SubtargetPredicate = HasBVHRayTracingInsts in { defm IMAGE_BVH_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x19, 0x19, 0xe6>, "image_bvh_intersect_ray", 0, 0, 0>; defm IMAGE_BVH_INTERSECT_RAY_a16 : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x19, 0x19, 0xe6>, "image_bvh_intersect_ray", 0, 1, 0>; defm IMAGE_BVH64_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x1a, 0x1a, 0xe7>, "image_bvh64_intersect_ray", 1, 0, 0>; defm IMAGE_BVH64_INTERSECT_RAY_a16 : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x1a, 0x1a, 0xe7>, "image_bvh64_intersect_ray", 1, 1, 0>; -} // End OtherPredicates = [HasImageInsts, HasBVHRayTracingInsts] +} // End SubtargetPredicate = HasBVHRayTracingInsts defm IMAGE_BVH_DUAL_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x80, MIMG.NOP, MIMG.NOP>, "image_bvh_dual_intersect_ray", 1, 0, 1>; defm IMAGE_BVH8_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x81, MIMG.NOP, MIMG.NOP>, "image_bvh8_intersect_ray", 1, 0, 0, 1>; diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll index a2518b90bb0f2..ee93f413566e5 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll @@ -5,6 +5,7 @@ ; RUN: llc -global-isel -mtriple=amdgpu11.00 -mattr=-real-true16 < %s | FileCheck -check-prefixes=GCN,GFX11,GFX11-FAKE16 %s ; RUN: not llc -global-isel -mtriple=amdgpu10.12 < %s -filetype=null 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel -mtriple=amdgpu12.50 < %s -filetype=null 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel -mtriple=amdgpu13.10 < %s -filetype=null 2>&1 | FileCheck -check-prefix=ERR %s ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr) ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr) diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll index e8ac7c139edc9..173739fa8f35f 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll @@ -17,6 +17,8 @@ ; RUN: not llc -global-isel=1 -mtriple=amdgpu10.12 < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=0 -mtriple=amdgpu12.50 < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=1 -mtriple=amdgpu12.50 < %s 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel=0 -mtriple=amdgpu13.10 < %s 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel=1 -mtriple=amdgpu13.10 < %s 2>&1 | FileCheck -check-prefix=ERR %s ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr) ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr) diff --git a/llvm/test/MC/AMDGPU/gfx13_asm_vimage_err.s b/llvm/test/MC/AMDGPU/gfx13_asm_vimage_err.s new file mode 100644 index 0000000000000..2107956f8bec8 --- /dev/null +++ b/llvm/test/MC/AMDGPU/gfx13_asm_vimage_err.s @@ -0,0 +1,13 @@ +// RUN: not llvm-mc -triple=amdgpu13.10 -filetype=null %s 2>&1 | FileCheck --check-prefix=GFX13-ERR --implicit-check-not=error: --strict-whitespace %s + +image_bvh_intersect_ray v[4:7], [v9, v10, v[11:13], v[14:16], v[17:19]], s[4:7] +// GFX13-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU + +image_bvh_intersect_ray v[4:7], [v9, v10, v[11:13], v[14:16]], s[4:7] a16 +// GFX13-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU + +image_bvh64_intersect_ray v[4:7], [v[9:10], v11, v[12:14], v[15:17], v[18:20]], s[4:7] +// GFX13-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU + +image_bvh64_intersect_ray v[4:7], [v[9:10], v11, v[12:14], v[15:17]], s[4:7] a16 +// GFX13-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
