1; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s 2; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %} 3 4target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64" 5target triple = "nvptx-nvidia-cuda" 6 7; Ensure global variables in address space 0 are promoted to address space 1 8 9; CHECK: .global .align 4 .u32 myglobal = 42; 10@myglobal = internal global i32 42, align 4 11; CHECK: .global .align 4 .u32 myconst = 420; 12@myconst = internal constant i32 420, align 4 13 14 15define ptx_kernel void @foo(ptr %a, ptr %b) { 16; Expect one load -- @myconst isn't loaded from, because we know its value 17; statically. 18; CHECK: ld.global.u32 19; CHECK: st.global.u32 20; CHECK: st.global.u32 21 %ld1 = load i32, ptr @myglobal 22 %ld2 = load i32, ptr @myconst 23 store i32 %ld1, ptr %a 24 store i32 %ld2, ptr %b 25 ret void 26} 27