pre-commit tests -- use update_llc_test_checks.py

Author: AlexMaclean

llvm/test/CodeGen/NVPTX/lower-args.ll

@@ -1,7 +1,8 @@
-; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-cuda | FileCheck %s --check-prefixes COMMON,IR,IRC
-; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-nvcl | FileCheck %s --check-prefixes COMMON,IR,IRO
-; RUN: llc < %s -mcpu=sm_20 --mtriple nvptx64-nvidia-cuda | FileCheck %s --check-prefixes COMMON,PTX,PTXC
-; RUN: llc < %s -mcpu=sm_20 --mtriple nvptx64-nvidia-nvcl| FileCheck %s --check-prefixes COMMON,PTX,PTXO
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-cuda | FileCheck %s --check-prefixes IR,IRC
+; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-nvcl | FileCheck %s --check-prefixes IR,IRO
+; RUN: llc < %s -mcpu=sm_20 --mtriple nvptx64-nvidia-cuda | FileCheck %s --check-prefixes PTX,PTXC
+; RUN: llc < %s -mcpu=sm_20 --mtriple nvptx64-nvidia-nvcl| FileCheck %s --check-prefixes PTX,PTXO
 ; RUN: %if ptxas %{ llc < %s -mcpu=sm_20 | %ptxas-verify %}
 
 target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
@@ -12,12 +13,60 @@ target triple = "nvptx64-nvidia-cuda"
 %class.padded = type { i8, i32 }
 
 ; Check that nvptx-lower-args preserves arg alignment
-; COMMON-LABEL: load_alignment
 define void @load_alignment(ptr nocapture readonly byval(%class.outer) align 8 %arg) {
+; IR-LABEL: define void @load_alignment(
+; IR-SAME: ptr readonly byval([[CLASS_OUTER:%.*]]) align 8 captures(none) [[ARG:%.*]]) {
+; IR-NEXT:  [[ENTRY:.*:]]
+; IR-NEXT:    [[ARG1:%.*]] = alloca [[CLASS_OUTER]], align 8
+; IR-NEXT:    [[ARG2:%.*]] = addrspacecast ptr [[ARG]] to ptr addrspace(101)
+; IR-NEXT:    call void @llvm.memcpy.p0.p101.i64(ptr align 8 [[ARG1]], ptr addrspace(101) align 8 [[ARG2]], i64 24, i1 false)
+; IR-NEXT:    [[ARG_IDX_VAL:%.*]] = load ptr, ptr [[ARG1]], align 8
+; IR-NEXT:    [[ARG_IDX1:%.*]] = getelementptr [[CLASS_OUTER]], ptr [[ARG1]], i64 0, i32 0, i32 1
+; IR-NEXT:    [[ARG_IDX1_VAL:%.*]] = load ptr, ptr [[ARG_IDX1]], align 8
+; IR-NEXT:    [[ARG_IDX2:%.*]] = getelementptr [[CLASS_OUTER]], ptr [[ARG1]], i64 0, i32 1
+; IR-NEXT:    [[ARG_IDX2_VAL:%.*]] = load i32, ptr [[ARG_IDX2]], align 8
+; IR-NEXT:    [[ARG_IDX_VAL_VAL:%.*]] = load i32, ptr [[ARG_IDX_VAL]], align 4
+; IR-NEXT:    [[ADD_I:%.*]] = add nsw i32 [[ARG_IDX_VAL_VAL]], [[ARG_IDX2_VAL]]
+; IR-NEXT:    store i32 [[ADD_I]], ptr [[ARG_IDX1_VAL]], align 4
+; IR-NEXT:    [[TMP:%.*]] = call ptr @escape(ptr nonnull [[ARG_IDX2]])
+; IR-NEXT:    ret void
+;
+; PTX-LABEL: load_alignment(
+; PTX:       {
+; PTX-NEXT:    .local .align 8 .b8 __local_depot0[24];
+; PTX-NEXT:    .reg .b64 %SP;
+; PTX-NEXT:    .reg .b64 %SPL;
+; PTX-NEXT:    .reg .b32 %r<4>;
+; PTX-NEXT:    .reg .b64 %rd<10>;
+; PTX-EMPTY:
+; PTX-NEXT:  // %bb.0: // %entry
+; PTX-NEXT:    mov.u64 %SPL, __local_depot0;
+; PTX-NEXT:    add.u64 %rd2, %SPL, 0;
+; PTX-NEXT:    ld.param.u64 %rd3, [load_alignment_param_0+16];
+; PTX-NEXT:    st.local.u64 [%rd2+16], %rd3;
+; PTX-NEXT:    ld.param.u64 %rd4, [load_alignment_param_0+8];
+; PTX-NEXT:    st.local.u64 [%rd2+8], %rd4;
+; PTX-NEXT:    ld.param.u64 %rd5, [load_alignment_param_0];
+; PTX-NEXT:    st.local.u64 [%rd2], %rd5;
+; PTX-NEXT:    add.s64 %rd6, %rd2, 16;
+; PTX-NEXT:    cvta.local.u64 %rd7, %rd6;
+; PTX-NEXT:    cvt.u32.u64 %r1, %rd3;
+; PTX-NEXT:    ld.u32 %r2, [%rd5];
+; PTX-NEXT:    add.s32 %r3, %r2, %r1;
+; PTX-NEXT:    st.u32 [%rd4], %r3;
+; PTX-NEXT:    { // callseq 0, 0
+; PTX-NEXT:    .param .b64 param0;
+; PTX-NEXT:    st.param.b64 [param0], %rd7;
+; PTX-NEXT:    .param .b64 retval0;
+; PTX-NEXT:    call.uni (retval0),
+; PTX-NEXT:    escape,
+; PTX-NEXT:    (
+; PTX-NEXT:    param0
+; PTX-NEXT:    );
+; PTX-NEXT:    ld.param.b64 %rd8, [retval0];
+; PTX-NEXT:    } // callseq 0
+; PTX-NEXT:    ret;
 entry:
-; IR: call void @llvm.memcpy.p0.p101.i64(ptr align 8
-; PTX: ld.param.u64
-; PTX-NOT: ld.param.u8
   %arg.idx.val = load ptr, ptr %arg, align 8
   %arg.idx1 = getelementptr %class.outer, ptr %arg, i64 0, i32 0, i32 1
   %arg.idx1.val = load ptr, ptr %arg.idx1, align 8
@@ -34,8 +83,16 @@ entry:
 }
 
 ; Check that nvptx-lower-args copies padding as the struct may have been a union
-; COMMON-LABEL: load_padding
 define void @load_padding(ptr nocapture readonly byval(%class.padded) %arg) {
+; IR-LABEL: define void @load_padding(
+; IR-SAME: ptr readonly byval([[CLASS_PADDED:%.*]]) captures(none) [[ARG:%.*]]) {
+; IR-NEXT:    [[ARG1:%.*]] = alloca [[CLASS_PADDED]], align 8
+; IR-NEXT:    [[ARG2:%.*]] = addrspacecast ptr [[ARG]] to ptr addrspace(101)
+; IR-NEXT:    call void @llvm.memcpy.p0.p101.i64(ptr align 8 [[ARG1]], ptr addrspace(101) align 8 [[ARG2]], i64 8, i1 false)
+; IR-NEXT:    [[TMP:%.*]] = call ptr @escape(ptr nonnull align 16 [[ARG1]])
+; IR-NEXT:    ret void
+;
+; PTX-LABEL: load_padding(
 ; PTX:       {
 ; PTX-NEXT:    .local .align 8 .b8 __local_depot1[8];
 ; PTX-NEXT:    .reg .b64 %SP;
@@ -45,8 +102,8 @@ define void @load_padding(ptr nocapture readonly byval(%class.padded) %arg) {
 ; PTX-NEXT:  // %bb.0:
 ; PTX-NEXT:    mov.u64 %SPL, __local_depot1;
 ; PTX-NEXT:    cvta.local.u64 %SP, %SPL;
-; PTX-NEXT:    add.u64         %rd1, %SP, 0;
-; PTX-NEXT:    add.u64         %rd2, %SPL, 0;
+; PTX-NEXT:    add.u64 %rd1, %SP, 0;
+; PTX-NEXT:    add.u64 %rd2, %SPL, 0;
 ; PTX-NEXT:    ld.param.u64 %rd3, [load_padding_param_0];
 ; PTX-NEXT:    st.local.u64 [%rd2], %rd3;
 ; PTX-NEXT:    { // callseq 1, 0
@@ -65,85 +122,168 @@ define void @load_padding(ptr nocapture readonly byval(%class.padded) %arg) {
   ret void
 }
 
-; COMMON-LABEL: ptr_generic
-define ptx_kernel void @ptr_generic(ptr %out, ptr %in) {
-; IRC:  %in3 = addrspacecast ptr %in to ptr addrspace(1)
-; IRC:  %in4 = addrspacecast ptr addrspace(1) %in3 to ptr
-; IRC:  %out1 = addrspacecast ptr %out to ptr addrspace(1)
-; IRC:  %out2 = addrspacecast ptr addrspace(1) %out1 to ptr
-; PTXC: cvta.to.global.u64
-; PTXC: cvta.to.global.u64
-; PTXC: ld.global.u32
-; PTXC: st.global.u32
-
 ; OpenCL can't make assumptions about incoming pointer, so we should generate
 ; generic pointers load/store.
-; IRO-NOT: addrspacecast
-; PTXO-NOT: cvta.to.global
-; PTXO: ld.u32
-; PTXO: st.u32
+define ptx_kernel void @ptr_generic(ptr %out, ptr %in) {
+; IRC-LABEL: define ptx_kernel void @ptr_generic(
+; IRC-SAME: ptr [[OUT:%.*]], ptr [[IN:%.*]]) {
+; IRC-NEXT:    [[IN3:%.*]] = addrspacecast ptr [[IN]] to ptr addrspace(1)
+; IRC-NEXT:    [[IN4:%.*]] = addrspacecast ptr addrspace(1) [[IN3]] to ptr
+; IRC-NEXT:    [[OUT1:%.*]] = addrspacecast ptr [[OUT]] to ptr addrspace(1)
+; IRC-NEXT:    [[OUT2:%.*]] = addrspacecast ptr addrspace(1) [[OUT1]] to ptr
+; IRC-NEXT:    [[V:%.*]] = load i32, ptr [[IN4]], align 4
+; IRC-NEXT:    store i32 [[V]], ptr [[OUT2]], align 4
+; IRC-NEXT:    ret void
+;
+; IRO-LABEL: define ptx_kernel void @ptr_generic(
+; IRO-SAME: ptr [[OUT:%.*]], ptr [[IN:%.*]]) {
+; IRO-NEXT:    [[V:%.*]] = load i32, ptr [[IN]], align 4
+; IRO-NEXT:    store i32 [[V]], ptr [[OUT]], align 4
+; IRO-NEXT:    ret void
+;
+; PTXC-LABEL: ptr_generic(
+; PTXC:       {
+; PTXC-NEXT:    .reg .b32 %r<2>;
+; PTXC-NEXT:    .reg .b64 %rd<5>;
+; PTXC-EMPTY:
+; PTXC-NEXT:  // %bb.0:
+; PTXC-NEXT:    ld.param.u64 %rd1, [ptr_generic_param_0];
+; PTXC-NEXT:    ld.param.u64 %rd2, [ptr_generic_param_1];
+; PTXC-NEXT:    cvta.to.global.u64 %rd3, %rd2;
+; PTXC-NEXT:    cvta.to.global.u64 %rd4, %rd1;
+; PTXC-NEXT:    ld.global.u32 %r1, [%rd3];
+; PTXC-NEXT:    st.global.u32 [%rd4], %r1;
+; PTXC-NEXT:    ret;
+;
+; PTXO-LABEL: ptr_generic(
+; PTXO:       {
+; PTXO-NEXT:    .reg .b32 %r<2>;
+; PTXO-NEXT:    .reg .b64 %rd<3>;
+; PTXO-EMPTY:
+; PTXO-NEXT:  // %bb.0:
+; PTXO-NEXT:    ld.param.u64 %rd1, [ptr_generic_param_0];
+; PTXO-NEXT:    ld.param.u64 %rd2, [ptr_generic_param_1];
+; PTXO-NEXT:    ld.u32 %r1, [%rd2];
+; PTXO-NEXT:    st.u32 [%rd1], %r1;
+; PTXO-NEXT:    ret;
   %v = load i32, ptr  %in, align 4
   store i32 %v, ptr %out, align 4
   ret void
 }
 
-; COMMON-LABEL: ptr_nongeneric
 define ptx_kernel void @ptr_nongeneric(ptr addrspace(1) %out, ptr addrspace(3) %in) {
-; IR-NOT: addrspacecast
-; PTX-NOT: cvta.to.global
-; PTX:  ld.shared.u32
-; PTX   st.global.u32
+; IR-LABEL: define ptx_kernel void @ptr_nongeneric(
+; IR-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(3) [[IN:%.*]]) {
+; IR-NEXT:    [[V:%.*]] = load i32, ptr addrspace(3) [[IN]], align 4
+; IR-NEXT:    store i32 [[V]], ptr addrspace(1) [[OUT]], align 4
+; IR-NEXT:    ret void
+;
+; PTX-LABEL: ptr_nongeneric(
+; PTX:       {
+; PTX-NEXT:    .reg .b32 %r<2>;
+; PTX-NEXT:    .reg .b64 %rd<3>;
+; PTX-EMPTY:
+; PTX-NEXT:  // %bb.0:
+; PTX-NEXT:    ld.param.u64 %rd1, [ptr_nongeneric_param_0];
+; PTX-NEXT:    ld.param.u64 %rd2, [ptr_nongeneric_param_1];
+; PTX-NEXT:    ld.shared.u32 %r1, [%rd2];
+; PTX-NEXT:    st.global.u32 [%rd1], %r1;
+; PTX-NEXT:    ret;
   %v = load i32, ptr addrspace(3) %in, align 4
   store i32 %v, ptr addrspace(1) %out, align 4
   ret void
 }
 
-; COMMON-LABEL: ptr_as_int
- define ptx_kernel void @ptr_as_int(i64 noundef %i, i32 noundef %v) {
-; IR:   [[P:%.*]] = inttoptr i64 %i to ptr
-; IRC:  [[P1:%.*]] = addrspacecast ptr [[P]] to ptr addrspace(1)
-; IRC:  addrspacecast ptr addrspace(1) [[P1]] to ptr
-; IRO-NOT: addrspacecast
-
-; PTXC-DAG:  ld.param.u64    [[I:%rd.*]], [ptr_as_int_param_0];
-; PTXC-DAG:  ld.param.u32    [[V:%r.*]], [ptr_as_int_param_1];
-; PTXC:      cvta.to.global.u64 %[[P:rd.*]], [[I]];
-; PTXC:      st.global.u32    [%[[P]]], [[V]];
-
-; PTXO-DAG:  ld.param.u64    %[[P:rd.*]], [ptr_as_int_param_0];
-; PTXO-DAG:  ld.param.u32    [[V:%r.*]], [ptr_as_int_param_1];
-; PTXO:      st.u32   [%[[P]]], [[V]];
-
+define ptx_kernel void @ptr_as_int(i64 noundef %i, i32 noundef %v) {
+; IRC-LABEL: define ptx_kernel void @ptr_as_int(
+; IRC-SAME: i64 noundef [[I:%.*]], i32 noundef [[V:%.*]]) {
+; IRC-NEXT:    [[P:%.*]] = inttoptr i64 [[I]] to ptr
+; IRC-NEXT:    [[P1:%.*]] = addrspacecast ptr [[P]] to ptr addrspace(1)
+; IRC-NEXT:    [[P2:%.*]] = addrspacecast ptr addrspace(1) [[P1]] to ptr
+; IRC-NEXT:    store i32 [[V]], ptr [[P2]], align 4
+; IRC-NEXT:    ret void
+;
+; IRO-LABEL: define ptx_kernel void @ptr_as_int(
+; IRO-SAME: i64 noundef [[I:%.*]], i32 noundef [[V:%.*]]) {
+; IRO-NEXT:    [[P:%.*]] = inttoptr i64 [[I]] to ptr
+; IRO-NEXT:    store i32 [[V]], ptr [[P]], align 4
+; IRO-NEXT:    ret void
+;
+; PTXC-LABEL: ptr_as_int(
+; PTXC:       {
+; PTXC-NEXT:    .reg .b32 %r<2>;
+; PTXC-NEXT:    .reg .b64 %rd<3>;
+; PTXC-EMPTY:
+; PTXC-NEXT:  // %bb.0:
+; PTXC-NEXT:    ld.param.u64 %rd1, [ptr_as_int_param_0];
+; PTXC-NEXT:    ld.param.u32 %r1, [ptr_as_int_param_1];
+; PTXC-NEXT:    cvta.to.global.u64 %rd2, %rd1;
+; PTXC-NEXT:    st.global.u32 [%rd2], %r1;
+; PTXC-NEXT:    ret;
+;
+; PTXO-LABEL: ptr_as_int(
+; PTXO:       {
+; PTXO-NEXT:    .reg .b32 %r<2>;
+; PTXO-NEXT:    .reg .b64 %rd<2>;
+; PTXO-EMPTY:
+; PTXO-NEXT:  // %bb.0:
+; PTXO-NEXT:    ld.param.u64 %rd1, [ptr_as_int_param_0];
+; PTXO-NEXT:    ld.param.u32 %r1, [ptr_as_int_param_1];
+; PTXO-NEXT:    st.u32 [%rd1], %r1;
+; PTXO-NEXT:    ret;
   %p = inttoptr i64 %i to ptr
   store i32 %v, ptr %p, align 4
   ret void
 }
 
 %struct.S = type { i64 }
 
-; COMMON-LABEL: ptr_as_int_aggr
 define ptx_kernel void @ptr_as_int_aggr(ptr nocapture noundef readonly byval(%struct.S) align 8 %s, i32 noundef %v) {
-; IR:   [[S:%.*]] = addrspacecast ptr %s to ptr addrspace(101)
-; IR:   [[I:%.*]] = load i64, ptr addrspace(101) [[S]], align 8
-; IR:   [[P0:%.*]] = inttoptr i64 [[I]] to ptr
-; IRC:  [[P1:%.*]] = addrspacecast ptr [[P]] to ptr addrspace(1)
-; IRC:  [[P:%.*]] = addrspacecast ptr addrspace(1) [[P1]] to ptr
-; IRO-NOT: addrspacecast
-
-; PTXC-DAG:  ld.param.u64    [[I:%rd.*]], [ptr_as_int_aggr_param_0];
-; PTXC-DAG:  ld.param.u32    [[V:%r.*]], [ptr_as_int_aggr_param_1];
-; PTXC:      cvta.to.global.u64 %[[P:rd.*]], [[I]];
-; PTXC:      st.global.u32    [%[[P]]], [[V]];
-
-; PTXO-DAG:  ld.param.u64    %[[P:rd.*]], [ptr_as_int_aggr_param_0];
-; PTXO-DAG:  ld.param.u32    [[V:%r.*]], [ptr_as_int_aggr_param_1];
-; PTXO:      st.u32   [%[[P]]], [[V]];
+; IRC-LABEL: define ptx_kernel void @ptr_as_int_aggr(
+; IRC-SAME: ptr noundef readonly byval([[STRUCT_S:%.*]]) align 8 captures(none) [[S:%.*]], i32 noundef [[V:%.*]]) {
+; IRC-NEXT:    [[S3:%.*]] = addrspacecast ptr [[S]] to ptr addrspace(101)
+; IRC-NEXT:    [[I:%.*]] = load i64, ptr addrspace(101) [[S3]], align 8
+; IRC-NEXT:    [[P:%.*]] = inttoptr i64 [[I]] to ptr
+; IRC-NEXT:    [[P1:%.*]] = addrspacecast ptr [[P]] to ptr addrspace(1)
+; IRC-NEXT:    [[P2:%.*]] = addrspacecast ptr addrspace(1) [[P1]] to ptr
+; IRC-NEXT:    store i32 [[V]], ptr [[P2]], align 4
+; IRC-NEXT:    ret void
+;
+; IRO-LABEL: define ptx_kernel void @ptr_as_int_aggr(
+; IRO-SAME: ptr noundef readonly byval([[STRUCT_S:%.*]]) align 8 captures(none) [[S:%.*]], i32 noundef [[V:%.*]]) {
+; IRO-NEXT:    [[S1:%.*]] = addrspacecast ptr [[S]] to ptr addrspace(101)
+; IRO-NEXT:    [[I:%.*]] = load i64, ptr addrspace(101) [[S1]], align 8
+; IRO-NEXT:    [[P:%.*]] = inttoptr i64 [[I]] to ptr
+; IRO-NEXT:    store i32 [[V]], ptr [[P]], align 4
+; IRO-NEXT:    ret void
+;
+; PTXC-LABEL: ptr_as_int_aggr(
+; PTXC:       {
+; PTXC-NEXT:    .reg .b32 %r<2>;
+; PTXC-NEXT:    .reg .b64 %rd<3>;
+; PTXC-EMPTY:
+; PTXC-NEXT:  // %bb.0:
+; PTXC-NEXT:    ld.param.u32 %r1, [ptr_as_int_aggr_param_1];
+; PTXC-NEXT:    ld.param.u64 %rd1, [ptr_as_int_aggr_param_0];
+; PTXC-NEXT:    cvta.to.global.u64 %rd2, %rd1;
+; PTXC-NEXT:    st.global.u32 [%rd2], %r1;
+; PTXC-NEXT:    ret;
+;
+; PTXO-LABEL: ptr_as_int_aggr(
+; PTXO:       {
+; PTXO-NEXT:    .reg .b32 %r<2>;
+; PTXO-NEXT:    .reg .b64 %rd<2>;
+; PTXO-EMPTY:
+; PTXO-NEXT:  // %bb.0:
+; PTXO-NEXT:    ld.param.u32 %r1, [ptr_as_int_aggr_param_1];
+; PTXO-NEXT:    ld.param.u64 %rd1, [ptr_as_int_aggr_param_0];
+; PTXO-NEXT:    st.u32 [%rd1], %r1;
+; PTXO-NEXT:    ret;
   %i = load i64, ptr %s, align 8
   %p = inttoptr i64 %i to ptr
   store i32 %v, ptr %p, align 4
   ret void
 }
 
-
 ; Function Attrs: convergent nounwind
 declare dso_local ptr @escape(ptr) local_unnamed_addr