llvm/test/CodeGen/NVPTX/lower-args.ll

   1 ; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-cuda | FileCheck %s --check-prefixes COMMON,IR,IRC
   2 ; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-nvcl | FileCheck %s --check-prefixes COMMON,IR,IRO
   3 ; RUN: llc < %s -mcpu=sm_20 --mtriple nvptx64-nvidia-cuda | FileCheck %s --check-prefixes COMMON,PTX,PTXC
   4 ; RUN: llc < %s -mcpu=sm_20 --mtriple nvptx64-nvidia-nvcl| FileCheck %s --check-prefixes COMMON,PTX,PTXO
   5 ; RUN: %if ptxas %{ llc < %s -mcpu=sm_20 | %ptxas-verify %}
   6
   7 target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
   8 target triple = "nvptx64-nvidia-cuda"
   9
  10 %class.outer = type <{ %class.inner, i32, [4 x i8] }>
  11 %class.inner = type { ptr, ptr }
  12
  13 ; Check that nvptx-lower-args preserves arg alignment
  14 ; COMMON-LABEL: load_alignment
  15 define void @load_alignment(ptr nocapture readonly byval(%class.outer) align 8 %arg) {
  16 entry:
  17 ; IR: load %class.outer, ptr addrspace(101)
  18 ; IR-SAME: align 8
  19 ; PTX: ld.param.u64
  20 ; PTX-NOT: ld.param.u8
  21   %arg.idx.val = load ptr, ptr %arg, align 8
  22   %arg.idx1 = getelementptr %class.outer, ptr %arg, i64 0, i32 0, i32 1
  23   %arg.idx1.val = load ptr, ptr %arg.idx1, align 8
  24   %arg.idx2 = getelementptr %class.outer, ptr %arg, i64 0, i32 1
  25   %arg.idx2.val = load i32, ptr %arg.idx2, align 8
  26   %arg.idx.val.val = load i32, ptr %arg.idx.val, align 4
  27   %add.i = add nsw i32 %arg.idx.val.val, %arg.idx2.val
  28   store i32 %add.i, ptr %arg.idx1.val, align 4
  29
  30   ; let the pointer escape so we still create a local copy this test uses to
  31   ; check the load alignment.
  32   %tmp = call ptr @escape(ptr nonnull %arg.idx2)
  33   ret void
  34 }
  35
  36
  37 ; COMMON-LABEL: ptr_generic
  38 define void @ptr_generic(ptr %out, ptr %in) {
  39 ; IRC:  %in3 = addrspacecast ptr %in to ptr addrspace(1)
  40 ; IRC:  %in4 = addrspacecast ptr addrspace(1) %in3 to ptr
  41 ; IRC:  %out1 = addrspacecast ptr %out to ptr addrspace(1)
  42 ; IRC:  %out2 = addrspacecast ptr addrspace(1) %out1 to ptr
  43 ; PTXC: cvta.to.global.u64
  44 ; PTXC: cvta.to.global.u64
  45 ; PTXC: ld.global.u32
  46 ; PTXC: st.global.u32
  47
  48 ; OpenCL can't make assumptions about incoming pointer, so we should generate
  49 ; generic pointers load/store.
  50 ; IRO-NOT: addrspacecast
  51 ; PTXO-NOT: cvta.to.global
  52 ; PTXO: ld.u32
  53 ; PTXO: st.u32
  54   %v = load i32, ptr  %in, align 4
  55   store i32 %v, ptr %out, align 4
  56   ret void
  57 }
  58
  59 ; COMMON-LABEL: ptr_nongeneric
  60 define void @ptr_nongeneric(ptr addrspace(1) %out, ptr addrspace(4) %in) {
  61 ; IR-NOT: addrspacecast
  62 ; PTX-NOT: cvta.to.global
  63 ; PTX:  ld.const.u32
  64 ; PTX   st.global.u32
  65   %v = load i32, ptr addrspace(4) %in, align 4
  66   store i32 %v, ptr addrspace(1) %out, align 4
  67   ret void
  68 }
  69
  70 ; COMMON-LABEL: ptr_as_int
  71  define void @ptr_as_int(i64 noundef %i, i32 noundef %v) {
  72 ; IR:   [[P:%.*]] = inttoptr i64 %i to ptr
  73 ; IRC:  [[P1:%.*]] = addrspacecast ptr [[P]] to ptr addrspace(1)
  74 ; IRC:  addrspacecast ptr addrspace(1) [[P1]] to ptr
  75 ; IRO-NOT: addrspacecast
  76
  77 ; PTXC-DAG:  ld.param.u64    [[I:%rd.*]], [ptr_as_int_param_0];
  78 ; PTXC-DAG:  ld.param.u32    [[V:%r.*]], [ptr_as_int_param_1];
  79 ; PTXC:      cvta.to.global.u64 %[[P:rd.*]], [[I]];
  80 ; PTXC:      st.global.u32    [%[[P]]], [[V]];
  81
  82 ; PTXO-DAG:  ld.param.u64    %[[P:rd.*]], [ptr_as_int_param_0];
  83 ; PTXO-DAG:  ld.param.u32    [[V:%r.*]], [ptr_as_int_param_1];
  84 ; PTXO:      st.u32   [%[[P]]], [[V]];
  85
  86   %p = inttoptr i64 %i to ptr
  87   store i32 %v, ptr %p, align 4
  88   ret void
  89 }
  90
  91 %struct.S = type { i64 }
  92
  93 ; COMMON-LABEL: ptr_as_int_aggr
  94 define void @ptr_as_int_aggr(ptr nocapture noundef readonly byval(%struct.S) align 8 %s, i32 noundef %v) {
  95 ; IR:   [[S:%.*]] = addrspacecast ptr %s to ptr addrspace(101)
  96 ; IR:   [[I:%.*]] = load i64, ptr addrspace(101) [[S]], align 8
  97 ; IR:   [[P0:%.*]] = inttoptr i64 [[I]] to ptr
  98 ; IRC:  [[P1:%.*]] = addrspacecast ptr [[P]] to ptr addrspace(1)
  99 ; IRC:  [[P:%.*]] = addrspacecast ptr addrspace(1) [[P1]] to ptr
 100 ; IRO-NOT: addrspacecast
 101
 102 ; PTXC-DAG:  ld.param.u64    [[I:%rd.*]], [ptr_as_int_aggr_param_0];
 103 ; PTXC-DAG:  ld.param.u32    [[V:%r.*]], [ptr_as_int_aggr_param_1];
 104 ; PTXC:      cvta.to.global.u64 %[[P:rd.*]], [[I]];
 105 ; PTXC:      st.global.u32    [%[[P]]], [[V]];
 106
 107 ; PTXO-DAG:  ld.param.u64    %[[P:rd.*]], [ptr_as_int_aggr_param_0];
 108 ; PTXO-DAG:  ld.param.u32    [[V:%r.*]], [ptr_as_int_aggr_param_1];
 109 ; PTXO:      st.u32   [%[[P]]], [[V]];
 110   %i = load i64, ptr %s, align 8
 111   %p = inttoptr i64 %i to ptr
 112   store i32 %v, ptr %p, align 4
 113   ret void
 114 }
 115
 116
 117 ; Function Attrs: convergent nounwind
 118 declare dso_local ptr @escape(ptr) local_unnamed_addr
 119 !nvvm.annotations = !{!0, !1, !2, !3}
 120 !0 = !{ptr @ptr_generic, !"kernel", i32 1}
 121 !1 = !{ptr @ptr_nongeneric, !"kernel", i32 1}
 122 !2 = !{ptr @ptr_as_int, !"kernel", i32 1}
 123 !3 = !{ptr @ptr_as_int_aggr, !"kernel", i32 1}