[AMDGPU] Adds target failure test cases and minor modificatinos

Author: ranapratap55

clang/lib/Sema/SemaAMDGPU.cpp

@@ -112,11 +112,12 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
+    StringRef FeatureList(getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
     bool HasImageInsts = Builtin::evaluateRequiredTargetFeatures("image-insts", CallerFeatureMap);
-
-    if(!HasImageInsts){
+    if (!Builtin::evaluateRequiredTargetFeatures(
+        FeatureList, CallerFeatureMap) && !HasImageInsts){
       Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
-        << FD->getDeclName() << "image-insts";
+        << FD->getDeclName() << FeatureList;
       return false;
     }
 
@@ -155,6 +156,15 @@ 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));
+    bool HasImageInsts = Builtin::evaluateRequiredTargetFeatures("image-insts", CallerFeatureMap);
+    if (!Builtin::evaluateRequiredTargetFeatures(
+        FeatureList, CallerFeatureMap) && !HasImageInsts){
+      Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
+        << FD->getDeclName() << FeatureList;
+      return false;
+    }
+
     unsigned ArgCount = TheCall->getNumArgs() - 1;
     llvm::APSInt Result;
 

clang/test/SemaOpenCL/builtins-image-load-param.cl

@@ -1,5 +1,6 @@
 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 -S -verify=expected -o - %s
 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -S -verify=expected -o - %s
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx942 -S -verify=GFX94 -o - %s
 // REQUIRES: amdgpu-registered-target
 
 typedef int int8 __attribute__((ext_vector_type(8)));
@@ -129,3 +130,16 @@ half4 test_builtin_image_load_mip_cube_2(half4 v4f16, int i32, int8 vec8i32) {
 
   return __builtin_amdgcn_image_load_mip_cube_v4f16_i32(100, i32, i32, i32, i32, vec8i32, 120, i32); //expected-error{{argument to '__builtin_amdgcn_image_load_mip_cube_v4f16_i32' must be a constant integer}}
 }
+
+float test_builtin_image_load_2d_gfx(float f32, int i32, int8 vec8i32) {
+
+  return __builtin_amdgcn_image_load_2d_f32_i32(12, i32, i32, vec8i32, 106, 103); //GFX94-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, int8 vec8i32) {
+
+  return __builtin_amdgcn_image_load_2d_v4f32_i32(100, i32, i32, vec8i32, 120, 110); //GFX94-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, int8 vec8i32) {
+
+  return __builtin_amdgcn_image_load_2d_v4f16_i32(100, i32, i32, vec8i32, 120, 110); //GFX94-error{{'__builtin_amdgcn_image_load_2d_v4f16_i32' needs target feature image-insts}}
+}

clang/test/SemaOpenCL/builtins-image-store-param.cl

@@ -1,5 +1,5 @@
 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -S -verify=expected -o - %s
-// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx701 -S -verify=GFX7 -o - %s
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx942 -S -verify=GFX94 -o - %s
 // REQUIRES: amdgpu-registered-target
 
 typedef int int8 __attribute__((ext_vector_type(8)));
@@ -130,3 +130,16 @@ void test_builtin_image_store_mip_cube_2(half4 v4f16, int i32, int8 vec8i32) {
 
   return __builtin_amdgcn_image_store_mip_cube_v4f16_i32(v4f16, 100, i32, i32, i32, i32, vec8i32, 120, i32); //expected-error{{argument to '__builtin_amdgcn_image_store_mip_cube_v4f16_i32' must be a constant integer}}
 }
+
+void test_builtin_image_store_2d_gfx(float f32, int i32, int8 vec8i32) {
+
+   __builtin_amdgcn_image_store_2d_f32_i32(f32, 12, i32, i32, vec8i32, 106, 103); //GFX94-error{{'__builtin_amdgcn_image_store_2d_f32_i32' needs target feature image-insts}}
+}
+void test_builtin_image_store_2d_gfx_1(float4 v4f32, int i32, int8 vec8i32) {
+
+   __builtin_amdgcn_image_store_2d_v4f32_i32(v4f32, 100, i32, i32, vec8i32, 120, 110); //GFX94-error{{'__builtin_amdgcn_image_store_2d_v4f32_i32' needs target feature image-insts}}
+ }
+ void test_builtin_image_store_2d_gfx_2(half4 v4f16, int i32, int8 vec8i32) {
+
+   __builtin_amdgcn_image_store_2d_v4f16_i32(v4f16, 100, i32, i32, vec8i32, 120, 110); //GFX94-error{{'__builtin_amdgcn_image_store_2d_v4f16_i32' needs target feature image-insts}}
+ }