https://github.com/ranapratap55 created https://github.com/llvm/llvm-project/pull/218358
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` >From 698f043089603628a24f62e7232189c2609e6dbd Mon Sep 17 00:00:00 2001 From: ranapratap55 <[email protected]> Date: Mon, 24 Aug 2026 15:02:48 +0530 Subject: [PATCH] [Clang][AMDGPU] Name the builtin in required target feature diagnostics --- clang/include/clang/Sema/SemaAMDGPU.h | 5 + clang/lib/Sema/SemaAMDGPU.cpp | 39 ++++---- .../builtins-amdgcn-d16-image-16bit-error.c | 96 +++++++++---------- .../SemaOpenCL/builtins-extended-image-err.cl | 90 ++++++++--------- .../builtins-image-load-image-feature-err.cl | 96 +++++++++---------- .../builtins-image-store-image-feature-err.cl | 56 +++++------ 6 files changed, 197 insertions(+), 185 deletions(-) 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) { - v = __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}} + v = __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32( // expected-error {{'__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32' needs target feature extended-image-insts,16-bit-insts}} 1, f32, f32, tex, vec4i32, 0, 0, 0); } diff --git a/clang/test/SemaOpenCL/builtins-extended-image-err.cl b/clang/test/SemaOpenCL/builtins-extended-image-err.cl index f295d03d93fa1..4f057b1a2841f 100644 --- a/clang/test/SemaOpenCL/builtins-extended-image-err.cl +++ b/clang/test/SemaOpenCL/builtins-extended-image-err.cl @@ -12,224 +12,224 @@ typedef half half4 __attribute__((ext_vector_type(4))); float4 test_amdgcn_image_gather4_lz_2d_v4f32_f32_r(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_gather4_lz_2d_v4f32_f32_r' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_gather4_lz_2d_v4f32_f32_g(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_gather4_lz_2d_v4f32_f32_g' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_gather4_lz_2d_v4f32_f32_b(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_gather4_lz_2d_v4f32_f32_b' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_gather4_lz_2d_v4f32_f32_a(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_gather4_lz_2d_v4f32_f32_a' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_gather4_lz_2d_v4f16_f32_r(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_gather4_lz_2d_v4f16_f32_r' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_lz_1d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_1d_v4f32_f32(15, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_1d_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_1d_v4f32_f32(15, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_1d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_l_1d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_1d_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_1d_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_1d_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_1d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_d_1d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_1d_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_1d_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_1d_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_1d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_lz_2d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_2d_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_2d_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_2d_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_2d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_l_2d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_2d_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_2d_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_2d_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_2d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_d_2d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_2d_v4f32_f32(15, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_2d_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_2d_v4f32_f32(15, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_2d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_lz_3d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_3d_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_3d_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_3d_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_3d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_l_3d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_3d_v4f32_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_3d_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_3d_v4f32_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_3d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_d_3d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_3d_v4f32_f32(15, f32, f32, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_3d_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_3d_v4f32_f32(15, f32, f32, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_3d_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_lz_cube_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_cube_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_cube_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_cube_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_cube_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_l_cube_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_cube_v4f32_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_cube_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_cube_v4f32_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_cube_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_lz_1darray_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_1darray_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_1darray_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_1darray_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_l_1darray_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_1darray_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_1darray_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_1darray_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_1darray_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_d_1darray_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_1darray_v4f32_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_1darray_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_1darray_v4f32_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_1darray_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_lz_2darray_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_2darray_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_2darray_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_2darray_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_l_2darray_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_2darray_v4f32_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_2darray_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_2darray_v4f32_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_2darray_v4f32_f32' needs target feature extended-image-insts}} } float4 test_amdgcn_image_sample_d_2darray_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_2darray_v4f32_f32(15, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_2darray_v4f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_2darray_v4f32_f32(15, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_2darray_v4f32_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_lz_1d_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_1d_v4f16_f32(15, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_1d_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_1d_v4f16_f32(15, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_1d_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_l_1d_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_1d_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_1d_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_1d_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_1d_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_d_1d_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_1d_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_1d_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_1d_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_1d_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_lz_2d_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_2d_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_2d_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_2d_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_2d_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_l_2d_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_2d_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_2d_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_2d_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_2d_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_d_2d_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_2d_v4f16_f32(15, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_2d_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_2d_v4f16_f32(15, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_2d_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_lz_3d_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_3d_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_3d_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_3d_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_3d_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_l_3d_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_3d_v4f16_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_3d_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_3d_v4f16_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_3d_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_d_3d_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_3d_v4f16_f32(15, f32, f32, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_3d_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_3d_v4f16_f32(15, f32, f32, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_3d_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_lz_cube_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_cube_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_cube_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_cube_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_cube_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_l_cube_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_cube_v4f16_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_cube_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_cube_v4f16_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_cube_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_lz_1darray_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_1darray_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_1darray_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_1darray_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_l_1darray_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_1darray_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_1darray_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_1darray_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_1darray_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_d_1darray_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_1darray_v4f16_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_1darray_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_1darray_v4f16_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_1darray_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_lz_2darray_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_2darray_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_2darray_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_2darray_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_l_2darray_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_2darray_v4f16_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_2darray_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_2darray_v4f16_f32(15, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_2darray_v4f16_f32' needs target feature extended-image-insts}} } half4 test_amdgcn_image_sample_d_2darray_v4f16_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_2darray_v4f16_f32(15, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_2darray_v4f16_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_2darray_v4f16_f32(15, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_2darray_v4f16_f32' needs target feature extended-image-insts}} } float test_amdgcn_image_sample_lz_2d_f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_2d_f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_2d_f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_2d_f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_2d_f32_f32' needs target feature extended-image-insts}} } float test_amdgcn_image_sample_l_2d_f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_2d_f32_f32(1, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_2d_f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_2d_f32_f32(1, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_2d_f32_f32' needs target feature extended-image-insts}} } float test_amdgcn_image_sample_d_2d_f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_2d_f32_f32(1, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_2d_f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_2d_f32_f32(1, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_2d_f32_f32' needs target feature extended-image-insts}} } float test_amdgcn_image_sample_lz_2darray_f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_lz_2darray_f32_f32(1, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_2darray_f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_lz_2darray_f32_f32(1, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_lz_2darray_f32_f32' needs target feature extended-image-insts}} } float test_amdgcn_image_sample_l_2darray_f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_l_2darray_f32_f32(1, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_l_2darray_f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_l_2darray_f32_f32(1, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_l_2darray_f32_f32' needs target feature extended-image-insts}} } float test_amdgcn_image_sample_d_2darray_f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_d_2darray_f32_f32(1, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_d_2darray_f32_f32' needs target feature extended-image-insts}} + return __builtin_amdgcn_image_sample_d_2darray_f32_f32(1, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'__builtin_amdgcn_image_sample_d_2darray_f32_f32' needs target feature extended-image-insts}} } diff --git a/clang/test/SemaOpenCL/builtins-image-load-image-feature-err.cl b/clang/test/SemaOpenCL/builtins-image-load-image-feature-err.cl index 55635f84366bf..fad2fdeda9c8d 100644 --- a/clang/test/SemaOpenCL/builtins-image-load-image-feature-err.cl +++ b/clang/test/SemaOpenCL/builtins-image-load-image-feature-err.cl @@ -7,213 +7,213 @@ typedef half half4 __attribute__((ext_vector_type(4))); float test_builtin_image_load_2d(float f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_2d_f32_i32(1, i32, i32, tex, 106, 103); //expected-error{{'test_builtin_image_load_2d' needs target feature image-insts}} + return __builtin_amdgcn_image_load_2d_f32_i32(1, i32, i32, tex, 106, 103); //expected-error{{'__builtin_amdgcn_image_load_2d_f32_i32' needs target feature image-insts}} } float4 test_builtin_image_load_2d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_2d_v4f32_i32(15, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_load_2d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_2d_v4f32_i32(15, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_load_2d_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_2d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_2d_v4f16_i32(15, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_2d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_2d_v4f16_i32(15, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_2d_v4f16_i32' needs target feature image-insts}} } float test_builtin_image_load_2darray(float f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_2darray_f32_i32(1, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_load_2darray' needs target feature image-insts}} + return __builtin_amdgcn_image_load_2darray_f32_i32(1, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_load_2darray_f32_i32' needs target feature image-insts}} } float4 test_builtin_image_load_2darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_2darray_v4f32_i32(15, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_load_2darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_2darray_v4f32_i32(15, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_load_2darray_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_2darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_2darray_v4f16_i32(15, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_2darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_2darray_v4f16_i32(15, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_2darray_v4f16_i32' needs target feature image-insts}} } float4 test_builtin_image_load_1d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_1d_v4f32_i32(15, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_1d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_1d_v4f32_i32(15, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_1d_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_1d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_1d_v4f16_i32(15, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_1d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_1d_v4f16_i32(15, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_1d_v4f16_i32' needs target feature image-insts}} } float4 test_builtin_image_load_1darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_1darray_v4f32_i32(15, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_load_1darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_1darray_v4f32_i32(15, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_load_1darray_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_1darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_1darray_v4f16_i32(15, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_load_1darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_1darray_v4f16_i32(15, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_load_1darray_v4f16_i32' needs target feature image-insts}} } float4 test_builtin_image_load_3d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_3d_v4f32_i32(15, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_3d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_3d_v4f32_i32(15, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_3d_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_3d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_3d_v4f16_i32(15, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_3d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_3d_v4f16_i32(15, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_3d_v4f16_i32' needs target feature image-insts}} } float4 test_builtin_image_load_cube_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_cube_v4f32_i32(15, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_load_cube_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_cube_v4f32_i32(15, i32, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_load_cube_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_cube_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_cube_v4f16_i32(15, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_load_cube_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_cube_v4f16_i32(15, i32, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_load_cube_v4f16_i32' needs target feature image-insts}} } float4 test_builtin_image_load_mip_1d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_1d_v4f32_i32(15, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_mip_1d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_1d_v4f32_i32(15, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_mip_1d_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_mip_1d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_1d_v4f16_i32(15, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_mip_1d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_1d_v4f16_i32(15, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_mip_1d_v4f16_i32' needs target feature image-insts}} } float4 test_builtin_image_load_mip_1darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_1darray_v4f32_i32(15, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_load_mip_1darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_1darray_v4f32_i32(15, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_load_mip_1darray_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_mip_1darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_1darray_v4f16_i32(15, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_load_mip_1darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_1darray_v4f16_i32(15, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_load_mip_1darray_v4f16_i32' needs target feature image-insts}} } float test_builtin_image_load_mip_2d(float f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_2d_f32_i32(1, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_mip_2d' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_2d_f32_i32(1, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_mip_2d_f32_i32' needs target feature image-insts}} } float4 test_builtin_image_load_mip_2d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_2d_v4f32_i32(15, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_mip_2d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_2d_v4f32_i32(15, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_mip_2d_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_mip_2d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_2d_v4f16_i32(15, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_load_mip_2d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_2d_v4f16_i32(15, i32, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_load_mip_2d_v4f16_i32' needs target feature image-insts}} } float test_builtin_image_load_mip_2darray(float f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_2darray_f32_i32(1, i32, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_load_mip_2darray' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_2darray_f32_i32(1, i32, i32, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_load_mip_2darray_f32_i32' needs target feature image-insts}} } float4 test_builtin_image_load_mip_2darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_2darray_v4f32_i32(15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_mip_2darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_2darray_v4f32_i32(15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_mip_2darray_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_mip_2darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_2darray_v4f16_i32(15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_mip_2darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_2darray_v4f16_i32(15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_mip_2darray_v4f16_i32' needs target feature image-insts}} } float4 test_builtin_image_load_mip_3d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_3d_v4f32_i32(15, i32, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_load_mip_3d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_3d_v4f32_i32(15, i32, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_load_mip_3d_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_mip_3d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_3d_v4f16_i32(15, i32, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_load_mip_3d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_3d_v4f16_i32(15, i32, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_load_mip_3d_v4f16_i32' needs target feature image-insts}} } float4 test_builtin_image_load_mip_cube_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_cube_v4f32_i32(15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_mip_cube_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_cube_v4f32_i32(15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_mip_cube_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_mip_cube_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_mip_cube_v4f16_i32(15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_load_mip_cube_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_mip_cube_v4f16_i32(15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_load_mip_cube_v4f16_i32' needs target feature image-insts}} } float test_builtin_image_load_2d_gfx(float f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_2d_f32_i32(1, i32, i32, tex, 106, 103); //expected-error{{'test_builtin_image_load_2d_gfx' needs target feature image-insts}} + return __builtin_amdgcn_image_load_2d_f32_i32(1, i32, i32, tex, 106, 103); //expected-error{{'__builtin_amdgcn_image_load_2d_f32_i32' needs target feature image-insts}} } float4 test_builtin_image_load_2d_gfx_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_2d_v4f32_i32(15, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_load_2d_gfx_1' needs target feature image-insts}} + return __builtin_amdgcn_image_load_2d_v4f32_i32(15, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_load_2d_v4f32_i32' needs target feature image-insts}} } half4 test_builtin_image_load_2d_gfx_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_load_2d_v4f16_i32(15, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_load_2d_gfx_2' needs target feature image-insts}} + return __builtin_amdgcn_image_load_2d_v4f16_i32(15, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_load_2d_v4f16_i32' needs target feature image-insts}} } float test_builtin_image_sample_2d(float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_2d_f32_f32(1, f32, f32, tex, vec4i32, 0, 106, 103); //expected-error{{'test_builtin_image_sample_2d' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_2d_f32_f32(1, f32, f32, tex, vec4i32, 0, 106, 103); //expected-error{{'__builtin_amdgcn_image_sample_2d_f32_f32' needs target feature image-insts}} } float4 test_builtin_image_sample_2d_1(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_2d_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'test_builtin_image_sample_2d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_2d_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'__builtin_amdgcn_image_sample_2d_v4f32_f32' needs target feature image-insts}} } half4 test_builtin_image_sample_2d_2(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_2d_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'test_builtin_image_sample_2d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_2d_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'__builtin_amdgcn_image_sample_2d_v4f16_f32' needs target feature image-insts}} } float test_builtin_image_sample_2darray(float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_2darray_f32_f32(1, f32, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'test_builtin_image_sample_2darray' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_2darray_f32_f32(1, f32, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'__builtin_amdgcn_image_sample_2darray_f32_f32' needs target feature image-insts}} } float4 test_builtin_image_sample_2darray_1(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_2darray_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'test_builtin_image_sample_2darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_2darray_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'__builtin_amdgcn_image_sample_2darray_v4f32_f32' needs target feature image-insts}} } half4 test_builtin_image_sample_2darray_2(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_2darray_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'test_builtin_image_sample_2darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_2darray_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'__builtin_amdgcn_image_sample_2darray_v4f16_f32' needs target feature image-insts}} } float4 test_builtin_image_sample_1d_1(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_1d_v4f32_f32(15, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'test_builtin_image_sample_1d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_1d_v4f32_f32(15, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'__builtin_amdgcn_image_sample_1d_v4f32_f32' needs target feature image-insts}} } half4 test_builtin_image_sample_1d_2(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_1d_v4f16_f32(15, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'test_builtin_image_sample_1d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_1d_v4f16_f32(15, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'__builtin_amdgcn_image_sample_1d_v4f16_f32' needs target feature image-insts}} } float4 test_builtin_image_sample_1darray_1(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_1darray_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'test_builtin_image_sample_1darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_1darray_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'__builtin_amdgcn_image_sample_1darray_v4f32_f32' needs target feature image-insts}} } half4 test_builtin_image_sample_1darray_2(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_1darray_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'test_builtin_image_sample_1darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_1darray_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 0, 110); //expected-error{{'__builtin_amdgcn_image_sample_1darray_v4f16_f32' needs target feature image-insts}} } float4 test_builtin_image_sample_3d_1(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_3d_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'test_builtin_image_sample_3d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_3d_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'__builtin_amdgcn_image_sample_3d_v4f32_f32' needs target feature image-insts}} } half4 test_builtin_image_sample_3d_2(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_3d_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'test_builtin_image_sample_3d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_3d_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 0); //expected-error{{'__builtin_amdgcn_image_sample_3d_v4f16_f32' needs target feature image-insts}} } float4 test_builtin_image_sample_cube_1(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_cube_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'test_builtin_image_sample_cube_1' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_cube_v4f32_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'__builtin_amdgcn_image_sample_cube_v4f32_f32' needs target feature image-insts}} } half4 test_builtin_image_sample_cube_2(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_cube_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'test_builtin_image_sample_cube_2' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_cube_v4f16_f32(15, f32, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'__builtin_amdgcn_image_sample_cube_v4f16_f32' needs target feature image-insts}} } float test_builtin_image_sample_2d_gfx(float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_2d_f32_f32(1, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'test_builtin_image_sample_2d_gfx' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_2d_f32_f32(1, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'__builtin_amdgcn_image_sample_2d_f32_f32' needs target feature image-insts}} } float4 test_builtin_image_sample_2d_gfx_1(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_2d_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'test_builtin_image_sample_2d_gfx_1' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_2d_v4f32_f32(15, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'__builtin_amdgcn_image_sample_2d_v4f32_f32' needs target feature image-insts}} } half4 test_builtin_image_sample_2d_gfx_2(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) { - return __builtin_amdgcn_image_sample_2d_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'test_builtin_image_sample_2d_gfx_2' needs target feature image-insts}} + return __builtin_amdgcn_image_sample_2d_v4f16_f32(15, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{'__builtin_amdgcn_image_sample_2d_v4f16_f32' needs target feature image-insts}} } diff --git a/clang/test/SemaOpenCL/builtins-image-store-image-feature-err.cl b/clang/test/SemaOpenCL/builtins-image-store-image-feature-err.cl index 5bc78e123d333..06f57d01f9336 100644 --- a/clang/test/SemaOpenCL/builtins-image-store-image-feature-err.cl +++ b/clang/test/SemaOpenCL/builtins-image-store-image-feature-err.cl @@ -6,124 +6,124 @@ typedef half half4 __attribute__((ext_vector_type(4))); void test_builtin_image_store_2d(float f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_2d_f32_i32(f32, 15, i32, i32, tex, 106, 103); //expected-error{{'test_builtin_image_store_2d' needs target feature image-insts}} + return __builtin_amdgcn_image_store_2d_f32_i32(f32, 15, i32, i32, tex, 106, 103); //expected-error{{'__builtin_amdgcn_image_store_2d_f32_i32' needs target feature image-insts}} } void test_builtin_image_store_2d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_2d_v4f32_i32(v4f32, 15, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_store_2d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_2d_v4f32_i32(v4f32, 15, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_store_2d_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_2d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_2d_v4f16_i32(v4f16, 15, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_2d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_2d_v4f16_i32(v4f16, 15, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_2d_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_2darray(float f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_2darray_f32_i32(f32, 15, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_store_2darray' needs target feature image-insts}} + return __builtin_amdgcn_image_store_2darray_f32_i32(f32, 15, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_store_2darray_f32_i32' needs target feature image-insts}} } void test_builtin_image_store_2darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_2darray_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_store_2darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_2darray_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_store_2darray_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_2darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_2darray_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_2darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_2darray_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_2darray_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_1d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_1d_v4f32_i32(v4f32, 15, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_1d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_1d_v4f32_i32(v4f32, 15, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_1d_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_1d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_1d_v4f16_i32(v4f16, 15, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_1d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_1d_v4f16_i32(v4f16, 15, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_1d_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_1darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_1darray_v4f32_i32(v4f32, 15, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_store_1darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_1darray_v4f32_i32(v4f32, 15, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_store_1darray_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_1darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_1darray_v4f16_i32(v4f16, 15, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_store_1darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_1darray_v4f16_i32(v4f16, 15, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_store_1darray_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_3d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_3d_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_3d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_3d_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_3d_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_3d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_3d_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_3d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_3d_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_3d_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_cube_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_cube_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_store_cube_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_cube_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_store_cube_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_cube_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_cube_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_store_cube_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_cube_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_store_cube_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_1d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_1d_v4f32_i32(v4f32, 15, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_mip_1d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_1d_v4f32_i32(v4f32, 15, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_mip_1d_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_1d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_1d_v4f16_i32(v4f16, 15, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_mip_1d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_1d_v4f16_i32(v4f16, 15, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_mip_1d_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_1darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_1darray_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_store_mip_1darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_1darray_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_store_mip_1darray_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_1darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_1darray_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_store_mip_1darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_1darray_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_store_mip_1darray_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_2d(float f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_2d_f32_i32(f32, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_mip_2d' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_2d_f32_i32(f32, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_mip_2d_f32_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_2d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_2d_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_mip_2d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_2d_v4f32_i32(v4f32, 15, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_mip_2d_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_2d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_2d_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_store_mip_2d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_2d_v4f16_i32(v4f16, 15, i32, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_store_mip_2d_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_2darray(float f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_2darray_f32_i32(f32, 15, i32, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_store_mip_2darray' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_2darray_f32_i32(f32, 15, i32, i32, i32, i32, tex, 120, 110); //expected-error{{'__builtin_amdgcn_image_store_mip_2darray_f32_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_2darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_2darray_v4f32_i32(v4f32, 15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_mip_2darray_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_2darray_v4f32_i32(v4f32, 15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_mip_2darray_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_2darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_2darray_v4f16_i32(v4f16, 15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_mip_2darray_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_2darray_v4f16_i32(v4f16, 15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_mip_2darray_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_3d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_3d_v4f32_i32(v4f32, 15, i32, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_store_mip_3d_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_3d_v4f32_i32(v4f32, 15, i32, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_store_mip_3d_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_3d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_3d_v4f16_i32(v4f16, 15, i32, i32, i32, i32, tex, 0, 110); //expected-error{{'test_builtin_image_store_mip_3d_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_3d_v4f16_i32(v4f16, 15, i32, i32, i32, i32, tex, 0, 110); //expected-error{{'__builtin_amdgcn_image_store_mip_3d_v4f16_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_cube_1(float4 v4f32, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_cube_v4f32_i32(v4f32, 15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_mip_cube_1' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_cube_v4f32_i32(v4f32, 15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_mip_cube_v4f32_i32' needs target feature image-insts}} } void test_builtin_image_store_mip_cube_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { - return __builtin_amdgcn_image_store_mip_cube_v4f16_i32(v4f16, 15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'test_builtin_image_store_mip_cube_2' needs target feature image-insts}} + return __builtin_amdgcn_image_store_mip_cube_v4f16_i32(v4f16, 15, i32, i32, i32, i32, tex, 120, 0); //expected-error{{'__builtin_amdgcn_image_store_mip_cube_v4f16_i32' needs target feature image-insts}} } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
