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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
