test/CodeGen/NVPTX/ld-addrspace.ll

   1 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS32
   2 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS64
   3 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS32
   4
   5
   6 ;; i8
   7 define i8 @ld_global_i8(i8 addrspace(1)* %ptr) {
   8 ; ALL-LABEL: ld_global_i8
   9 ; G32: ld.global.u8 %{{.*}}, [%r{{[0-9]+}}]
  10 ; G64: ld.global.u8 %{{.*}}, [%rd{{[0-9]+}}]
  11 ; ALL: ret
  12   %a = load i8, i8 addrspace(1)* %ptr
  13   ret i8 %a
  14 }
  15 define i8 @ld_shared_i8(i8 addrspace(3)* %ptr) {
  16 ; ALL-LABEL: ld_shared_i8
  17 ; LS32: ld.shared.u8 %{{.*}}, [%r{{[0-9]+}}]
  18 ; LS64: ld.shared.u8 %{{.*}}, [%rd{{[0-9]+}}]
  19 ; ALL: ret
  20   %a = load i8, i8 addrspace(3)* %ptr
  21   ret i8 %a
  22 }
  23 define i8 @ld_local_i8(i8 addrspace(5)* %ptr) {
  24 ; ALL-LABEL: ld_local_i8
  25 ; LS32: ld.local.u8 %{{.*}}, [%r{{[0-9]+}}]
  26 ; LS64: ld.local.u8 %{{.*}}, [%rd{{[0-9]+}}]
  27 ; ALL: ret
  28   %a = load i8, i8 addrspace(5)* %ptr
  29   ret i8 %a
  30 }
  31
  32 ;; i16
  33 define i16 @ld_global_i16(i16 addrspace(1)* %ptr) {
  34 ; ALL-LABEL: ld_global_i16
  35 ; G32: ld.global.u16 %{{.*}}, [%r{{[0-9]+}}]
  36 ; G64: ld.global.u16 %{{.*}}, [%rd{{[0-9]+}}]
  37 ; ALL: ret
  38   %a = load i16, i16 addrspace(1)* %ptr
  39   ret i16 %a
  40 }
  41 define i16 @ld_shared_i16(i16 addrspace(3)* %ptr) {
  42 ; ALL-LABEL: ld_shared_i16
  43 ; LS32: ld.shared.u16 %{{.*}}, [%r{{[0-9]+}}]
  44 ; LS64: ld.shared.u16 %{{.*}}, [%rd{{[0-9]+}}]
  45 ; ALL: ret
  46   %a = load i16, i16 addrspace(3)* %ptr
  47   ret i16 %a
  48 }
  49 define i16 @ld_local_i16(i16 addrspace(5)* %ptr) {
  50 ; ALL-LABEL: ld_local_i16
  51 ; LS32: ld.local.u16 %{{.*}}, [%r{{[0-9]+}}]
  52 ; LS64: ld.local.u16 %{{.*}}, [%rd{{[0-9]+}}]
  53 ; ALL: ret
  54   %a = load i16, i16 addrspace(5)* %ptr
  55   ret i16 %a
  56 }
  57
  58 ;; i32
  59 define i32 @ld_global_i32(i32 addrspace(1)* %ptr) {
  60 ; ALL-LABEL: ld_global_i32
  61 ; G32: ld.global.u32 %{{.*}}, [%r{{[0-9]+}}]
  62 ; G64: ld.global.u32 %{{.*}}, [%rd{{[0-9]+}}]
  63 ; ALL: ret
  64   %a = load i32, i32 addrspace(1)* %ptr
  65   ret i32 %a
  66 }
  67 define i32 @ld_shared_i32(i32 addrspace(3)* %ptr) {
  68 ; ALL-LABEL: ld_shared_i32
  69 ; LS32: ld.shared.u32 %{{.*}}, [%r{{[0-9]+}}]
  70 ; LS64: ld.shared.u32 %{{.*}}, [%rd{{[0-9]+}}]
  71 ; PTX64: ret
  72   %a = load i32, i32 addrspace(3)* %ptr
  73   ret i32 %a
  74 }
  75 define i32 @ld_local_i32(i32 addrspace(5)* %ptr) {
  76 ; ALL-LABEL: ld_local_i32
  77 ; LS32: ld.local.u32 %{{.*}}, [%r{{[0-9]+}}]
  78 ; LS64: ld.local.u32 %{{.*}}, [%rd{{[0-9]+}}]
  79 ; ALL: ret
  80   %a = load i32, i32 addrspace(5)* %ptr
  81   ret i32 %a
  82 }
  83
  84 ;; i64
  85 define i64 @ld_global_i64(i64 addrspace(1)* %ptr) {
  86 ; ALL-LABEL: ld_global_i64
  87 ; G32: ld.global.u64 %{{.*}}, [%r{{[0-9]+}}]
  88 ; G64: ld.global.u64 %{{.*}}, [%rd{{[0-9]+}}]
  89 ; ALL: ret
  90   %a = load i64, i64 addrspace(1)* %ptr
  91   ret i64 %a
  92 }
  93 define i64 @ld_shared_i64(i64 addrspace(3)* %ptr) {
  94 ; ALL-LABEL: ld_shared_i64
  95 ; LS32: ld.shared.u64 %{{.*}}, [%r{{[0-9]+}}]
  96 ; LS64: ld.shared.u64 %{{.*}}, [%rd{{[0-9]+}}]
  97 ; ALL: ret
  98   %a = load i64, i64 addrspace(3)* %ptr
  99   ret i64 %a
 100 }
 101 define i64 @ld_local_i64(i64 addrspace(5)* %ptr) {
 102 ; ALL-LABEL: ld_local_i64
 103 ; LS32: ld.local.u64 %{{.*}}, [%r{{[0-9]+}}]
 104 ; LS64: ld.local.u64 %{{.*}}, [%rd{{[0-9]+}}]
 105 ; ALL: ret
 106   %a = load i64, i64 addrspace(5)* %ptr
 107   ret i64 %a
 108 }
 109
 110 ;; f32
 111 define float @ld_global_f32(float addrspace(1)* %ptr) {
 112 ; ALL-LABEL: ld_global_f32
 113 ; G32: ld.global.f32 %{{.*}}, [%r{{[0-9]+}}]
 114 ; G64: ld.global.f32 %{{.*}}, [%rd{{[0-9]+}}]
 115 ; ALL: ret
 116   %a = load float, float addrspace(1)* %ptr
 117   ret float %a
 118 }
 119 define float @ld_shared_f32(float addrspace(3)* %ptr) {
 120 ; ALL-LABEL: ld_shared_f32
 121 ; LS32: ld.shared.f32 %{{.*}}, [%r{{[0-9]+}}]
 122 ; LS64: ld.shared.f32 %{{.*}}, [%rd{{[0-9]+}}]
 123 ; ALL: ret
 124   %a = load float, float addrspace(3)* %ptr
 125   ret float %a
 126 }
 127 define float @ld_local_f32(float addrspace(5)* %ptr) {
 128 ; ALL-LABEL: ld_local_f32
 129 ; LS32: ld.local.f32 %{{.*}}, [%r{{[0-9]+}}]
 130 ; LS64: ld.local.f32 %{{.*}}, [%rd{{[0-9]+}}]
 131 ; ALL: ret
 132   %a = load float, float addrspace(5)* %ptr
 133   ret float %a
 134 }
 135
 136 ;; f64
 137 define double @ld_global_f64(double addrspace(1)* %ptr) {
 138 ; ALL-LABEL: ld_global_f64
 139 ; G32: ld.global.f64 %{{.*}}, [%r{{[0-9]+}}]
 140 ; G64: ld.global.f64 %{{.*}}, [%rd{{[0-9]+}}]
 141 ; ALL: ret
 142   %a = load double, double addrspace(1)* %ptr
 143   ret double %a
 144 }
 145 define double @ld_shared_f64(double addrspace(3)* %ptr) {
 146 ; ALL-LABEL: ld_shared_f64
 147 ; LS32: ld.shared.f64 %{{.*}}, [%r{{[0-9]+}}]
 148 ; LS64: ld.shared.f64 %{{.*}}, [%rd{{[0-9]+}}]
 149 ; ALL: ret
 150   %a = load double, double addrspace(3)* %ptr
 151   ret double %a
 152 }
 153 define double @ld_local_f64(double addrspace(5)* %ptr) {
 154 ; ALL-LABEL: ld_local_f64
 155 ; LS32: ld.local.f64 %{{.*}}, [%r{{[0-9]+}}]
 156 ; LS64: ld.local.f64 %{{.*}}, [%rd{{[0-9]+}}]
 157 ; ALL: ret
 158   %a = load double, double addrspace(5)* %ptr
 159   ret double %a
 160 }