llvm/test/CodeGen/NVPTX/addrspacecast.ll

   1 ; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,CLS32,G32
   2 ; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,NOPTRCONV,CLS64,G64
   3 ; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr| FileCheck %s -check-prefixes=ALL,PTRCONV,CLS64,G64
   4
   5 ; ALL-LABEL: conv1
   6 define i32 @conv1(i32 addrspace(1)* %ptr) {
   7 ; G32: cvta.global.u32
   8 ; ALL-NOT: cvt.u64.u32
   9 ; G64: cvta.global.u64
  10 ; ALL: ld.u32
  11   %genptr = addrspacecast i32 addrspace(1)* %ptr to i32*
  12   %val = load i32, i32* %genptr
  13   ret i32 %val
  14 }
  15
  16 ; ALL-LABEL: conv2
  17 define i32 @conv2(i32 addrspace(3)* %ptr) {
  18 ; CLS32: cvta.shared.u32
  19 ; PTRCONV: cvt.u64.u32
  20 ; NOPTRCONV-NOT: cvt.u64.u32
  21 ; CLS64: cvta.shared.u64
  22 ; ALL: ld.u32
  23   %genptr = addrspacecast i32 addrspace(3)* %ptr to i32*
  24   %val = load i32, i32* %genptr
  25   ret i32 %val
  26 }
  27
  28 ; ALL-LABEL: conv3
  29 define i32 @conv3(i32 addrspace(4)* %ptr) {
  30 ; CLS32: cvta.const.u32
  31 ; PTRCONV: cvt.u64.u32
  32 ; NOPTRCONV-NOT: cvt.u64.u32
  33 ; CLS64: cvta.const.u64
  34 ; ALL: ld.u32
  35   %genptr = addrspacecast i32 addrspace(4)* %ptr to i32*
  36   %val = load i32, i32* %genptr
  37   ret i32 %val
  38 }
  39
  40 ; ALL-LABEL: conv4
  41 define i32 @conv4(i32 addrspace(5)* %ptr) {
  42 ; CLS32: cvta.local.u32
  43 ; PTRCONV: cvt.u64.u32
  44 ; NOPTRCONV-NOT: cvt.u64.u32
  45 ; CLS64: cvta.local.u64
  46 ; ALL: ld.u32
  47   %genptr = addrspacecast i32 addrspace(5)* %ptr to i32*
  48   %val = load i32, i32* %genptr
  49   ret i32 %val
  50 }
  51
  52 ; ALL-LABEL: conv5
  53 define i32 @conv5(i32* %ptr) {
  54 ; CLS32: cvta.to.global.u32
  55 ; ALL-NOT: cvt.u64.u32
  56 ; CLS64: cvta.to.global.u64
  57 ; ALL: ld.global.u32
  58   %specptr = addrspacecast i32* %ptr to i32 addrspace(1)*
  59   %val = load i32, i32 addrspace(1)* %specptr
  60   ret i32 %val
  61 }
  62
  63 ; ALL-LABEL: conv6
  64 define i32 @conv6(i32* %ptr) {
  65 ; CLS32: cvta.to.shared.u32
  66 ; CLS64: cvta.to.shared.u64
  67 ; PTRCONV: cvt.u32.u64
  68 ; NOPTRCONV-NOT: cvt.u32.u64
  69 ; ALL: ld.shared.u32
  70   %specptr = addrspacecast i32* %ptr to i32 addrspace(3)*
  71   %val = load i32, i32 addrspace(3)* %specptr
  72   ret i32 %val
  73 }
  74
  75 ; ALL-LABEL: conv7
  76 define i32 @conv7(i32* %ptr) {
  77 ; CLS32: cvta.to.const.u32
  78 ; CLS64: cvta.to.const.u64
  79 ; PTRCONV: cvt.u32.u64
  80 ; NOPTRCONV-NOT: cvt.u32.u64
  81 ; ALL: ld.const.u32
  82   %specptr = addrspacecast i32* %ptr to i32 addrspace(4)*
  83   %val = load i32, i32 addrspace(4)* %specptr
  84   ret i32 %val
  85 }
  86
  87 ; ALL-LABEL: conv8
  88 define i32 @conv8(i32* %ptr) {
  89 ; CLS32: cvta.to.local.u32
  90 ; CLS64: cvta.to.local.u64
  91 ; PTRCONV: cvt.u32.u64
  92 ; NOPTRCONV-NOT: cvt.u32.u64
  93 ; ALL: ld.local.u32
  94   %specptr = addrspacecast i32* %ptr to i32 addrspace(5)*
  95   %val = load i32, i32 addrspace(5)* %specptr
  96   ret i32 %val
  97 }