llvm/test/CodeGen/NVPTX/st_async_release_multimem.ll

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

define void @test_multimem_st_async_release_sys_b8(ptr addrspace(1) %addr, i8 %value) {
; CHECK-LABEL: test_multimem_st_async_release_sys_b8(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [test_multimem_st_async_release_sys_b8_param_0];
; CHECK-NEXT:    ld.param.b8 %rs1, [test_multimem_st_async_release_sys_b8_param_1];
; CHECK-NEXT:    {
; CHECK-NEXT:    .reg .b8 %st_async_val;
; CHECK-NEXT:    cvt.u8.u16 %st_async_val, %rs1;
; CHECK-NEXT:    multimem.st.async.release.sys.global.b8 [%rd1], %st_async_val;
; CHECK-NEXT:    }
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.st.async.sys.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 true)
  ret void
}

define void @test_multimem_st_async_release_sys_b16(ptr addrspace(1) %addr, i16 %value) {
; CHECK-LABEL: test_multimem_st_async_release_sys_b16(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [test_multimem_st_async_release_sys_b16_param_0];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_multimem_st_async_release_sys_b16_param_1];
; CHECK-NEXT:    multimem.st.async.release.sys.global.b16 [%rd1], %rs1;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.st.async.sys.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 true)
  ret void
}

define void @test_multimem_st_async_release_sys_b32(ptr addrspace(1) %addr, i32 %value) {
; CHECK-LABEL: test_multimem_st_async_release_sys_b32(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [test_multimem_st_async_release_sys_b32_param_0];
; CHECK-NEXT:    ld.param.b32 %r1, [test_multimem_st_async_release_sys_b32_param_1];
; CHECK-NEXT:    multimem.st.async.release.sys.global.b32 [%rd1], %r1;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.st.async.sys.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 true)
  ret void
}

define void @test_multimem_st_async_release_sys_b64(ptr addrspace(1) %addr, i64 %value) {
; CHECK-LABEL: test_multimem_st_async_release_sys_b64(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [test_multimem_st_async_release_sys_b64_param_0];
; CHECK-NEXT:    ld.param.b64 %rd2, [test_multimem_st_async_release_sys_b64_param_1];
; CHECK-NEXT:    multimem.st.async.release.sys.global.b64 [%rd1], %rd2;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.st.async.sys.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 true)
  ret void
}

define void @test_multimem_st_async_release_gpu_b8(ptr addrspace(1) %addr, i8 %value) {
; CHECK-LABEL: test_multimem_st_async_release_gpu_b8(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [test_multimem_st_async_release_gpu_b8_param_0];
; CHECK-NEXT:    ld.param.b8 %rs1, [test_multimem_st_async_release_gpu_b8_param_1];
; CHECK-NEXT:    {
; CHECK-NEXT:    .reg .b8 %st_async_val;
; CHECK-NEXT:    cvt.u8.u16 %st_async_val, %rs1;
; CHECK-NEXT:    multimem.st.async.release.gpu.global.b8 [%rd1], %st_async_val;
; CHECK-NEXT:    }
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.st.async.gpu.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 true)
  ret void
}

define void @test_multimem_st_async_release_gpu_b16(ptr addrspace(1) %addr, i16 %value) {
; CHECK-LABEL: test_multimem_st_async_release_gpu_b16(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [test_multimem_st_async_release_gpu_b16_param_0];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_multimem_st_async_release_gpu_b16_param_1];
; CHECK-NEXT:    multimem.st.async.release.gpu.global.b16 [%rd1], %rs1;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.st.async.gpu.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 true)
  ret void
}

define void @test_multimem_st_async_release_gpu_b32(ptr addrspace(1) %addr, i32 %value) {
; CHECK-LABEL: test_multimem_st_async_release_gpu_b32(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [test_multimem_st_async_release_gpu_b32_param_0];
; CHECK-NEXT:    ld.param.b32 %r1, [test_multimem_st_async_release_gpu_b32_param_1];
; CHECK-NEXT:    multimem.st.async.release.gpu.global.b32 [%rd1], %r1;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.st.async.gpu.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 true)
  ret void
}

define void @test_multimem_st_async_release_gpu_b64(ptr addrspace(1) %addr, i64 %value) {
; CHECK-LABEL: test_multimem_st_async_release_gpu_b64(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [test_multimem_st_async_release_gpu_b64_param_0];
; CHECK-NEXT:    ld.param.b64 %rd2, [test_multimem_st_async_release_gpu_b64_param_1];
; CHECK-NEXT:    multimem.st.async.release.gpu.global.b64 [%rd1], %rd2;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.st.async.gpu.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 true)
  ret void
}