Author: Matt Arsenault Date: 2026-07-31T17:23:42+02:00 New Revision: 61b16cb275a4823059b34f67c41c265098b6d35f
URL: https://github.com/llvm/llvm-project/commit/61b16cb275a4823059b34f67c41c265098b6d35f DIFF: https://github.com/llvm/llvm-project/commit/61b16cb275a4823059b34f67c41c265098b6d35f.diff LOG: AMDGPU: Do not give gfx12.5 bvh-ray-tracing-insts (#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. This stops clang from wrongly accepting __builtin_amdgcn_image_bvh_intersect_ray* on gfx1250. Co-authored-by: Claude (Claude-Opus-4.8) Added: clang/test/SemaOpenCL/builtins-amdgcn-image-bvh-intersect-ray-err.cl Modified: llvm/lib/Target/AMDGPU/AMDGPU.td llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll Removed: ################################################################################ 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 fb2221e671758..994c465aa0bed 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -1689,7 +1689,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, @@ -2217,6 +2217,7 @@ def FeatureISAVersion12 : FeatureSet< FeatureFlatAtomicFaddF32Inst, FeatureImageInsts, FeatureExtendedImageInsts, + FeatureBVHRayTracingInsts, FeatureFP8ConversionInsts, FeatureCvtFP8ByteSel, FeatureOCPFP8ConversionInsts, 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) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
