1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
|
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mattr=+ptx73 -mcpu=sm_52 | FileCheck %s
target triple = "nvptx64-nvidia-cuda"
define void @foo(i64 %a, ptr %p0, ptr %p1) {
; CHECK-LABEL: foo(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<10>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [foo_param_0];
; CHECK-NEXT: add.s64 %rd2, %rd1, 7;
; CHECK-NEXT: and.b64 %rd3, %rd2, -8;
; CHECK-NEXT: alloca.u64 %rd4, %rd3, 16;
; CHECK-NEXT: cvta.local.u64 %rd5, %rd4;
; CHECK-NEXT: ld.param.b64 %rd6, [foo_param_1];
; CHECK-NEXT: alloca.u64 %rd7, %rd3, 16;
; CHECK-NEXT: cvta.local.u64 %rd8, %rd7;
; CHECK-NEXT: ld.param.b64 %rd9, [foo_param_2];
; CHECK-NEXT: st.b64 [%rd6], %rd5;
; CHECK-NEXT: st.b64 [%rd9], %rd8;
; CHECK-NEXT: ret;
%b = alloca i8, i64 %a, align 16
%c = alloca i8, i64 %a, align 16
store ptr %b, ptr %p0, align 8
store ptr %c, ptr %p1, align 8
ret void
}
|