diff options
| author | Changpeng Fang <changpeng.fang@amd.com> | 2025-07-24 16:23:33 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2025-07-24 16:23:33 -0700 |
| commit | d7a38a94cd89e831fb17c5d45d11d0539e1e898d (patch) | |
| tree | a196f5d6071f8a0456bead16cacc9cf15bb52c8a /clang/lib/CodeGen | |
| parent | af98a245f8604ed5db9ef08bed96f0f545e8b46b (diff) | |
| download | llvm-d7a38a94cd89e831fb17c5d45d11d0539e1e898d.zip llvm-d7a38a94cd89e831fb17c5d45d11d0539e1e898d.tar.gz llvm-d7a38a94cd89e831fb17c5d45d11d0539e1e898d.tar.bz2 | |
[AMDGPU] Support builtin/intrinsics for load monitors on gfx1250 (#150540)
Diffstat (limited to 'clang/lib/CodeGen')
| -rw-r--r-- | clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 35 |
1 files changed, 35 insertions, 0 deletions
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index e469a01..70f510a 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -633,6 +633,41 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy}); return Builder.CreateCall(F, {Addr}); } + 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: { + + Intrinsic::ID IID; + switch (BuiltinID) { + case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32: + IID = Intrinsic::amdgcn_global_load_monitor_b32; + break; + case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64: + IID = Intrinsic::amdgcn_global_load_monitor_b64; + break; + case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128: + IID = Intrinsic::amdgcn_global_load_monitor_b128; + break; + case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32: + IID = Intrinsic::amdgcn_flat_load_monitor_b32; + break; + case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64: + IID = Intrinsic::amdgcn_flat_load_monitor_b64; + break; + case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: + IID = Intrinsic::amdgcn_flat_load_monitor_b128; + break; + } + + llvm::Type *LoadTy = ConvertType(E->getType()); + llvm::Value *Addr = EmitScalarExpr(E->getArg(0)); + llvm::Value *Val = EmitScalarExpr(E->getArg(1)); + llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy}); + return Builder.CreateCall(F, {Addr, Val}); + } case AMDGPU::BI__builtin_amdgcn_load_to_lds: { // Should this have asan instrumentation? return emitBuiltinWithOneOverloadedType<5>(*this, E, |