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