194 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK,CHECK-PTX64 %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s4; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}5; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| %ptxas-verify -arch=sm_90 %}6 7target triple = "nvptx64-nvidia-cuda"8 9declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7), ptr addrspace(3), ptr addrspace(1), i32, i16, i64, i1, i1)10declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.global(ptr addrspace(1), ptr addrspace(3), i32, i64, i1)11declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(7), ptr addrspace(3), ptr addrspace(3), i32)12declare void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1), i32, i64, i1)13 14define void @cp_async_bulk_g2s(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(7) %dst, i32 %size, i16 %mc, i64 %ch) {15; CHECK-PTX64-LABEL: cp_async_bulk_g2s(16; CHECK-PTX64: {17; CHECK-PTX64-NEXT: .reg .b16 %rs<2>;18; CHECK-PTX64-NEXT: .reg .b32 %r<2>;19; CHECK-PTX64-NEXT: .reg .b64 %rd<5>;20; CHECK-PTX64-EMPTY:21; CHECK-PTX64-NEXT: // %bb.0:22; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [cp_async_bulk_g2s_param_0];23; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [cp_async_bulk_g2s_param_1];24; CHECK-PTX64-NEXT: ld.param.b64 %rd3, [cp_async_bulk_g2s_param_2];25; CHECK-PTX64-NEXT: ld.param.b32 %r1, [cp_async_bulk_g2s_param_3];26; CHECK-PTX64-NEXT: ld.param.b16 %rs1, [cp_async_bulk_g2s_param_4];27; CHECK-PTX64-NEXT: ld.param.b64 %rd4, [cp_async_bulk_g2s_param_5];28; CHECK-PTX64-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%rd3], [%rd1], %r1, [%rd2];29; CHECK-PTX64-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %rd4;30; CHECK-PTX64-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster [%rd3], [%rd1], %r1, [%rd2], %rs1;31; CHECK-PTX64-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %rs1, %rd4;32; CHECK-PTX64-NEXT: ret;33;34; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_g2s(35; CHECK-PTX-SHARED32: {36; CHECK-PTX-SHARED32-NEXT: .reg .b16 %rs<2>;37; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<4>;38; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<3>;39; CHECK-PTX-SHARED32-EMPTY:40; CHECK-PTX-SHARED32-NEXT: // %bb.0:41; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd1, [cp_async_bulk_g2s_param_0];42; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [cp_async_bulk_g2s_param_1];43; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [cp_async_bulk_g2s_param_2];44; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r3, [cp_async_bulk_g2s_param_3];45; CHECK-PTX-SHARED32-NEXT: ld.param.b16 %rs1, [cp_async_bulk_g2s_param_4];46; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd2, [cp_async_bulk_g2s_param_5];47; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%r2], [%rd1], %r3, [%r1];48; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.L2::cache_hint [%r2], [%rd1], %r3, [%r1], %rd2;49; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster [%r2], [%rd1], %r3, [%r1], %rs1;50; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster.L2::cache_hint [%r2], [%rd1], %r3, [%r1], %rs1, %rd2;51; CHECK-PTX-SHARED32-NEXT: ret;52 tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 0, i1 0)53 tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 0, i1 1)54 tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 1, i1 0)55 tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 1, i1 1)56 ret void57}58 59; Tests to verify that the immediate values for "mc and ch" land correctly in PTX.60; The values of 16 and 64 are arbitrary and do not have any significance.61define void @cp_async_bulk_g2s_imm_mc_ch(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(7) %dst, i32 %size, i16 %mc, i64 %ch) {62; CHECK-PTX64-LABEL: cp_async_bulk_g2s_imm_mc_ch(63; CHECK-PTX64: {64; CHECK-PTX64-NEXT: .reg .b16 %rs<3>;65; CHECK-PTX64-NEXT: .reg .b32 %r<2>;66; CHECK-PTX64-NEXT: .reg .b64 %rd<6>;67; CHECK-PTX64-EMPTY:68; CHECK-PTX64-NEXT: // %bb.0:69; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [cp_async_bulk_g2s_imm_mc_ch_param_0];70; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [cp_async_bulk_g2s_imm_mc_ch_param_1];71; CHECK-PTX64-NEXT: ld.param.b64 %rd3, [cp_async_bulk_g2s_imm_mc_ch_param_2];72; CHECK-PTX64-NEXT: ld.param.b32 %r1, [cp_async_bulk_g2s_imm_mc_ch_param_3];73; CHECK-PTX64-NEXT: mov.b64 %rd4, 64;74; CHECK-PTX64-NEXT: mov.b16 %rs1, 16;75; CHECK-PTX64-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %rs1, %rd4;76; CHECK-PTX64-NEXT: ld.param.b64 %rd5, [cp_async_bulk_g2s_imm_mc_ch_param_5];77; CHECK-PTX64-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster [%rd3], [%rd1], %r1, [%rd2], %rs1;78; CHECK-PTX64-NEXT: ld.param.b16 %rs2, [cp_async_bulk_g2s_imm_mc_ch_param_4];79; CHECK-PTX64-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %rd4;80; CHECK-PTX64-NEXT: ret;81;82; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_g2s_imm_mc_ch(83; CHECK-PTX-SHARED32: {84; CHECK-PTX-SHARED32-NEXT: .reg .b16 %rs<3>;85; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<4>;86; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<4>;87; CHECK-PTX-SHARED32-EMPTY:88; CHECK-PTX-SHARED32-NEXT: // %bb.0:89; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd1, [cp_async_bulk_g2s_imm_mc_ch_param_0];90; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [cp_async_bulk_g2s_imm_mc_ch_param_1];91; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [cp_async_bulk_g2s_imm_mc_ch_param_2];92; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r3, [cp_async_bulk_g2s_imm_mc_ch_param_3];93; CHECK-PTX-SHARED32-NEXT: mov.b64 %rd2, 64;94; CHECK-PTX-SHARED32-NEXT: mov.b16 %rs1, 16;95; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster.L2::cache_hint [%r2], [%rd1], %r3, [%r1], %rs1, %rd2;96; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd3, [cp_async_bulk_g2s_imm_mc_ch_param_5];97; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster [%r2], [%rd1], %r3, [%r1], %rs1;98; CHECK-PTX-SHARED32-NEXT: ld.param.b16 %rs2, [cp_async_bulk_g2s_imm_mc_ch_param_4];99; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.L2::cache_hint [%r2], [%rd1], %r3, [%r1], %rd2;100; CHECK-PTX-SHARED32-NEXT: ret;101 tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 16, i64 64, i1 1, i1 1)102 tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 16, i64 %ch, i1 1, i1 0)103 tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 64, i1 0, i1 1)104 ret void105}106 107define void @cp_async_bulk_s2g(ptr addrspace(3) %src, ptr addrspace(1) %dst, i32 %size, i64 %ch) {108; CHECK-PTX64-LABEL: cp_async_bulk_s2g(109; CHECK-PTX64: {110; CHECK-PTX64-NEXT: .reg .b32 %r<2>;111; CHECK-PTX64-NEXT: .reg .b64 %rd<5>;112; CHECK-PTX64-EMPTY:113; CHECK-PTX64-NEXT: // %bb.0:114; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [cp_async_bulk_s2g_param_0];115; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [cp_async_bulk_s2g_param_1];116; CHECK-PTX64-NEXT: ld.param.b32 %r1, [cp_async_bulk_s2g_param_2];117; CHECK-PTX64-NEXT: mov.b64 %rd3, 64;118; CHECK-PTX64-NEXT: cp.async.bulk.global.shared::cta.bulk_group.L2::cache_hint [%rd2], [%rd1], %r1, %rd3;119; CHECK-PTX64-NEXT: ld.param.b64 %rd4, [cp_async_bulk_s2g_param_3];120; CHECK-PTX64-NEXT: cp.async.bulk.global.shared::cta.bulk_group [%rd2], [%rd1], %r1;121; CHECK-PTX64-NEXT: cp.async.bulk.global.shared::cta.bulk_group.L2::cache_hint [%rd2], [%rd1], %r1, %rd4;122; CHECK-PTX64-NEXT: ret;123;124; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_s2g(125; CHECK-PTX-SHARED32: {126; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;127; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<4>;128; CHECK-PTX-SHARED32-EMPTY:129; CHECK-PTX-SHARED32-NEXT: // %bb.0:130; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [cp_async_bulk_s2g_param_0];131; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd1, [cp_async_bulk_s2g_param_1];132; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [cp_async_bulk_s2g_param_2];133; CHECK-PTX-SHARED32-NEXT: mov.b64 %rd2, 64;134; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.global.shared::cta.bulk_group.L2::cache_hint [%rd1], [%r1], %r2, %rd2;135; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd3, [cp_async_bulk_s2g_param_3];136; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.global.shared::cta.bulk_group [%rd1], [%r1], %r2;137; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.global.shared::cta.bulk_group.L2::cache_hint [%rd1], [%r1], %r2, %rd3;138; CHECK-PTX-SHARED32-NEXT: ret;139 tail call void @llvm.nvvm.cp.async.bulk.shared.cta.to.global(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 64, i1 1)140 tail call void @llvm.nvvm.cp.async.bulk.shared.cta.to.global(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 %ch, i1 0)141 tail call void @llvm.nvvm.cp.async.bulk.shared.cta.to.global(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 %ch, i1 1)142 ret void143}144 145define void @cp_async_bulk_cta_to_cluster(ptr addrspace(3) %src, ptr addrspace(3) %bar, ptr addrspace(7) %dst, i32 %size) {146; CHECK-PTX64-LABEL: cp_async_bulk_cta_to_cluster(147; CHECK-PTX64: {148; CHECK-PTX64-NEXT: .reg .b32 %r<2>;149; CHECK-PTX64-NEXT: .reg .b64 %rd<4>;150; CHECK-PTX64-EMPTY:151; CHECK-PTX64-NEXT: // %bb.0:152; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [cp_async_bulk_cta_to_cluster_param_0];153; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [cp_async_bulk_cta_to_cluster_param_1];154; CHECK-PTX64-NEXT: ld.param.b64 %rd3, [cp_async_bulk_cta_to_cluster_param_2];155; CHECK-PTX64-NEXT: ld.param.b32 %r1, [cp_async_bulk_cta_to_cluster_param_3];156; CHECK-PTX64-NEXT: cp.async.bulk.shared::cluster.shared::cta.mbarrier::complete_tx::bytes [%rd3], [%rd1], %r1, [%rd2];157; CHECK-PTX64-NEXT: ret;158;159; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_cta_to_cluster(160; CHECK-PTX-SHARED32: {161; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<5>;162; CHECK-PTX-SHARED32-EMPTY:163; CHECK-PTX-SHARED32-NEXT: // %bb.0:164; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [cp_async_bulk_cta_to_cluster_param_0];165; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [cp_async_bulk_cta_to_cluster_param_1];166; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r3, [cp_async_bulk_cta_to_cluster_param_2];167; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r4, [cp_async_bulk_cta_to_cluster_param_3];168; CHECK-PTX-SHARED32-NEXT: cp.async.bulk.shared::cluster.shared::cta.mbarrier::complete_tx::bytes [%r3], [%r1], %r4, [%r2];169; CHECK-PTX-SHARED32-NEXT: ret;170 tail call void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(3) %src, i32 %size)171 ret void172}173 174define void @cp_async_bulk_prefetch(ptr addrspace(1) %src, i32 %size, i64 %ch) {175; CHECK-LABEL: cp_async_bulk_prefetch(176; CHECK: {177; CHECK-NEXT: .reg .b32 %r<2>;178; CHECK-NEXT: .reg .b64 %rd<4>;179; CHECK-EMPTY:180; CHECK-NEXT: // %bb.0:181; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_prefetch_param_0];182; CHECK-NEXT: ld.param.b32 %r1, [cp_async_bulk_prefetch_param_1];183; CHECK-NEXT: mov.b64 %rd2, 64;184; CHECK-NEXT: cp.async.bulk.prefetch.L2.global.L2::cache_hint [%rd1], %r1, %rd2;185; CHECK-NEXT: ld.param.b64 %rd3, [cp_async_bulk_prefetch_param_2];186; CHECK-NEXT: cp.async.bulk.prefetch.L2.global.L2::cache_hint [%rd1], %r1, %rd3;187; CHECK-NEXT: cp.async.bulk.prefetch.L2.global [%rd1], %r1;188; CHECK-NEXT: ret;189 tail call void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1) %src, i32 %size, i64 64, i1 1)190 tail call void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1) %src, i32 %size, i64 %ch, i1 1)191 tail call void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1) %src, i32 %size, i64 %ch, i1 0)192 ret void193}194