xref: /llvm-project/llvm/test/CodeGen/NVPTX/ldu-ldg.ll (revision b279f6b098d3849f7f1c1f539b108307d5f8ae2d) 
1; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_32 | FileCheck %s
2; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_32 | %ptxas-verify %}
3
4
5declare i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 %align)
6declare i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 %align)
7declare i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 %align)
8declare i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 %align)
9declare ptr @llvm.nvvm.ldu.global.p.p1(ptr addrspace(1) %ptr, i32 %align)
10declare float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 %align)
11declare double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)
12declare half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)
13declare <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)
14
15declare i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 %align)
16declare i16 @llvm.nvvm.ldg.global.i.i16.p1(ptr addrspace(1) %ptr, i32 %align)
17declare i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %ptr, i32 %align)
18declare i64 @llvm.nvvm.ldg.global.i.i64.p1(ptr addrspace(1) %ptr, i32 %align)
19declare ptr @llvm.nvvm.ldg.global.p.p1(ptr addrspace(1) %ptr, i32 %align)
20declare float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %ptr, i32 %align)
21declare double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)
22declare half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)
23declare <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)
24
25; CHECK-LABEL: test_ldu_i8
26define i8 @test_ldu_i8(ptr addrspace(1) %ptr) {
27  ; CHECK: ldu.global.u8
28  %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
29  ret i8 %val
30}
31
32; CHECK-LABEL: test_ldu_i16
33define i16 @test_ldu_i16(ptr addrspace(1) %ptr) {
34  ; CHECK: ldu.global.u16
35  %val = tail call i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)
36  ret i16 %val
37}
38
39; CHECK-LABEL: test_ldu_i32
40define i32 @test_ldu_i32(ptr addrspace(1) %ptr) {
41  ; CHECK: ldu.global.u32
42  %val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)
43  ret i32 %val
44}
45
46; CHECK-LABEL: test_ldu_i64
47define i64 @test_ldu_i64(ptr addrspace(1) %ptr) {
48  ; CHECK: ldu.global.u64
49  %val = tail call i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)
50  ret i64 %val
51}
52
53; CHECK-LABEL: test_ldu_p
54define ptr @test_ldu_p(ptr addrspace(1) %ptr) {
55  ; CHECK: ldu.global.u64
56  %val = tail call ptr @llvm.nvvm.ldu.global.p.p1(ptr addrspace(1) %ptr, i32 8)
57  ret ptr %val
58}
59
60
61; CHECK-LABEL: test_ldu_f32
62define float @test_ldu_f32(ptr addrspace(1) %ptr) {
63  ; CHECK: ldu.global.f32
64  %val = tail call float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)
65  ret float %val
66}
67
68; CHECK-LABEL: test_ldu_f64
69define double @test_ldu_f64(ptr addrspace(1) %ptr) {
70  ; CHECK: ldu.global.f64
71  %val = tail call double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)
72  ret double %val
73}
74
75; CHECK-LABEL: test_ldu_f16
76define half @test_ldu_f16(ptr addrspace(1) %ptr) {
77  ; CHECK: ldu.global.u16
78  %val = tail call half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)
79  ret half %val
80}
81
82; CHECK-LABEL: test_ldu_v2f16
83define <2 x half> @test_ldu_v2f16(ptr addrspace(1) %ptr) {
84  ; CHECK: ldu.global.u32
85  %val = tail call <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
86  ret <2 x half> %val
87}
88
89; CHECK-LABEL: test_ldg_i8
90define i8 @test_ldg_i8(ptr addrspace(1) %ptr) {
91  ; CHECK: ld.global.nc.u8
92  %val = tail call i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
93  ret i8 %val
94}
95
96; CHECK-LABEL: test_ldg_i16
97define i16 @test_ldg_i16(ptr addrspace(1) %ptr) {
98  ; CHECK: ld.global.nc.u16
99  %val = tail call i16 @llvm.nvvm.ldg.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)
100  ret i16 %val
101}
102
103; CHECK-LABEL: test_ldg_i32
104define i32 @test_ldg_i32(ptr addrspace(1) %ptr) {
105  ; CHECK: ld.global.nc.u32
106  %val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)
107  ret i32 %val
108}
109
110; CHECK-LABEL: test_ldg_i64
111define i64 @test_ldg_i64(ptr addrspace(1) %ptr) {
112  ; CHECK: ld.global.nc.u64
113  %val = tail call i64 @llvm.nvvm.ldg.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)
114  ret i64 %val
115}
116
117; CHECK-LABEL: test_ldg_p
118define ptr @test_ldg_p(ptr addrspace(1) %ptr) {
119  ; CHECK: ld.global.nc.u64
120  %val = tail call ptr @llvm.nvvm.ldg.global.p.p1(ptr addrspace(1) %ptr, i32 8)
121  ret ptr %val
122}
123
124; CHECK-LABEL: test_ldg_f32
125define float @test_ldg_f32(ptr addrspace(1) %ptr) {
126  ; CHECK: ld.global.nc.f32
127  %val = tail call float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)
128  ret float %val
129}
130
131; CHECK-LABEL: test_ldg_f64
132define double @test_ldg_f64(ptr addrspace(1) %ptr) {
133  ; CHECK: ld.global.nc.f64
134  %val = tail call double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)
135  ret double %val
136}
137
138; CHECK-LABEL: test_ldg_f16
139define half @test_ldg_f16(ptr addrspace(1) %ptr) {
140  ; CHECK: ld.global.nc.u16
141  %val = tail call half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)
142  ret half %val
143}
144
145; CHECK-LABEL: test_ldg_v2f16
146define <2 x half> @test_ldg_v2f16(ptr addrspace(1) %ptr) {
147  ; CHECK: ld.global.nc.u32
148  %val = tail call <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
149  ret <2 x half> %val
150}
151