| 
1 | 1 | ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5  | 
2 |  | -; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK-PTX64 %s  | 
3 |  | -; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| FileCheck --check-prefixes=CHECK-PTX-SHARED32 %s  | 
 | 2 | +; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK,CHECK-PTX64 %s  | 
 | 3 | +; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s  | 
4 | 4 | ; RUN: %if ptxas-12.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}  | 
5 | 5 | ; RUN: %if ptxas-12.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| %ptxas-verify -arch=sm_90 %}  | 
6 | 6 | 
 
  | 
@@ -119,18 +119,18 @@ define void @cp_async_bulk_cta_to_cluster(ptr addrspace(3) %src, ptr addrspace(3  | 
119 | 119 | }  | 
120 | 120 | 
 
  | 
121 | 121 | define void @cp_async_bulk_prefetch(ptr addrspace(1) %src, i32 %size, i64 %ch) {  | 
122 |  | -; CHECK-PTX64-LABEL: cp_async_bulk_prefetch(  | 
123 |  | -; CHECK-PTX64:       {  | 
124 |  | -; CHECK-PTX64-NEXT:    .reg .b32 %r<2>;  | 
125 |  | -; CHECK-PTX64-NEXT:    .reg .b64 %rd<3>;  | 
126 |  | -; CHECK-PTX64-EMPTY:  | 
127 |  | -; CHECK-PTX64-NEXT:  // %bb.0:  | 
128 |  | -; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_prefetch_param_0];  | 
129 |  | -; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_prefetch_param_1];  | 
130 |  | -; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_prefetch_param_2];  | 
131 |  | -; CHECK-PTX64-NEXT:    cp.async.bulk.prefetch.L2.global.L2::cache_hint [%rd1], %r1, %rd2;  | 
132 |  | -; CHECK-PTX64-NEXT:    cp.async.bulk.prefetch.L2.global [%rd1], %r1;  | 
133 |  | -; CHECK-PTX64-NEXT:    ret;  | 
 | 122 | +; CHECK-LABEL: cp_async_bulk_prefetch(  | 
 | 123 | +; CHECK:       {  | 
 | 124 | +; CHECK-NEXT:    .reg .b32 %r<2>;  | 
 | 125 | +; CHECK-NEXT:    .reg .b64 %rd<3>;  | 
 | 126 | +; CHECK-EMPTY:  | 
 | 127 | +; CHECK-NEXT:  // %bb.0:  | 
 | 128 | +; CHECK-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_prefetch_param_0];  | 
 | 129 | +; CHECK-NEXT:    ld.param.u32 %r1, [cp_async_bulk_prefetch_param_1];  | 
 | 130 | +; CHECK-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_prefetch_param_2];  | 
 | 131 | +; CHECK-NEXT:    cp.async.bulk.prefetch.L2.global.L2::cache_hint [%rd1], %r1, %rd2;  | 
 | 132 | +; CHECK-NEXT:    cp.async.bulk.prefetch.L2.global [%rd1], %r1;  | 
 | 133 | +; CHECK-NEXT:    ret;  | 
134 | 134 |   tail call void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1) %src, i32 %size, i64 %ch, i1 1)  | 
135 | 135 |   tail call void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1) %src, i32 %size, i64 0, i1 0)  | 
136 | 136 |   ret void  | 
 | 
0 commit comments