| Justin Holewinski | 9a2350e | 2014-07-17 11:59:04 +0000 | [diff] [blame] | 1 | ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=SM20 |
| 2 | ; RUN: llc < %s -march=nvptx -mcpu=sm_30 | FileCheck %s --check-prefix=SM30 |
| 3 | |
| 4 | |
| 5 | target triple = "nvptx-unknown-cuda" |
| 6 | |
| 7 | declare { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64, i32) |
| 8 | declare i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)*) |
| 9 | |
| 10 | ; SM20-LABEL: .entry foo |
| 11 | ; SM30-LABEL: .entry foo |
| 12 | define void @foo(i64 %img, float* %red, i32 %idx) { |
| 13 | ; SM20: ld.param.u64 %rd[[TEXREG:[0-9]+]], [foo_param_0]; |
| 14 | ; SM20: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXREG]], {%r{{[0-9]+}}}] |
| 15 | ; SM30: ld.param.u64 %rd[[TEXREG:[0-9]+]], [foo_param_0]; |
| 16 | ; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXREG]], {%r{{[0-9]+}}}] |
| 17 | %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx) |
| 18 | %ret = extractvalue { float, float, float, float } %val, 0 |
| Jingyue Wu | a2f6027 | 2015-06-04 21:28:26 +0000 | [diff] [blame^] | 19 | ; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] |
| 20 | ; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] |
| Justin Holewinski | 9a2350e | 2014-07-17 11:59:04 +0000 | [diff] [blame] | 21 | store float %ret, float* %red |
| 22 | ret void |
| 23 | } |
| 24 | |
| 25 | |
| 26 | @tex0 = internal addrspace(1) global i64 0, align 8 |
| 27 | |
| 28 | ; SM20-LABEL: .entry bar |
| 29 | ; SM30-LABEL: .entry bar |
| 30 | define void @bar(float* %red, i32 %idx) { |
| 31 | ; SM30: mov.u64 %rd[[TEXHANDLE:[0-9]+]], tex0 |
| 32 | %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @tex0) |
| 33 | ; SM20: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [tex0, {%r{{[0-9]+}}}] |
| 34 | ; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXHANDLE]], {%r{{[0-9]+}}}] |
| 35 | %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx) |
| 36 | %ret = extractvalue { float, float, float, float } %val, 0 |
| Jingyue Wu | a2f6027 | 2015-06-04 21:28:26 +0000 | [diff] [blame^] | 37 | ; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] |
| 38 | ; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] |
| Justin Holewinski | 9a2350e | 2014-07-17 11:59:04 +0000 | [diff] [blame] | 39 | store float %ret, float* %red |
| 40 | ret void |
| 41 | } |
| 42 | |
| 43 | !nvvm.annotations = !{!1, !2, !3} |
| Duncan P. N. Exon Smith | be7ea19 | 2014-12-15 19:07:53 +0000 | [diff] [blame] | 44 | !1 = !{void (i64, float*, i32)* @foo, !"kernel", i32 1} |
| 45 | !2 = !{void (float*, i32)* @bar, !"kernel", i32 1} |
| 46 | !3 = !{i64 addrspace(1)* @tex0, !"texture", i32 1} |