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