6a961dc03d
\[NVPTX\] Add Prefetch intrinsics This PR adds prefetch intrinsics with the relevant eviction priorities. * Lit tests are added as part of prefetch.ll * The generated PTX is verified with a 12.3 ptxas executable. * Added docs for these intrinsics in NVPTXUsage.rst. For more information, refer PTX ISA `<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-prefetch-prefetchu>`_. --------- Co-authored-by: abmajumder <abmajumder@nvidia.com>
81 lines
3.5 KiB
LLVM
81 lines
3.5 KiB
LLVM
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
|
|
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK-PTX64 %s
|
|
; RUN: %if ptxas-12.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}
|
|
|
|
target triple = "nvptx64-nvidia-cuda"
|
|
|
|
declare void @llvm.nvvm.prefetch.local.L1.evictnormal(ptr addrspace(5) %local_ptr)
|
|
declare void @llvm.nvvm.prefetch.local.L2.evictnormal(ptr addrspace(5) %local_ptr)
|
|
|
|
declare void @llvm.nvvm.prefetch.global.L1.evictnormal(ptr addrspace(1) %global_ptr)
|
|
declare void @llvm.nvvm.prefetch.global.L2.evictnormal(ptr addrspace(1) %global_ptr)
|
|
declare void @llvm.nvvm.prefetch.global.L1.evictlast(ptr addrspace(1) %global_ptr)
|
|
declare void @llvm.nvvm.prefetch.global.L2.evictlast(ptr addrspace(1) %global_ptr)
|
|
|
|
declare void @llvm.nvvm.prefetch.L1.evictnormal(ptr %ptr)
|
|
declare void @llvm.nvvm.prefetch.L2.evictnormal(ptr %ptr)
|
|
|
|
declare void @llvm.nvvm.prefetchu.L1.evictnormal(ptr %ptr)
|
|
|
|
define void @prefetch_local(ptr addrspace(5) %local_ptr) {
|
|
; CHECK-PTX64-LABEL: prefetch_local(
|
|
; CHECK-PTX64: {
|
|
; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
|
|
; CHECK-PTX64-EMPTY:
|
|
; CHECK-PTX64-NEXT: // %bb.0:
|
|
; CHECK-PTX64-NEXT: ld.param.u64 %rd1, [prefetch_local_param_0];
|
|
; CHECK-PTX64-NEXT: prefetch.local.L1.evictnormal [%rd1];
|
|
; CHECK-PTX64-NEXT: prefetch.local.L2.evictnormal [%rd1];
|
|
; CHECK-PTX64-NEXT: ret;
|
|
tail call void @llvm.nvvm.prefetch.local.L1.evictnormal(ptr addrspace(5) %local_ptr)
|
|
tail call void @llvm.nvvm.prefetch.local.L2.evictnormal(ptr addrspace(5) %local_ptr)
|
|
ret void
|
|
}
|
|
|
|
define void @prefetch_global(ptr addrspace(1) %global_ptr) {
|
|
; CHECK-PTX64-LABEL: prefetch_global(
|
|
; CHECK-PTX64: {
|
|
; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
|
|
; CHECK-PTX64-EMPTY:
|
|
; CHECK-PTX64-NEXT: // %bb.0:
|
|
; CHECK-PTX64-NEXT: ld.param.u64 %rd1, [prefetch_global_param_0];
|
|
; CHECK-PTX64-NEXT: prefetch.global.L1.evictnormal [%rd1];
|
|
; CHECK-PTX64-NEXT: prefetch.global.L2.evictnormal [%rd1];
|
|
; CHECK-PTX64-NEXT: prefetch.global.L1.evictlast [%rd1];
|
|
; CHECK-PTX64-NEXT: prefetch.global.L2.evictlast [%rd1];
|
|
; CHECK-PTX64-NEXT: ret;
|
|
tail call void @llvm.nvvm.prefetch.global.L1.evictnormal(ptr addrspace(1) %global_ptr)
|
|
tail call void @llvm.nvvm.prefetch.global.L2.evictnormal(ptr addrspace(1) %global_ptr)
|
|
tail call void @llvm.nvvm.prefetch.global.L1.evictlast(ptr addrspace(1) %global_ptr)
|
|
tail call void @llvm.nvvm.prefetch.global.L2.evictlast(ptr addrspace(1) %global_ptr)
|
|
ret void
|
|
}
|
|
|
|
|
|
define void @prefetch_(ptr %ptr) {
|
|
; CHECK-PTX64-LABEL: prefetch_(
|
|
; CHECK-PTX64: {
|
|
; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
|
|
; CHECK-PTX64-EMPTY:
|
|
; CHECK-PTX64-NEXT: // %bb.0:
|
|
; CHECK-PTX64-NEXT: ld.param.u64 %rd1, [prefetch__param_0];
|
|
; CHECK-PTX64-NEXT: prefetch.L1.evictnormal [%rd1];
|
|
; CHECK-PTX64-NEXT: prefetch.L2.evictnormal [%rd1];
|
|
; CHECK-PTX64-NEXT: ret;
|
|
tail call void @llvm.nvvm.prefetch.L1.evictnormal(ptr %ptr)
|
|
tail call void @llvm.nvvm.prefetch.L2.evictnormal(ptr %ptr)
|
|
ret void
|
|
}
|
|
|
|
define void @prefetchu_l1(ptr %ptr) {
|
|
; CHECK-PTX64-LABEL: prefetchu_l1(
|
|
; CHECK-PTX64: {
|
|
; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
|
|
; CHECK-PTX64-EMPTY:
|
|
; CHECK-PTX64-NEXT: // %bb.0:
|
|
; CHECK-PTX64-NEXT: ld.param.u64 %rd1, [prefetchu_l1_param_0];
|
|
; CHECK-PTX64-NEXT: prefetchu.L1.evictnormal [%rd1];
|
|
; CHECK-PTX64-NEXT: ret;
|
|
tail call void @llvm.nvvm.prefetchu.L1.evictnormal(ptr %ptr)
|
|
ret void
|
|
} |