|
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