[clang] [CIR][AMDGPU][NFC] Add CIRGenBuiltinAMDGPU file to support AMDGPU builtins codegen (PR #179237)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Feb 2 21:11:20 PST 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
@llvm/pr-subscribers-clangir
Author: Chaitanya (skc7)
<details>
<summary>Changes</summary>
Upstreaming ClangIR PR : https://github.com/llvm/clangir/pull/2030
This PR adds CIRGenBuiltinAMDGPU.cpp file for AMDGPU specific builtin codegen support.
Lists out all the builtins that are currently supported for codegen in `clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp`.
All builtins codegen are currently "NYI".
---
Patch is 31.00 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/179237.diff
4 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp (+1)
- (added) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+609)
- (modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+3)
- (modified) clang/lib/CIR/CodeGen/CMakeLists.txt (+1)
``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
index 88d37d56fcd78..e9d5f8d688c65 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
@@ -1885,6 +1885,7 @@ emitTargetArchBuiltinExpr(CIRGenFunction *cgf, unsigned builtinID,
case llvm::Triple::ppc64le:
case llvm::Triple::r600:
case llvm::Triple::amdgcn:
+ return cgf->emitAMDGPUBuiltinExpr(builtinID, e);
case llvm::Triple::systemz:
case llvm::Triple::nvptx:
case llvm::Triple::nvptx64:
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
new file mode 100644
index 0000000000000..438f0e1317786
--- /dev/null
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -0,0 +1,609 @@
+//===---- CIRGenBuiltinAMDGPU.cpp - Emit CIR for AMDGPU builtins ----------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This contains code to emit AMDGPU Builtin calls.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CIRGenFunction.h"
+
+#include "mlir/IR/Value.h"
+#include "clang/Basic/TargetBuiltins.h"
+#include "llvm/Support/ErrorHandling.h"
+
+using namespace clang;
+using namespace clang::CIRGen;
+using namespace cir;
+
+mlir::Value CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
+ const CallExpr *expr) {
+ switch (builtinId) {
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64: {
+ llvm_unreachable("wave_reduce_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_div_scale:
+ case AMDGPU::BI__builtin_amdgcn_div_scalef: {
+ llvm_unreachable("div_scale_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_div_fmas:
+ case AMDGPU::BI__builtin_amdgcn_div_fmasf: {
+ llvm_unreachable("div_fmas_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
+ case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
+ case AMDGPU::BI__builtin_amdgcn_mov_dpp:
+ case AMDGPU::BI__builtin_amdgcn_update_dpp: {
+ llvm_unreachable("mov_dpp_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_permlane16:
+ case AMDGPU::BI__builtin_amdgcn_permlanex16:
+ case AMDGPU::BI__builtin_amdgcn_permlane64: {
+ llvm_unreachable("permlane_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_readlane:
+ case AMDGPU::BI__builtin_amdgcn_readfirstlane: {
+ llvm_unreachable("readlane_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_div_fixup:
+ case AMDGPU::BI__builtin_amdgcn_div_fixupf:
+ case AMDGPU::BI__builtin_amdgcn_div_fixuph: {
+ llvm_unreachable("div_fixup_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_trig_preop:
+ case AMDGPU::BI__builtin_amdgcn_trig_preopf: {
+ llvm_unreachable("trig_preop_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_rcp:
+ case AMDGPU::BI__builtin_amdgcn_rcpf:
+ case AMDGPU::BI__builtin_amdgcn_rcph:
+ case AMDGPU::BI__builtin_amdgcn_rcp_bf16: {
+ llvm_unreachable("rcp_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_sqrt:
+ case AMDGPU::BI__builtin_amdgcn_sqrtf:
+ case AMDGPU::BI__builtin_amdgcn_sqrth:
+ case AMDGPU::BI__builtin_amdgcn_sqrt_bf16: {
+ llvm_unreachable("sqrt_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_rsq:
+ case AMDGPU::BI__builtin_amdgcn_rsqf:
+ case AMDGPU::BI__builtin_amdgcn_rsqh:
+ case AMDGPU::BI__builtin_amdgcn_rsq_bf16: {
+ llvm_unreachable("rsq_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
+ case AMDGPU::BI__builtin_amdgcn_rsq_clampf: {
+ llvm_unreachable("rsq_clamp_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_sinf:
+ case AMDGPU::BI__builtin_amdgcn_sinh:
+ case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
+ llvm_unreachable("sinf_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_cosf:
+ case AMDGPU::BI__builtin_amdgcn_cosh:
+ case AMDGPU::BI__builtin_amdgcn_cos_bf16: {
+ llvm_unreachable("cosf_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_dispatch_ptr: {
+ llvm_unreachable("dispatch_ptr_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_logf:
+ case AMDGPU::BI__builtin_amdgcn_log_bf16: {
+ llvm_unreachable("logf_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_exp2f:
+ case AMDGPU::BI__builtin_amdgcn_exp2_bf16: {
+ llvm_unreachable("exp2f_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_log_clampf: {
+ llvm_unreachable("log_clampf_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_ldexp:
+ case AMDGPU::BI__builtin_amdgcn_ldexpf:
+ case AMDGPU::BI__builtin_amdgcn_ldexph: {
+ llvm_unreachable("ldexp_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_frexp_mant:
+ case AMDGPU::BI__builtin_amdgcn_frexp_mantf:
+ case AMDGPU::BI__builtin_amdgcn_frexp_manth: {
+ llvm_unreachable("frexp_mant_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_frexp_exp:
+ case AMDGPU::BI__builtin_amdgcn_frexp_expf:
+ case AMDGPU::BI__builtin_amdgcn_frexp_exph: {
+ llvm_unreachable("frexp_exp_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_fract:
+ case AMDGPU::BI__builtin_amdgcn_fractf:
+ case AMDGPU::BI__builtin_amdgcn_fracth: {
+ llvm_unreachable("fract_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_lerp: {
+ llvm_unreachable("lerp_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_ubfe: {
+ llvm_unreachable("ubfe_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_sbfe: {
+ llvm_unreachable("sbfe_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_ballot_w32:
+ case AMDGPU::BI__builtin_amdgcn_ballot_w64: {
+ llvm_unreachable("ballot_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32:
+ case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64: {
+ llvm_unreachable("inverse_ballot_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_tanhf:
+ case AMDGPU::BI__builtin_amdgcn_tanhh:
+ case AMDGPU::BI__builtin_amdgcn_tanh_bf16: {
+ llvm_unreachable("tanh_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_uicmp:
+ case AMDGPU::BI__builtin_amdgcn_uicmpl:
+ case AMDGPU::BI__builtin_amdgcn_sicmp:
+ case AMDGPU::BI__builtin_amdgcn_sicmpl: {
+ llvm_unreachable("uicmp_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_fcmp:
+ case AMDGPU::BI__builtin_amdgcn_fcmpf: {
+ llvm_unreachable("fcmp_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_class:
+ case AMDGPU::BI__builtin_amdgcn_classf:
+ case AMDGPU::BI__builtin_amdgcn_classh: {
+ llvm_unreachable("class_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_fmed3f:
+ case AMDGPU::BI__builtin_amdgcn_fmed3h: {
+ llvm_unreachable("fmed3_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_ds_append:
+ case AMDGPU::BI__builtin_amdgcn_ds_consume: {
+ llvm_unreachable("ds_append_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:
+ case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16: {
+ llvm_unreachable("global_load_tr_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:
+ case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:
+ case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:
+ case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:
+ case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:
+ case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16: {
+ llvm_unreachable("ds_load_tr_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:
+ case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:
+ case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:
+ case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:
+ case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:
+ case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: {
+ llvm_unreachable("ds_read_tr_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:
+ case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:
+ case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:
+ case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:
+ case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:
+ case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: {
+ llvm_unreachable("global_load_monitor_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
+ case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
+ case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: {
+ llvm_unreachable("cluster_load_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_load_to_lds: {
+ llvm_unreachable("load_to_lds_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
+ case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
+ case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
+ case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
+ case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
+ case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B: {
+ llvm_unreachable("cooperative_atomic_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_get_fpenv:
+ case AMDGPU::BI__builtin_amdgcn_set_fpenv: {
+ llvm_unreachable("fpenv_* builtins NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_read_exec:
+ case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
+ case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
+ llvm_unreachable("read_exec_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:
+ case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:
+ case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l:
+ case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: {
+ llvm_unreachable("image_bvh_intersect_ray_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray:
+ case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: {
+ llvm_unreachable("image_bvh8_intersect_ray_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn:
+ case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:
+ case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:
+ case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: {
+ llvm_unreachable("ds_bvh_stack_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_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_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_unreachable("image_load_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
+ 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: {
+ llvm_unreachable("image_store_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32: {
+ llvm_unreachable("image_sample_d_2darray_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32: {
+ llvm_unreachable("image_gather4_lz_2d_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
+ case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
+ llvm_unreachable("mfma_scale_* NYI");
+ }
+ case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:
+ case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:
+ case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:
+ case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:
+ case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:
+ case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:
+ case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:
+ case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:
+ case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:
+ case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:
+ case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:
+ case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:
+ case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:
+ case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:
+ case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:
+ case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:
+ case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:
+ case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:
+ case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:
+ case A...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/179237
More information about the cfe-commits
mailing list