[AMDGPU] Changes image desc macro to use opaque ptr type macro and minor changes

Author: ranapratap55

clang/include/clang/Basic/AMDGPUTypes.def

@@ -20,17 +20,11 @@
   AMDGPU_TYPE(Name, Id, SingletonId, Width, Align)
 #endif
 
-#ifndef AMDGPU_IMAGE_RSRC_TYPE
-#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId) \
-  AMDGPU_TYPE(Name, Id, SingletonId, 256, 256)
-#endif
-
 AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_buffer_rsrc_t", AMDGPUBufferRsrc, AMDGPUBufferRsrcTy, 128, 128, 8)
-AMDGPU_IMAGE_RSRC_TYPE("__amdgpu_image_rsrc_t", AMDGPUImageDescRsrc, AMDGPUImageDescRsrcTy)
+AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_texture_t", AMDGPUTexture, AMDGPUTextureTy, 256, 256, 0)
 
 AMDGPU_NAMED_BARRIER_TYPE("__amdgpu_named_workgroup_barrier_t", AMDGPUNamedWorkgroupBarrier, AMDGPUNamedWorkgroupBarrierTy, 128, 32, 0)
 
 #undef AMDGPU_TYPE
 #undef AMDGPU_OPAQUE_PTR_TYPE
 #undef AMDGPU_NAMED_BARRIER_TYPE
-#undef AMDGPU_IMAGE_RSRC_TYPE

clang/include/clang/Basic/Builtins.def

@@ -34,7 +34,7 @@
 //  Q -> target builtin type, followed by a character to distinguish the builtin type
 //    Qa -> AArch64 svcount_t builtin type.
 //    Qb -> AMDGPU __amdgpu_buffer_rsrc_t builtin type.
-//    Qt -> AMDGPU __amdgpu_image_desc_t builtin type.
+//    Qt -> AMDGPU __amdgpu_texture_t builtin type.
 //  E -> ext_vector, followed by the number of elements and the base type.
 //  X -> _Complex, followed by the base type.
 //  Y -> ptrdiff_t

clang/lib/AST/ASTContext.cpp

@@ -12581,7 +12581,7 @@ static QualType DecodeTypeFromStr(const char *&Str, const ASTContext &Context,
       break;
     }
     case 't': {
-      Type = Context.AMDGPUImageDescRsrcTy;
+      Type = Context.AMDGPUTextureTy;
       break;
     }
     default:

clang/lib/CodeGen/CGDebugInfo.cpp

@@ -1020,14 +1020,6 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) {
           DBuilder.createBasicType(Name, Width, llvm::dwarf::DW_ATE_unsigned); \
     return SingletonId;                                                        \
   }
-#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId)                          \
-  case BuiltinType::Id: {                                                      \
-    if (!SingletonId)                                                          \
-      SingletonId =                                                            \
-          DBuilder.createForwardDecl(llvm::dwarf::DW_TAG_structure_type, Name, \
-                                     TheCU, TheCU->getFile(), 0);              \
-    return SingletonId;                                                        \
-  }
 #include "clang/Basic/AMDGPUTypes.def"
   case BuiltinType::UChar:
   case BuiltinType::Char_U:

clang/lib/CodeGen/CodeGenTypes.cpp

@@ -581,10 +581,6 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
   case BuiltinType::Id:                                                        \
     return llvm::TargetExtType::get(getLLVMContext(), "amdgcn.named.barrier",  \
                                     {}, {Scope});
-#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId)                          \
-  case BuiltinType::Id:                                                        \
-    return llvm::VectorType::get(llvm::Type::getInt32Ty(getLLVMContext()), 8,  \
-                                 false);
 #include "clang/Basic/AMDGPUTypes.def"
 #define HLSL_INTANGIBLE_TYPE(Name, Id, SingletonId) case BuiltinType::Id:
 #include "clang/Basic/HLSLIntangibleTypes.def"

clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c

@@ -5,14 +5,13 @@
 // CHECK-LABEL: define dso_local void @test_locals(
 // CHECK-SAME: ) #[[ATTR0:[0-9]+]] !dbg [[DBG6:![0-9]+]] {
 // CHECK-NEXT:  [[ENTRY:.*:]]
-// CHECK-NEXT:    [[IMG:%.*]] = alloca <8 x i32>, align 32, addrspace(5)
+// CHECK-NEXT:    [[IMG:%.*]] = alloca ptr, align 32, addrspace(5)
 // CHECK-NEXT:    [[IMG_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IMG]] to ptr
 // CHECK-NEXT:      #dbg_declare(ptr addrspace(5) [[IMG]], [[META11:![0-9]+]], !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef), [[META14:![0-9]+]])
-// CHECK-NEXT:    [[TMP0:%.*]] = load <8 x i32>, ptr [[IMG_ASCAST]], align 32, !dbg [[DBG15:![0-9]+]]
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[IMG_ASCAST]], align 32, !dbg [[DBG15:![0-9]+]]
 // CHECK-NEXT:    ret void, !dbg [[DBG16:![0-9]+]]
 //
 void test_locals(void) {
-  __amdgpu_image_rsrc_t img;
+  __amdgpu_texture_t img;
   (void)img;
 }
-

clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp

@@ -2,6 +2,6 @@
 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s
 namespace std { class type_info; }
-auto &a = typeid(__amdgpu_image_rsrc_t);
+auto &a = typeid(__amdgpu_texture_t);
 //// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
 // CHECK: {{.*}}

clang/test/SemaCXX/amdgpu-image-rsrc.cpp

@@ -3,19 +3,15 @@
 // RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn -Wno-unused-value %s
 
 void foo() {
-  int n = 1;
-  __amdgpu_image_rsrc_t v = 0;            // expected-error {{cannot initialize a variable of type '__amdgpu_image_rsrc_t' with an rvalue of type 'int'}}
-  static_cast<__amdgpu_image_rsrc_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_image_rsrc_t' is not allowed}}
-  reinterpret_cast<__amdgpu_image_rsrc_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_image_rsrc_t' is not allowed}}
-  (void)(v + v); // expected-error {{invalid operands}}
-  int x(v);      // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_image_rsrc_t'}}
-  __amdgpu_image_rsrc_t k;
+  int n = 100;
+  __amdgpu_texture_t v = 0;           // expected-error {{cannot initialize a variable of type '__amdgpu_texture_t' with an rvalue of type 'int'}}
+  static_cast<__amdgpu_texture_t>(n);  // expected-error {{static_cast from 'int' to '__amdgpu_texture_t' is not allowed}}
+  reinterpret_cast<__amdgpu_texture_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_texture_t' is not allowed}}
+  (void)(v + v); // expected-error {{invalid operands to binary expression ('__amdgpu_texture_t' and '__amdgpu_texture_t')}}
+  int x(v);      // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_texture_t'}}
+  __amdgpu_texture_t k;
 }
 
-static_assert(sizeof(__amdgpu_image_rsrc_t) == 32, "size");
-static_assert(alignof(__amdgpu_image_rsrc_t) == 32, "align");
-
 template<class T> void bar(T);
-void use(__amdgpu_image_rsrc_t r) { bar(r); }
-struct S { __amdgpu_image_rsrc_t r; int a; };
-static_assert(sizeof(S) == 64, "struct layout");
+void use(__amdgpu_texture_t r) { bar(r); }
+struct S { __amdgpu_texture_t r; int a; };

clang/test/SemaOpenCL/amdgpu-image-rsrc.cl

@@ -5,10 +5,9 @@
 
 void f() {
     int n = 3;
-    __amdgpu_image_rsrc_t v = 0; // expected-error {{initializing '__private __amdgpu_image_rsrc_t' with an expression of incompatible type 'int'}}
-    int k = v;                   // expected-error {{initializing '__private int' with an expression of incompatible type '__private __amdgpu_image_rsrc_t'}}
-    (void)(v + v);               // expected-error {{invalid operands}}
-    __amdgpu_image_rsrc_t r;
-    int *p = (int*)r;            // expected-error {{operand of type '__amdgpu_image_rsrc_t' where arithmetic or pointer type is required}}
-    (void)p;
+    __amdgpu_texture_t v = (__amdgpu_texture_t)0; // expected-error {{used type '__amdgpu_texture_t' where arithmetic or pointer type is required}}
+    int k = v;          // expected-error {{initializing '__private int' with an expression of incompatible type '__private __amdgpu_texture_t'}}
+    (void)(v + v);      // expected-error {{invalid operands}}
+    __amdgpu_texture_t r;
+    int *p = (int*)r;   // expected-error {{operand of type '__amdgpu_texture_t' where arithmetic or pointer type is required}}
 }

clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp

@@ -6,7 +6,7 @@ void foo() {
 #pragma omp target
   {
     int n = 5;
-    __amdgpu_image_rsrc_t v = 0; // expected-error {{cannot initialize a variable of type '__amdgpu_image_rsrc_t' with an rvalue of type 'int'}}
-    (void)(v + v);               // expected-error {{invalid operands to binary expression ('__amdgpu_image_rsrc_t' and '__amdgpu_image_rsrc_t'}}
+    __amdgpu_texture_t v = 0; // expected-error {{cannot initialize a variable of type '__amdgpu_texture_t' with an rvalue of type 'int'}}
+    (void)(v + v);            // expected-error {{invalid operands to binary expression}}
   }
 }