[llvm] r239100 - [NVPTX] roll forward r239082

[Tobias Grosser]

On 06/04/2015 11:28 PM, Jingyue Wu wrote:
> Author: jingyue
> Date: Thu Jun  4 16:28:26 2015
> New Revision: 239100
>
> URL: http://llvm.org/viewvc/llvm-project?rev=239100&view=rev
> Log:
> [NVPTX] roll forward r239082
>
> NVPTXISelDAGToDAG translates "addrspacecast to param" to
> NVPTX::nvvm_ptr_gen_to_param

This commit breaks kernels where the arguments are already in the global 
address space (test case attached). Are such kernels invalid in some way?

Best,
Tobias
-------------- next part --------------
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s

target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-unknown-unknown"

; Verify that both %input and %output are converted to global pointers and then
; addrspacecast'ed back to the original type.
define void @kernel(float  addrspace(1)* %input, float  addrspace(1)* %output) {
; CHECK-LABEL: .visible .entry kernel(
; CHECK: cvta.to.global.u64
; CHECK: cvta.to.global.u64
  %1 = load float, float  addrspace(1)* %input, align 4
; CHECK: ld.global.f32
  store float %1, float addrspace(1)* %output, align 4
; CHECK: st.global.f32
  ret void
}

!nvvm.annotations = !{!0}
!0 = !{void (float addrspace(1)*, float addrspace(1)*)* @kernel, !"kernel", i32 1}