27 lines · plain
1; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s2; 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 18 9; CHECK: .global .align 4 .u32 myglobal = 42;10@myglobal = internal global i32 42, align 411; CHECK: .global .align 4 .u32 myconst = 420;12@myconst = internal constant i32 420, align 413 14 15define ptx_kernel void @foo(ptr %a, ptr %b) {16; Expect one load -- @myconst isn't loaded from, because we know its value17; statically.18; CHECK: ld.global.b3219; CHECK: st.global.b3220; CHECK: st.global.b3221 %ld1 = load i32, ptr @myglobal22 %ld2 = load i32, ptr @myconst23 store i32 %ld1, ptr %a24 store i32 %ld2, ptr %b25 ret void26}27