112 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx86 | FileCheck %s --check-prefixes=CHECK,CHECK-PTX-SHARED643; RUN: llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s4; RUN: %if ptxas-sm_100 && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 | %ptxas-verify -arch=sm_100 %}5; RUN: %if ptxas-sm_100 && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_100 %}6 7define void @nvvm_clusterlaunchcontrol_try_cancel(8; CHECK-PTX-SHARED64-LABEL: nvvm_clusterlaunchcontrol_try_cancel(9; CHECK-PTX-SHARED64: {10; CHECK-PTX-SHARED64-NEXT: .reg .b64 %rd<3>;11; CHECK-PTX-SHARED64-EMPTY:12; CHECK-PTX-SHARED64-NEXT: // %bb.0:13; CHECK-PTX-SHARED64-NEXT: ld.param.b64 %rd1, [nvvm_clusterlaunchcontrol_try_cancel_param_0];14; CHECK-PTX-SHARED64-NEXT: ld.param.b64 %rd2, [nvvm_clusterlaunchcontrol_try_cancel_param_1];15; CHECK-PTX-SHARED64-NEXT: clusterlaunchcontrol.try_cancel.async.shared::cta.mbarrier::complete_tx::bytes.b128 [%rd1], [%rd2];16; CHECK-PTX-SHARED64-NEXT: ret;17;18; CHECK-PTX-SHARED32-LABEL: nvvm_clusterlaunchcontrol_try_cancel(19; CHECK-PTX-SHARED32: {20; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;21; CHECK-PTX-SHARED32-EMPTY:22; CHECK-PTX-SHARED32-NEXT: // %bb.0:23; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [nvvm_clusterlaunchcontrol_try_cancel_param_0];24; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [nvvm_clusterlaunchcontrol_try_cancel_param_1];25; CHECK-PTX-SHARED32-NEXT: clusterlaunchcontrol.try_cancel.async.shared::cta.mbarrier::complete_tx::bytes.b128 [%r1], [%r2];26; CHECK-PTX-SHARED32-NEXT: ret;27 ptr addrspace(3) %saddr, ptr addrspace(3) %smbar,28 i128 %try_cancel_response) {29 30 tail call void @llvm.nvvm.clusterlaunchcontrol.try_cancel.async.shared(ptr addrspace(3) %saddr, ptr addrspace(3) %smbar)31 ret void;32}33 34define i32 @nvvm_clusterlaunchcontrol_query_cancel_is_canceled(i128 %try_cancel_response) local_unnamed_addr #0 {35; CHECK-LABEL: nvvm_clusterlaunchcontrol_query_cancel_is_canceled(36; CHECK: {37; CHECK-NEXT: .reg .pred %p<2>;38; CHECK-NEXT: .reg .b32 %r<2>;39; CHECK-NEXT: .reg .b64 %rd<3>;40; CHECK-EMPTY:41; CHECK-NEXT: // %bb.0:42; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [nvvm_clusterlaunchcontrol_query_cancel_is_canceled_param_0];43; CHECK-NEXT: {44; CHECK-NEXT: .reg .b128 %clc_handle;45; CHECK-NEXT: mov.b128 %clc_handle, {%rd1, %rd2};46; CHECK-NEXT: clusterlaunchcontrol.query_cancel.is_canceled.pred.b128 %p1, %clc_handle;47; CHECK-NEXT: }48; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;49; CHECK-NEXT: st.param.b32 [func_retval0], %r1;50; CHECK-NEXT: ret;51 %v0 = call i1 @llvm.nvvm.clusterlaunchcontrol.query_cancel.is_canceled(i128 %try_cancel_response)52 %v2 = zext i1 %v0 to i3253 ret i32 %v2;54}55 56define i32 @nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_x(i128 %try_cancel_response) local_unnamed_addr #0 {57; CHECK-LABEL: nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_x(58; CHECK: {59; CHECK-NEXT: .reg .b32 %r<2>;60; CHECK-NEXT: .reg .b64 %rd<3>;61; CHECK-EMPTY:62; CHECK-NEXT: // %bb.0:63; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_x_param_0];64; CHECK-NEXT: {65; CHECK-NEXT: .reg .b128 %clc_handle;66; CHECK-NEXT: mov.b128 %clc_handle, {%rd1, %rd2};67; CHECK-NEXT: clusterlaunchcontrol.query_cancel.get_first_ctaid::x.b32.b128 %r1, %clc_handle;68; CHECK-NEXT: }69; CHECK-NEXT: st.param.b32 [func_retval0], %r1;70; CHECK-NEXT: ret;71 %v0 = call i32 @llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.x(i128 %try_cancel_response)72 ret i32 %v0;73}74 75define i32 @nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_y(i128 %try_cancel_response) local_unnamed_addr #0 {76; CHECK-LABEL: nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_y(77; CHECK: {78; CHECK-NEXT: .reg .b32 %r<2>;79; CHECK-NEXT: .reg .b64 %rd<3>;80; CHECK-EMPTY:81; CHECK-NEXT: // %bb.0:82; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_y_param_0];83; CHECK-NEXT: {84; CHECK-NEXT: .reg .b128 %clc_handle;85; CHECK-NEXT: mov.b128 %clc_handle, {%rd1, %rd2};86; CHECK-NEXT: clusterlaunchcontrol.query_cancel.get_first_ctaid::y.b32.b128 %r1, %clc_handle;87; CHECK-NEXT: }88; CHECK-NEXT: st.param.b32 [func_retval0], %r1;89; CHECK-NEXT: ret;90 %v0 = call i32 @llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.y(i128 %try_cancel_response)91 ret i32 %v0;92}93 94define i32 @nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_z(i128 %try_cancel_response) local_unnamed_addr #0 {95; CHECK-LABEL: nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_z(96; CHECK: {97; CHECK-NEXT: .reg .b32 %r<2>;98; CHECK-NEXT: .reg .b64 %rd<3>;99; CHECK-EMPTY:100; CHECK-NEXT: // %bb.0:101; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_z_param_0];102; CHECK-NEXT: {103; CHECK-NEXT: .reg .b128 %clc_handle;104; CHECK-NEXT: mov.b128 %clc_handle, {%rd1, %rd2};105; CHECK-NEXT: clusterlaunchcontrol.query_cancel.get_first_ctaid::z.b32.b128 %r1, %clc_handle;106; CHECK-NEXT: }107; CHECK-NEXT: st.param.b32 [func_retval0], %r1;108; CHECK-NEXT: ret;109 %v0 = call i32 @llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.z(i128 %try_cancel_response)110 ret i32 %v0;111}112