; PTX64: cvta.global.u64
; PTX64: ld.u32
%genptr = addrspacecast i32 addrspace(1)* %ptr to i32*
- %val = load i32* %genptr
+ %val = load i32, i32* %genptr
ret i32 %val
}
; PTX64: cvta.shared.u64
; PTX64: ld.u32
%genptr = addrspacecast i32 addrspace(3)* %ptr to i32*
- %val = load i32* %genptr
+ %val = load i32, i32* %genptr
ret i32 %val
}
; PTX64: cvta.const.u64
; PTX64: ld.u32
%genptr = addrspacecast i32 addrspace(4)* %ptr to i32*
- %val = load i32* %genptr
+ %val = load i32, i32* %genptr
ret i32 %val
}
; PTX64: cvta.local.u64
; PTX64: ld.u32
%genptr = addrspacecast i32 addrspace(5)* %ptr to i32*
- %val = load i32* %genptr
+ %val = load i32, i32* %genptr
ret i32 %val
}
; PTX64: cvta.to.global.u64
; PTX64: ld.global.u32
%specptr = addrspacecast i32* %ptr to i32 addrspace(1)*
- %val = load i32 addrspace(1)* %specptr
+ %val = load i32, i32 addrspace(1)* %specptr
ret i32 %val
}
; PTX64: cvta.to.shared.u64
; PTX64: ld.shared.u32
%specptr = addrspacecast i32* %ptr to i32 addrspace(3)*
- %val = load i32 addrspace(3)* %specptr
+ %val = load i32, i32 addrspace(3)* %specptr
ret i32 %val
}
; PTX64: cvta.to.const.u64
; PTX64: ld.const.u32
%specptr = addrspacecast i32* %ptr to i32 addrspace(4)*
- %val = load i32 addrspace(4)* %specptr
+ %val = load i32, i32 addrspace(4)* %specptr
ret i32 %val
}
; PTX64: cvta.to.local.u64
; PTX64: ld.local.u32
%specptr = addrspacecast i32* %ptr to i32 addrspace(5)*
- %val = load i32 addrspace(5)* %specptr
+ %val = load i32, i32 addrspace(5)* %specptr
ret i32 %val
}