intel/llvm
1
0
Fork 0
mirror of https://github.com/intel/llvm.git synced 2026-01-25 10:55:58 +08:00
Code Issues Packages Projects Releases Wiki Activity

[NVPTX] Add builtin support for 'nanosleep' PTX instrunction (#79888)

Browse Source 
Summary:
This patch adds a builtin for the `nanosleep` PTX function. It takes
either an immediate or a register and sleeps for [0, 2t] nanoseconds
given t. More information at the documentation:

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-nanosleep
This commit is contained in:
Joseph Huber
2024-01-29 14:07:58 -06:00
committed by GitHub
parent d492faa7aa
commit 5f12cc912a
5 changed files with 42 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
clang/include/clang/Basic/BuiltinsNVPTX.def
View File

@@ -157,6 +157,7 @@ BUILTIN(__nvvm_read_ptx_sreg_pm3, "i", "n")
// MISC
BUILTIN(__nvvm_prmt, "UiUiUiUi", "")
TARGET_BUILTIN(__nvvm_nanosleep, "vi", "n", AND(SM_70, PTX63))
// Min Max

11
clang/test/CodeGen/builtins-nvptx.c
View File

@@ -818,6 +818,17 @@ __device__ void nvvm_vote(int pred) {
// CHECK: ret void
}
// CHECK-LABEL: nvvm_nanosleep
__device__ void nvvm_nanosleep(int d) {
#if __CUDA_ARCH__ >= 700
// CHECK_PTX70_SM80: call void @llvm.nvvm.nanosleep
__nvvm_nanosleep(d);
// CHECK_PTX70_SM80: call void @llvm.nvvm.nanosleep
__nvvm_nanosleep(1);
#endif
}
// CHECK-LABEL: nvvm_mbarrier
__device__ void nvvm_mbarrier(long long* addr, __attribute__((address_space(3))) long long* sharedAddr, int count, long long state) {
#if __CUDA_ARCH__ >= 800

4
llvm/include/llvm/IR/IntrinsicsNVVM.td
View File

@@ -557,6 +557,10 @@ let TargetPrefix = "nvvm" in {
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]>;
def int_nvvm_nanosleep : ClangBuiltin<"__nvvm_nanosleep">,
DefaultAttrsIntrinsic<[], [llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrHasSideEffects]>;
//
// Min Max
//

6
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
View File

@@ -640,6 +640,12 @@ class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass,
def INT_NVVM_PRMT : F_MATH_3<"prmt.b32 \t$dst, $src0, $src1, $src2;", Int32Regs,
Int32Regs, Int32Regs, Int32Regs, int_nvvm_prmt>;
def INT_NVVM_NANOSLEEP_I : NVPTXInst<(outs), (ins i32imm:$i), "nanosleep.u32 \t$i;",
[(int_nvvm_nanosleep imm:$i)]>,
Requires<[hasPTX<63>, hasSM<70>]>;
def INT_NVVM_NANOSLEEP_R : NVPTXInst<(outs), (ins Int32Regs:$i), "nanosleep.u32 \t$i;",
[(int_nvvm_nanosleep Int32Regs:$i)]>,
Requires<[hasPTX<63>, hasSM<70>]>;
//
// Min Max
//

20
llvm/test/CodeGen/NVPTX/nanosleep.ll Normal file
View File

@@ -0,0 +1,20 @@
; RUN: llc < %s -march=nvptx64 -O2 -mcpu=sm_70 -mattr=+ptx63 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | %ptxas-verify %}
declare void @llvm.nvvm.nanosleep(i32)
; CHECK-LABEL: test_nanosleep_r
define void @test_nanosleep_r(i32 noundef %d) {
entry:
; CHECK: nanosleep.u32 %[[REG:.+]];
call void @llvm.nvvm.nanosleep(i32 %d)
ret void
}
; CHECK-LABEL: test_nanosleep_i
define void @test_nanosleep_i() {
entry:
; CHECK: nanosleep.u32 42;
call void @llvm.nvvm.nanosleep(i32 42)
ret void
}