[clang] [llvm] AMDGPU: Do not give gfx12.5 bvh-ray-tracing-insts (PR #213246)

[Matt Arsenault via llvm-commits]

https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/213246

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)

>From 69fb339399ba292f23dfc17d4ebff943f937a242 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Fri, 31 Jul 2026 07:12:57 +0200
Subject: [PATCH] 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)