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 } 98
