36 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74| FileCheck --check-prefixes=CHECK-PTX64 %s
3; RUN: %if ptxas-sm_80 && ptxas-isa-7.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74| %ptxas-verify -arch=sm_80 %}
4
5target triple = "nvptx64-nvidia-cuda"
6
7declare void @llvm.nvvm.discard.global.L2(ptr addrspace(1) %global_ptr, i64 immarg %size)
8declare void @llvm.nvvm.discard.L2(ptr %ptr, i64 immarg %size)
9
10define void @discard_global_L2(ptr addrspace(1) %global_ptr) {
11; CHECK-PTX64-LABEL: discard_global_L2(
12; CHECK-PTX64: {
13; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
14; CHECK-PTX64-EMPTY:
15; CHECK-PTX64-NEXT: // %bb.0:
16; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [discard_global_L2_param_0];
17; CHECK-PTX64-NEXT: discard.global.L2 [%rd1], 128;
18; CHECK-PTX64-NEXT: ret;
19 tail call void @llvm.nvvm.discard.global.L2(ptr addrspace(1) %global_ptr, i64 128)
20 ret void
21}
22
23define void @discard_L2(ptr %ptr) {
24; CHECK-PTX64-LABEL: discard_L2(
25; CHECK-PTX64: {
26; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
27; CHECK-PTX64-EMPTY:
28; CHECK-PTX64-NEXT: // %bb.0:
29; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [discard_L2_param_0];
30; CHECK-PTX64-NEXT: discard.L2 [%rd1], 128;
31; CHECK-PTX64-NEXT: ret;
32 tail call void @llvm.nvvm.discard.L2(ptr %ptr, i64 128)
33 ret void
34}
35
36