test/CodeGen/NVPTX/disjoint-or-addr.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
target triple = "nvptx64-nvidia-cuda"

@a = external global ptr align 16

define i32  @test_disjoint_or_addr(i16 %a) {
; CHECK-LABEL: test_disjoint_or_addr(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    mov.b64 %rd1, a;
; CHECK-NEXT:    cvta.global.u64 %rd2, %rd1;
; CHECK-NEXT:    ld.b32 %r1, [%rd2+8];
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
  %a1 = ptrtoint ptr @a to i64
  %a2 = or disjoint i64 %a1, 8
  %a3 = inttoptr i64 %a2 to ptr
  %v = load i32, ptr %a3
  ret i32 %v
}