llvm/test/CodeGen/NVPTX/volatile-local.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx90 | FileCheck %s --check-prefix=PTX90
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx91 | FileCheck %s --check-prefix=PTX91
; RUN: %if ptxas-sm_90 && ptxas-isa-9.1 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx91 | %ptxas-verify -arch=sm_90 %}

define i32 @load_local_volatile(ptr addrspace(5) %p) {
; PTX90-LABEL: load_local_volatile(
; PTX90:       {
; PTX90-NEXT:    .reg .b32 %r<2>;
; PTX90-NEXT:    .reg .b64 %rd<2>;
; PTX90-EMPTY:
; PTX90-NEXT:  // %bb.0:
; PTX90-NEXT:    ld.param::func.b64 %rd1, [load_local_volatile_param_0];
; PTX90-NEXT:    ld.local.b32 %r1, [%rd1];
; PTX90-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX90-NEXT:    ret;
;
; PTX91-LABEL: load_local_volatile(
; PTX91:       {
; PTX91-NEXT:    .reg .b32 %r<2>;
; PTX91-NEXT:    .reg .b64 %rd<2>;
; PTX91-EMPTY:
; PTX91-NEXT:  // %bb.0:
; PTX91-NEXT:    ld.param::func.b64 %rd1, [load_local_volatile_param_0];
; PTX91-NEXT:    ld.volatile.local.b32 %r1, [%rd1];
; PTX91-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX91-NEXT:    ret;
  %v = load volatile i32, ptr addrspace(5) %p, align 4
  ret i32 %v
}

define void @store_local_volatile(ptr addrspace(5) %p, i32 %v) {
; PTX90-LABEL: store_local_volatile(
; PTX90:       {
; PTX90-NEXT:    .reg .b32 %r<2>;
; PTX90-NEXT:    .reg .b64 %rd<2>;
; PTX90-EMPTY:
; PTX90-NEXT:  // %bb.0:
; PTX90-NEXT:    ld.param::func.b64 %rd1, [store_local_volatile_param_0];
; PTX90-NEXT:    ld.param::func.b32 %r1, [store_local_volatile_param_1];
; PTX90-NEXT:    st.local.b32 [%rd1], %r1;
; PTX90-NEXT:    ret;
;
; PTX91-LABEL: store_local_volatile(
; PTX91:       {
; PTX91-NEXT:    .reg .b32 %r<2>;
; PTX91-NEXT:    .reg .b64 %rd<2>;
; PTX91-EMPTY:
; PTX91-NEXT:  // %bb.0:
; PTX91-NEXT:    ld.param::func.b64 %rd1, [store_local_volatile_param_0];
; PTX91-NEXT:    ld.param::func.b32 %r1, [store_local_volatile_param_1];
; PTX91-NEXT:    st.volatile.local.b32 [%rd1], %r1;
; PTX91-NEXT:    ret;
  store volatile i32 %v, ptr addrspace(5) %p, align 4
  ret void
}

; Local accesses discard atomic ordering but preserve volatile in PTX 9.1.
define i32 @local_volatile_unordered(ptr addrspace(5) %p) {
; PTX90-LABEL: local_volatile_unordered(
; PTX90:       {
; PTX90-NEXT:    .reg .b32 %r<2>;
; PTX90-NEXT:    .reg .b64 %rd<2>;
; PTX90-EMPTY:
; PTX90-NEXT:  // %bb.0:
; PTX90-NEXT:    ld.param::func.b64 %rd1, [local_volatile_unordered_param_0];
; PTX90-NEXT:    ld.local.b32 %r1, [%rd1];
; PTX90-NEXT:    st.local.b32 [%rd1], %r1;
; PTX90-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX90-NEXT:    ret;
;
; PTX91-LABEL: local_volatile_unordered(
; PTX91:       {
; PTX91-NEXT:    .reg .b32 %r<2>;
; PTX91-NEXT:    .reg .b64 %rd<2>;
; PTX91-EMPTY:
; PTX91-NEXT:  // %bb.0:
; PTX91-NEXT:    ld.param::func.b64 %rd1, [local_volatile_unordered_param_0];
; PTX91-NEXT:    ld.volatile.local.b32 %r1, [%rd1];
; PTX91-NEXT:    st.volatile.local.b32 [%rd1], %r1;
; PTX91-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX91-NEXT:    ret;
  %v = load atomic volatile i32, ptr addrspace(5) %p unordered, align 4
  store atomic volatile i32 %v, ptr addrspace(5) %p unordered, align 4
  ret i32 %v
}

define i32 @local_volatile_monotonic(ptr addrspace(5) %p) {
; PTX90-LABEL: local_volatile_monotonic(
; PTX90:       {
; PTX90-NEXT:    .reg .b32 %r<2>;
; PTX90-NEXT:    .reg .b64 %rd<2>;
; PTX90-EMPTY:
; PTX90-NEXT:  // %bb.0:
; PTX90-NEXT:    ld.param::func.b64 %rd1, [local_volatile_monotonic_param_0];
; PTX90-NEXT:    ld.local.b32 %r1, [%rd1];
; PTX90-NEXT:    st.local.b32 [%rd1], %r1;
; PTX90-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX90-NEXT:    ret;
;
; PTX91-LABEL: local_volatile_monotonic(
; PTX91:       {
; PTX91-NEXT:    .reg .b32 %r<2>;
; PTX91-NEXT:    .reg .b64 %rd<2>;
; PTX91-EMPTY:
; PTX91-NEXT:  // %bb.0:
; PTX91-NEXT:    ld.param::func.b64 %rd1, [local_volatile_monotonic_param_0];
; PTX91-NEXT:    ld.volatile.local.b32 %r1, [%rd1];
; PTX91-NEXT:    st.volatile.local.b32 [%rd1], %r1;
; PTX91-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX91-NEXT:    ret;
  %v = load atomic volatile i32, ptr addrspace(5) %p monotonic, align 4
  store atomic volatile i32 %v, ptr addrspace(5) %p monotonic, align 4
  ret i32 %v
}

define i32 @local_volatile_acquire_release(ptr addrspace(5) %p) {
; PTX90-LABEL: local_volatile_acquire_release(
; PTX90:       {
; PTX90-NEXT:    .reg .b32 %r<2>;
; PTX90-NEXT:    .reg .b64 %rd<2>;
; PTX90-EMPTY:
; PTX90-NEXT:  // %bb.0:
; PTX90-NEXT:    ld.param::func.b64 %rd1, [local_volatile_acquire_release_param_0];
; PTX90-NEXT:    ld.local.b32 %r1, [%rd1];
; PTX90-NEXT:    st.local.b32 [%rd1], %r1;
; PTX90-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX90-NEXT:    ret;
;
; PTX91-LABEL: local_volatile_acquire_release(
; PTX91:       {
; PTX91-NEXT:    .reg .b32 %r<2>;
; PTX91-NEXT:    .reg .b64 %rd<2>;
; PTX91-EMPTY:
; PTX91-NEXT:  // %bb.0:
; PTX91-NEXT:    ld.param::func.b64 %rd1, [local_volatile_acquire_release_param_0];
; PTX91-NEXT:    ld.volatile.local.b32 %r1, [%rd1];
; PTX91-NEXT:    st.volatile.local.b32 [%rd1], %r1;
; PTX91-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX91-NEXT:    ret;
  %v = load atomic volatile i32, ptr addrspace(5) %p acquire, align 4
  store atomic volatile i32 %v, ptr addrspace(5) %p release, align 4
  ret i32 %v
}

define i32 @local_volatile_seq_cst(ptr addrspace(5) %p) {
; PTX90-LABEL: local_volatile_seq_cst(
; PTX90:       {
; PTX90-NEXT:    .reg .b32 %r<2>;
; PTX90-NEXT:    .reg .b64 %rd<2>;
; PTX90-EMPTY:
; PTX90-NEXT:  // %bb.0:
; PTX90-NEXT:    ld.param::func.b64 %rd1, [local_volatile_seq_cst_param_0];
; PTX90-NEXT:    ld.local.b32 %r1, [%rd1];
; PTX90-NEXT:    st.local.b32 [%rd1], %r1;
; PTX90-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX90-NEXT:    ret;
;
; PTX91-LABEL: local_volatile_seq_cst(
; PTX91:       {
; PTX91-NEXT:    .reg .b32 %r<2>;
; PTX91-NEXT:    .reg .b64 %rd<2>;
; PTX91-EMPTY:
; PTX91-NEXT:  // %bb.0:
; PTX91-NEXT:    ld.param::func.b64 %rd1, [local_volatile_seq_cst_param_0];
; PTX91-NEXT:    ld.volatile.local.b32 %r1, [%rd1];
; PTX91-NEXT:    st.volatile.local.b32 [%rd1], %r1;
; PTX91-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX91-NEXT:    ret;
  %v = load atomic volatile i32, ptr addrspace(5) %p seq_cst, align 4
  store atomic volatile i32 %v, ptr addrspace(5) %p seq_cst, align 4
  ret i32 %v
}