|
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