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

Rana Pratap Reddy via cfe-commits Thu, 26 Jun 2025 02:09:22 -0700

================
@@ -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;
----------------
ranapratap55 wrote:


updated to `nullptr`

https://github.com/llvm/llvm-project/pull/140210
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

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

Reply via email to