[llvm] [NVPTX] Add asynchronous store intrinsics (PR #200768)

Durgadoss R via llvm-commits llvm-commits at lists.llvm.org
Tue Jun 9 05:17:35 PDT 2026


================
@@ -0,0 +1,192 @@
+; 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=+ptx87 | FileCheck %s
+; RUN: %if ptxas-sm_100 && ptxas-isa-8.7 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx87 | %ptxas-verify -arch=sm_100 %}
+
+define void @test_st_async_release_sys_b8(ptr addrspace(1) %addr, i8 %value) {
+; CHECK-LABEL: test_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_st_async_release_sys_b8_param_0];
+; CHECK-NEXT:    ld.param.b8 %rs1, [test_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:    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 false)
----------------
durga4github wrote:

nice to see the pretty-printing usage..

https://github.com/llvm/llvm-project/pull/200768


More information about the llvm-commits mailing list