Skip to content

[AMDGPU] Support for type inferring image load/store builtins for AMDGPU #140210

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 8 commits into
base: main
Choose a base branch
from
+2,982 −0
Open

[AMDGPU] Support for type inferring image load/store builtins for AMDGPU #140210

Show file tree
Hide file tree
Changes from 7 commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
4fbfa90
[WIP][AMDGPU] Support for type inferring image load/store builtins fo…
ranapratap55 May 16, 2025
90065ee
[AMDGPU] Adds builtins for image load and sema checking for image load
ranapratap55 Jun 18, 2025
21c6829
[AMDGPU] Adds builtins for image store and sema checking for image store
ranapratap55 Jun 26, 2025
78caf4f
[AMDGPU] Adds EmitAMDGCNImageOverloadedReturnType for amdgcn_image_lo…
ranapratap55 Aug 12, 2025
f27e89a
[AMDGPU] Extends builtin support for amdgcn_image_sample and adds sem…
ranapratap55 Sep 9, 2025
46c833f
Using image desc as opaque ptr type for image load/store/sample
ranapratap55 Oct 7, 2025
44e69e1
update clang format
ranapratap55 Oct 7, 2025
fb0dc56
Updated getTextureDescIndex() to lambda and minor changes
ranapratap55 Oct 8, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
74 changes: 74 additions & 0 deletions clang/include/clang/Basic/BuiltinsAMDGPU.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -885,5 +885,79 @@ TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_store_16x8B, "vV2i*V2iIicC*",
TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_load_8x16B, "V4iV4i*IicC*", "nc", "gfx1250-insts,wavefrontsize32")
TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_store_8x16B, "vV4i*V4iIicC*", "nc", "gfx1250-insts,wavefrontsize32")

//===----------------------------------------------------------------------===//
// Image builtins
//===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_f32_i32, "fiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_f32_i32, "fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_f32_i32, "fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_f32_i32, "fiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f32_i32, "vV4fiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4hiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_f32_i32, "vfiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_f32_i32, "vfiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_f32_i32, "vfiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_f32_i32, "vfiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4hifQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_f32_f32, "fiffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_f32_f32, "fifffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")

#undef BUILTIN
#undef TARGET_BUILTIN
219 changes: 219 additions & 0 deletions clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,8 +11,11 @@
//===----------------------------------------------------------------------===//

#include "CGBuiltin.h"
#include "CodeGenFunction.h"
#include "clang/Basic/TargetBuiltins.h"
#include "clang/Frontend/FrontendDiagnostic.h"
#include "llvm/Analysis/ValueTracking.h"
#include "llvm/CodeGen/MachineFunction.h"
#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/IntrinsicsR600.h"
#include "llvm/IR/MemoryModelRelaxationAnnotations.h"
Expand Down Expand Up @@ -181,6 +184,92 @@ static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E,
return Call;
}

static bool IsImageSampleBuiltIn(unsigned BuiltinID) {
switch (BuiltinID) {
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
case clang::AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
return true;
default:
return false;
}
}

static llvm::Value *LoadTextureDescPtorAsVec8I32(CodeGenFunction &CGF,
llvm::Value *RsrcPtr) {
auto &B = CGF.Builder;
auto *VecTy = llvm::FixedVectorType::get(B.getInt32Ty(), 8);

if (RsrcPtr->getType() == VecTy)
return RsrcPtr;

if (RsrcPtr->getType()->isIntegerTy(32)) {
unsigned AS = 8;
llvm::PointerType *VecPtrTy = llvm::PointerType::get(VecTy, AS);
llvm::Value *Ptr = B.CreateIntToPtr(RsrcPtr, VecPtrTy, "tex.rsrc.from.int");
return B.CreateAlignedLoad(VecTy, Ptr, llvm::Align(32), "tex.rsrc.val");
}

if (RsrcPtr->getType()->isPointerTy()) {
unsigned AS = RsrcPtr->getType()->getPointerAddressSpace();
auto *VecPtrTy = llvm::PointerType::get(VecTy, AS);
llvm::Value *Typed = B.CreateBitCast(RsrcPtr, VecPtrTy, "tex.rsrc.typed");
return B.CreateAlignedLoad(VecTy, Typed, llvm::Align(32), "tex.rsrc.val");
}

const auto &DL = CGF.CGM.getDataLayout();
if (DL.getTypeSizeInBits(RsrcPtr->getType()) == 256)
return B.CreateBitCast(RsrcPtr, VecTy, "tex.rsrc.val");

RsrcPtr->getType()->print(llvm::errs());
llvm::report_fatal_error(": Unexpected texture resource argument form");
}

static unsigned GetTextureDescIndex(unsigned BuiltinID, const CallExpr *E) {
unsigned N = E->getNumArgs();
if (IsImageSampleBuiltIn(BuiltinID)) {
if (N < 5)
return (unsigned)-1;
return N - 5;
}

if (N < 3)
return (unsigned)-1;
return N - 3;
}

llvm::CallInst *
EmitAMDGCNImageOverloadedReturnType(clang::CodeGen::CodeGenFunction &CGF,
const clang::CallExpr *E,
unsigned IntrinsicID, bool IsImageStore) {
clang::SmallVector<llvm::Value *, 10> Args;
unsigned RsrcIndex = GetTextureDescIndex(E->getBuiltinCallee(), E);

for (unsigned I = 0; I < E->getNumArgs(); ++I) {
llvm::Value *V = CGF.EmitScalarExpr(E->getArg(I));
if (I == RsrcIndex)
V = LoadTextureDescPtorAsVec8I32(CGF, V);
Args.push_back(V);
}

llvm::Type *RetTy = CGF.ConvertType(E->getType());
if (IsImageStore)
RetTy = CGF.VoidTy;
llvm::CallInst *Call = CGF.Builder.CreateIntrinsic(RetTy, IntrinsicID, Args);
return Call;
}

// Emit an intrinsic that has 1 float or double operand, and 1 integer.
static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
const CallExpr *E,
Expand Down Expand Up @@ -937,6 +1026,136 @@ 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_1d_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_1d, false);
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_1darray, false);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_2d, false);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_2darray, false);
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_3d, false);
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_cube, false);
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_mip_1d, false);
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_mip_1darray, false);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_mip_2d, false);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_mip_2darray, false);
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_mip_3d, false);
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_load_mip_cube, false);
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_1d, true);
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_1darray, true);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_2d, true);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_2darray, true);
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_3d, true);
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_cube, true);
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_mip_1d, true);
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_mip_1darray, true);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_mip_2d, true);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_mip_2darray, true);
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_mip_3d, true);
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_store_mip_cube, true);
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_sample_1d, false);
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_sample_1darray, false);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_sample_2d, false);
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:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_sample_2darray, false);
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_sample_3d, false);
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
return EmitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_sample_cube, false);
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
llvm::FixedVectorType *VT = FixedVectorType::get(Builder.getInt32Ty(), 8);
Expand Down
Loading
Loading