[clang] [WIP][AMDGPU] Support for type inferring image load/store builtins for AMDGPU (PR #140210)

[Shilei Tian via cfe-commits]

================
@@ -683,6 +683,206 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
 
     return Builder.CreateInsertElement(I0, A, 1);
   }
+  case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
+  case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
+  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: {
+    llvm::Type *RetTy = nullptr;
+    switch (BuiltinID) {
+    case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
+      RetTy = llvm::Type::getFloatTy(Builder.getContext());
+      break;
+    case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
+      RetTy =
+          FixedVectorType::get(llvm::Type::getFloatTy(Builder.getContext()), 4);
+      break;
+    case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
+      RetTy =
+          FixedVectorType::get(llvm::Type::getHalfTy(Builder.getContext()), 4);
+      break;
+    }
+
+    llvm::Value *Dmask = EmitScalarExpr(E->getArg(0));
+    llvm::Value *S = EmitScalarExpr(E->getArg(1));
+    llvm::Value *T = EmitScalarExpr(E->getArg(2));
+    llvm::Value *Slice;
+    llvm::Value *Mip;
+    llvm::Value *Rsrc;
+    llvm::Value *Tfe;
+    llvm::Value *Cpol;
+
+    SmallVector<Value *, 10> ArgTys;
+
+    Intrinsic::ID IID;
+    llvm::CallInst *Call;
+
+    switch (BuiltinID) {
+    case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
+    case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32: {
+      Rsrc = EmitScalarExpr(E->getArg(2));
+      Tfe = EmitScalarExpr(E->getArg(3));
+      Cpol = EmitScalarExpr(E->getArg(4));
+
+      ArgTys = {Dmask, S, Rsrc, Tfe, Cpol};
+      IID = Intrinsic::amdgcn_image_load_1d;
+      Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
+      break;
+    }
+    case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
----------------
shiltian wrote:

can you reorganize the code to avoid going through the switch twice here?

https://github.com/llvm/llvm-project/pull/140210