llvm/test/CodeGen/NVPTX/reserved-smem-offset.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_80 -mattr=+ptx76 | FileCheck %s
; RUN: %if ptxas-sm_80 && ptxas-isa-7.6 %{ llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_80 -mattr=+ptx76 | %ptxas-verify -arch=sm_80 %}

declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_begin()
declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_end()
declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_cap()
declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_0()
declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_1()

define i32 @test() {
; CHECK-LABEL: test(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<10>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    mov.u32 %r1, %reserved_smem_offset_begin;
; CHECK-NEXT:    mov.u32 %r2, %reserved_smem_offset_end;
; CHECK-NEXT:    mov.u32 %r3, %reserved_smem_offset_cap;
; CHECK-NEXT:    mov.u32 %r4, %reserved_smem_offset_0;
; CHECK-NEXT:    mov.u32 %r5, %reserved_smem_offset_1;
; CHECK-NEXT:    add.s32 %r6, %r1, %r2;
; CHECK-NEXT:    add.s32 %r7, %r6, %r3;
; CHECK-NEXT:    add.s32 %r8, %r7, %r4;
; CHECK-NEXT:    add.s32 %r9, %r8, %r5;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r9;
; CHECK-NEXT:    ret;
  %begin = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_begin()
  %end = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_end()
  %cap = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_cap()
  %offset0 = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_0()
  %offset1 = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_1()

  %ret0 = add i32 %begin, %end
  %ret1 = add i32 %ret0, %cap
  %ret2 = add i32 %ret1, %offset0
  %ret3 = add i32 %ret2, %offset1

  ret i32 %ret3
}