[clang] [Clang][AMDGPU] Name the builtin in required target feature diagnostics (PR #218358)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Aug 24 02:43:09 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Rana Pratap Reddy (ranapratap55)
<details>
<summary>Changes</summary>
The Sema target feature checks for the builtins passed `getCurFunctionDecl()` to `err_builtin_needs_feature`, so the diagnostic named the function containing the call rather than the builtin that actually requires the feature. Now, we are taking the name from `TheCall->getDirectCallee()`.
Before:
`error: 'test_builtin_image_load_2d' needs target feature image-insts`
After:
`error: '__builtin_amdgcn_image_load_2d_f32_i32' needs target feature image-insts`
---
Patch is 80.87 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/218358.diff
6 Files Affected:
- (modified) clang/include/clang/Sema/SemaAMDGPU.h (+5)
- (modified) clang/lib/Sema/SemaAMDGPU.cpp (+23-16)
- (modified) clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c (+48-48)
- (modified) clang/test/SemaOpenCL/builtins-extended-image-err.cl (+45-45)
- (modified) clang/test/SemaOpenCL/builtins-image-load-image-feature-err.cl (+48-48)
- (modified) clang/test/SemaOpenCL/builtins-image-store-image-feature-err.cl (+28-28)
``````````diff
diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h
index 0cf4f8adcf108..bd18245102894 100644
--- a/clang/include/clang/Sema/SemaAMDGPU.h
+++ b/clang/include/clang/Sema/SemaAMDGPU.h
@@ -16,6 +16,7 @@
#include "clang/AST/ASTFwd.h"
#include "clang/Sema/SemaBase.h"
#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/ADT/StringMap.h"
namespace clang {
class AttributeCommonInfo;
@@ -31,6 +32,10 @@ class SemaAMDGPU : public SemaBase {
bool CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall);
+ bool checkBuiltinRequiredTargetFeatures(
+ unsigned BuiltinID, CallExpr *TheCall,
+ const llvm::StringMap<bool> &CallerFeatureMap);
+
/// Emits a diagnostic if the \p E is not an atomic ordering encoded in the C
/// ABI format, or if the atomic ordering is not valid for the operation type
/// as defined by \p MayLoad and \p MayStore. \returns true if a diagnostic
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index 0bda851b0c1cf..399e606753596 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -34,6 +34,23 @@ namespace clang {
SemaAMDGPU::SemaAMDGPU(Sema &S) : SemaBase(S) {}
+bool SemaAMDGPU::checkBuiltinRequiredTargetFeatures(
+ unsigned BuiltinID, CallExpr *TheCall,
+ const llvm::StringMap<bool> &CallerFeatureMap) {
+ ASTContext &Ctx = getASTContext();
+ StringRef FeatureList(Ctx.BuiltinInfo.getRequiredFeatures(BuiltinID));
+ if (Builtin::evaluateRequiredTargetFeatures(FeatureList, CallerFeatureMap))
+ return false;
+
+ const FunctionDecl *BuiltinDecl = TheCall->getDirectCallee();
+ Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
+ << (BuiltinDecl ? BuiltinDecl->getDeclName()
+ : DeclarationName(&Ctx.Idents.get(
+ Ctx.BuiltinInfo.getName(BuiltinID))))
+ << FeatureList;
+ return true;
+}
+
bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
CallExpr *TheCall) {
const auto *FD = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
@@ -306,14 +323,9 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32: {
- StringRef FeatureList(
- getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
- if (!Builtin::evaluateRequiredTargetFeatures(FeatureList,
- CallerFeatureMap)) {
- Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
- << FD->getDeclName() << FeatureList;
- return false;
- }
+ if (checkBuiltinRequiredTargetFeatures(BuiltinID, TheCall,
+ CallerFeatureMap))
+ return true;
unsigned ArgCount = TheCall->getNumArgs() - 1;
llvm::APSInt Result;
@@ -383,14 +395,9 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
- StringRef FeatureList(
- getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
- if (!Builtin::evaluateRequiredTargetFeatures(FeatureList,
- CallerFeatureMap)) {
- Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
- << FD->getDeclName() << FeatureList;
- return false;
- }
+ if (checkBuiltinRequiredTargetFeatures(BuiltinID, TheCall,
+ CallerFeatureMap))
+ return true;
unsigned ArgCount = TheCall->getNumArgs() - 1;
llvm::APSInt Result;
diff --git a/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c b/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c
index c7680b907f37e..abfddf2a413da 100644
--- a/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c
+++ b/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c
@@ -9,106 +9,106 @@ typedef half half4 __attribute__((ext_vector_type(4)));
typedef int int4 __attribute__((ext_vector_type(4)));
void test(half4 v, __amdgpu_texture_t tex) {
- v = __builtin_amdgcn_image_load_1d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_1d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_1d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_1darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_1darray_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_1darray_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_2d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_2d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_2d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_2darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_2darray_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_2darray_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_3d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_3d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_3d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_cube_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_cube_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_cube_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_mip_1d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_mip_1d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_mip_1d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_mip_1darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_mip_1darray_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_mip_1darray_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_mip_2d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_mip_2d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_mip_2d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_mip_2darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_mip_2darray_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_mip_2darray_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_mip_3d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_mip_3d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_mip_3d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, 0, 0, tex, 0, 0);
- v = __builtin_amdgcn_image_load_mip_cube_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_load_mip_cube_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_load_mip_cube_v4f16_i32' needs target feature image-insts,16-bit-insts}}
0, 0, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_1d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_1d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_1d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_1darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_1darray_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_1darray_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_2d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_2d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_2d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_2darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_2darray_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_2darray_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_3d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_3d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_3d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_cube_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_cube_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_cube_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_mip_1d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_mip_1d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_mip_1d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_mip_1darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_mip_1darray_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_mip_1darray_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_mip_2d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_mip_2d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_mip_2d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_mip_2darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_mip_2darray_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_mip_2darray_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_mip_3d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_mip_3d_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_mip_3d_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, 0, 0, tex, 0, 0);
- __builtin_amdgcn_image_store_mip_cube_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ __builtin_amdgcn_image_store_mip_cube_v4f16_i32( // expected-error {{'__builtin_amdgcn_image_store_mip_cube_v4f16_i32' needs target feature image-insts,16-bit-insts}}
v, 0, 0, 0, 0, 0, tex, 0, 0);
}
void test_sample(half4 v, float f32, __amdgpu_texture_t tex, int4 vec4i32) {
- v = __builtin_amdgcn_image_sample_1d_v4f16_f32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_1d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_1d_v4f16_f32' needs target feature image-insts,16-bit-insts}}
15, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_1darray_v4f16_f32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_1darray_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_1darray_v4f16_f32' needs target feature image-insts,16-bit-insts}}
15, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_2d_v4f16_f32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_2d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_2d_v4f16_f32' needs target feature image-insts,16-bit-insts}}
15, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_2darray_v4f16_f32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_2darray_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_2darray_v4f16_f32' needs target feature image-insts,16-bit-insts}}
15, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_3d_v4f16_f32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_3d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_3d_v4f16_f32' needs target feature image-insts,16-bit-insts}}
15, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_cube_v4f16_f32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_cube_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_cube_v4f16_f32' needs target feature image-insts,16-bit-insts}}
15, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_lz_1d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_lz_1d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_lz_1d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_lz_1darray_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_lz_1darray_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_lz_2d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_lz_2d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_lz_2d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_lz_2darray_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_lz_2darray_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_lz_3d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_lz_3d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_lz_3d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_lz_cube_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_lz_cube_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_lz_cube_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_l_1d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_l_1d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_l_1d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_l_1darray_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_l_1darray_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_l_1darray_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_l_2d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_l_2d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_l_2d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_l_2darray_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_l_2darray_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_l_2darray_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_l_3d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_l_3d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_l_3d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_l_cube_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_l_cube_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_l_cube_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_d_1d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_d_1d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_d_1d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_d_1darray_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_d_1darray_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_d_1darray_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_d_2d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_d_2d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_d_2d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_d_2darray_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_d_2darray_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_d_2darray_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 0, 0);
- v = __builtin_amdgcn_image_sample_d_3d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ v = __builtin_amdgcn_image_sample_d_3d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_sample_d_3d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 0, 0);
}
void test_gather4(half4 v, float f32, __amdgpu_texture_t tex, int4 vec4i32)...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/218358
More information about the cfe-commits
mailing list