[clang] [Clang][AMDGPU] Add builtins for instrinsic `llvm.amdgcn.raw.ptr.buffer.load` (PR #99258)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Thu Jul 18 01:50:21 PDT 2024
================
@@ -19177,6 +19177,42 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128:
return emitBuiltinWithOneOverloadedType<5>(
*this, E, Intrinsic::amdgcn_raw_ptr_buffer_store);
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: {
+ llvm::Type *RetTy = nullptr;
+ switch (BuiltinID) {
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:
+ RetTy = Int8Ty;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:
+ RetTy = Int16Ty;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:
+ RetTy = Int32Ty;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:
+ RetTy = llvm::FixedVectorType::get(Int32Ty, /*NumElements=*/2);
+ break;
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:
+ RetTy = llvm::FixedVectorType::get(Int32Ty, /*NumElements=*/3);
+ break;
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128:
+ RetTy = llvm::FixedVectorType::get(Int32Ty, /*NumElements=*/4);
+ break;
+ }
+ SmallVector<Value *, 4> Args;
+ Args.push_back(EmitScalarExpr(E->getArg(0)));
+ Args.push_back(EmitScalarExpr(E->getArg(1)));
+ Args.push_back(EmitScalarExpr(E->getArg(2)));
+ Args.push_back(EmitScalarExpr(E->getArg(3)));
----------------
arsenm wrote:
Can do this in initializer list
https://github.com/llvm/llvm-project/pull/99258
More information about the cfe-commits
mailing list