| Justin Lebar | 6d6b11a | 2016-09-11 01:39:04 +0000 | [diff] [blame] | 1 | ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s | 
|  | 2 |  | 
|  | 3 | ; Check that invariant loads from the global addrspace are lowered to | 
|  | 4 | ; ld.global.nc. | 
|  | 5 |  | 
|  | 6 | ; CHECK-LABEL: @ld_global | 
|  | 7 | define i32 @ld_global(i32 addrspace(1)* %ptr) { | 
|  | 8 | ; CHECK: ld.global.nc.{{[a-z]}}32 | 
|  | 9 | %a = load i32, i32 addrspace(1)* %ptr, !invariant.load !0 | 
|  | 10 | ret i32 %a | 
|  | 11 | } | 
|  | 12 |  | 
|  | 13 | ; CHECK-LABEL: @ld_not_invariant | 
|  | 14 | define i32 @ld_not_invariant(i32 addrspace(1)* %ptr) { | 
|  | 15 | ; CHECK: ld.global.{{[a-z]}}32 | 
|  | 16 | %a = load i32, i32 addrspace(1)* %ptr | 
|  | 17 | ret i32 %a | 
|  | 18 | } | 
|  | 19 |  | 
|  | 20 | ; CHECK-LABEL: @ld_not_global_addrspace | 
|  | 21 | define i32 @ld_not_global_addrspace(i32 addrspace(0)* %ptr) { | 
|  | 22 | ; CHECK: ld.{{[a-z]}}32 | 
|  | 23 | %a = load i32, i32 addrspace(0)* %ptr | 
|  | 24 | ret i32 %a | 
|  | 25 | } | 
|  | 26 |  | 
|  | 27 | !0 = !{} |