derek.holloway
7502bf64a4
CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
197 lines
6.4 KiB
LLVM
197 lines
6.4 KiB
LLVM
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
|
|
; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX32
|
|
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX64
|
|
; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
|
|
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
|
|
|
|
; Ensure we access the local stack properly
|
|
|
|
define void @foo(i32 %a) {
|
|
; PTX32-LABEL: foo(
|
|
; PTX32: {
|
|
; PTX32-NEXT: .local .align 4 .b8 __local_depot0[4];
|
|
; PTX32-NEXT: .reg .b32 %SP;
|
|
; PTX32-NEXT: .reg .b32 %SPL;
|
|
; PTX32-NEXT: .reg .b32 %r<3>;
|
|
; PTX32-EMPTY:
|
|
; PTX32-NEXT: // %bb.0:
|
|
; PTX32-NEXT: mov.b32 %SPL, __local_depot0;
|
|
; PTX32-NEXT: ld.param.b32 %r1, [foo_param_0];
|
|
; PTX32-NEXT: add.u32 %r2, %SPL, 0;
|
|
; PTX32-NEXT: st.local.b32 [%r2], %r1;
|
|
; PTX32-NEXT: ret;
|
|
;
|
|
; PTX64-LABEL: foo(
|
|
; PTX64: {
|
|
; PTX64-NEXT: .local .align 4 .b8 __local_depot0[4];
|
|
; PTX64-NEXT: .reg .b64 %SP;
|
|
; PTX64-NEXT: .reg .b64 %SPL;
|
|
; PTX64-NEXT: .reg .b32 %r<2>;
|
|
; PTX64-NEXT: .reg .b64 %rd<2>;
|
|
; PTX64-EMPTY:
|
|
; PTX64-NEXT: // %bb.0:
|
|
; PTX64-NEXT: mov.b64 %SPL, __local_depot0;
|
|
; PTX64-NEXT: ld.param.b32 %r1, [foo_param_0];
|
|
; PTX64-NEXT: add.u64 %rd1, %SPL, 0;
|
|
; PTX64-NEXT: st.local.b32 [%rd1], %r1;
|
|
; PTX64-NEXT: ret;
|
|
%local = alloca i32, align 4
|
|
store volatile i32 %a, ptr %local
|
|
ret void
|
|
}
|
|
|
|
define ptx_kernel void @foo2(i32 %a) {
|
|
; PTX32-LABEL: foo2(
|
|
; PTX32: {
|
|
; PTX32-NEXT: .local .align 4 .b8 __local_depot1[4];
|
|
; PTX32-NEXT: .reg .b32 %SP;
|
|
; PTX32-NEXT: .reg .b32 %SPL;
|
|
; PTX32-NEXT: .reg .b32 %r<4>;
|
|
; PTX32-EMPTY:
|
|
; PTX32-NEXT: // %bb.0:
|
|
; PTX32-NEXT: mov.b32 %SPL, __local_depot1;
|
|
; PTX32-NEXT: cvta.local.u32 %SP, %SPL;
|
|
; PTX32-NEXT: ld.param.b32 %r1, [foo2_param_0];
|
|
; PTX32-NEXT: add.u32 %r2, %SP, 0;
|
|
; PTX32-NEXT: add.u32 %r3, %SPL, 0;
|
|
; PTX32-NEXT: st.local.b32 [%r3], %r1;
|
|
; PTX32-NEXT: { // callseq 0, 0
|
|
; PTX32-NEXT: .param .b32 param0;
|
|
; PTX32-NEXT: st.param.b32 [param0], %r2;
|
|
; PTX32-NEXT: call.uni bar, (param0);
|
|
; PTX32-NEXT: } // callseq 0
|
|
; PTX32-NEXT: ret;
|
|
;
|
|
; PTX64-LABEL: foo2(
|
|
; PTX64: {
|
|
; PTX64-NEXT: .local .align 4 .b8 __local_depot1[4];
|
|
; PTX64-NEXT: .reg .b64 %SP;
|
|
; PTX64-NEXT: .reg .b64 %SPL;
|
|
; PTX64-NEXT: .reg .b32 %r<2>;
|
|
; PTX64-NEXT: .reg .b64 %rd<3>;
|
|
; PTX64-EMPTY:
|
|
; PTX64-NEXT: // %bb.0:
|
|
; PTX64-NEXT: mov.b64 %SPL, __local_depot1;
|
|
; PTX64-NEXT: cvta.local.u64 %SP, %SPL;
|
|
; PTX64-NEXT: ld.param.b32 %r1, [foo2_param_0];
|
|
; PTX64-NEXT: add.u64 %rd1, %SP, 0;
|
|
; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
|
|
; PTX64-NEXT: st.local.b32 [%rd2], %r1;
|
|
; PTX64-NEXT: { // callseq 0, 0
|
|
; PTX64-NEXT: .param .b64 param0;
|
|
; PTX64-NEXT: st.param.b64 [param0], %rd1;
|
|
; PTX64-NEXT: call.uni bar, (param0);
|
|
; PTX64-NEXT: } // callseq 0
|
|
; PTX64-NEXT: ret;
|
|
%local = alloca i32, align 4
|
|
store i32 %a, ptr %local
|
|
call void @bar(ptr %local)
|
|
ret void
|
|
}
|
|
|
|
declare void @bar(ptr %a)
|
|
|
|
define void @foo3(i32 %a) {
|
|
; PTX32-LABEL: foo3(
|
|
; PTX32: {
|
|
; PTX32-NEXT: .local .align 4 .b8 __local_depot2[12];
|
|
; PTX32-NEXT: .reg .b32 %SP;
|
|
; PTX32-NEXT: .reg .b32 %SPL;
|
|
; PTX32-NEXT: .reg .b32 %r<5>;
|
|
; PTX32-EMPTY:
|
|
; PTX32-NEXT: // %bb.0:
|
|
; PTX32-NEXT: mov.b32 %SPL, __local_depot2;
|
|
; PTX32-NEXT: ld.param.b32 %r1, [foo3_param_0];
|
|
; PTX32-NEXT: add.u32 %r2, %SPL, 0;
|
|
; PTX32-NEXT: shl.b32 %r3, %r1, 2;
|
|
; PTX32-NEXT: add.s32 %r4, %r2, %r3;
|
|
; PTX32-NEXT: st.local.b32 [%r4], %r1;
|
|
; PTX32-NEXT: ret;
|
|
;
|
|
; PTX64-LABEL: foo3(
|
|
; PTX64: {
|
|
; PTX64-NEXT: .local .align 4 .b8 __local_depot2[12];
|
|
; PTX64-NEXT: .reg .b64 %SP;
|
|
; PTX64-NEXT: .reg .b64 %SPL;
|
|
; PTX64-NEXT: .reg .b32 %r<2>;
|
|
; PTX64-NEXT: .reg .b64 %rd<4>;
|
|
; PTX64-EMPTY:
|
|
; PTX64-NEXT: // %bb.0:
|
|
; PTX64-NEXT: mov.b64 %SPL, __local_depot2;
|
|
; PTX64-NEXT: ld.param.b32 %r1, [foo3_param_0];
|
|
; PTX64-NEXT: add.u64 %rd1, %SPL, 0;
|
|
; PTX64-NEXT: mul.wide.s32 %rd2, %r1, 4;
|
|
; PTX64-NEXT: add.s64 %rd3, %rd1, %rd2;
|
|
; PTX64-NEXT: st.local.b32 [%rd3], %r1;
|
|
; PTX64-NEXT: ret;
|
|
%local = alloca [3 x i32], align 4
|
|
%1 = getelementptr inbounds i32, ptr %local, i32 %a
|
|
store i32 %a, ptr %1
|
|
ret void
|
|
}
|
|
|
|
define void @foo4() {
|
|
; PTX32-LABEL: foo4(
|
|
; PTX32: {
|
|
; PTX32-NEXT: .local .align 4 .b8 __local_depot3[8];
|
|
; PTX32-NEXT: .reg .b32 %SP;
|
|
; PTX32-NEXT: .reg .b32 %SPL;
|
|
; PTX32-NEXT: .reg .b32 %r<5>;
|
|
; PTX32-EMPTY:
|
|
; PTX32-NEXT: // %bb.0:
|
|
; PTX32-NEXT: mov.b32 %SPL, __local_depot3;
|
|
; PTX32-NEXT: cvta.local.u32 %SP, %SPL;
|
|
; PTX32-NEXT: add.u32 %r1, %SP, 0;
|
|
; PTX32-NEXT: add.u32 %r2, %SPL, 0;
|
|
; PTX32-NEXT: add.u32 %r3, %SP, 4;
|
|
; PTX32-NEXT: add.u32 %r4, %SPL, 4;
|
|
; PTX32-NEXT: st.local.b32 [%r2], 0;
|
|
; PTX32-NEXT: st.local.b32 [%r4], 0;
|
|
; PTX32-NEXT: { // callseq 1, 0
|
|
; PTX32-NEXT: .param .b32 param0;
|
|
; PTX32-NEXT: st.param.b32 [param0], %r1;
|
|
; PTX32-NEXT: call.uni bar, (param0);
|
|
; PTX32-NEXT: } // callseq 1
|
|
; PTX32-NEXT: { // callseq 2, 0
|
|
; PTX32-NEXT: .param .b32 param0;
|
|
; PTX32-NEXT: st.param.b32 [param0], %r3;
|
|
; PTX32-NEXT: call.uni bar, (param0);
|
|
; PTX32-NEXT: } // callseq 2
|
|
; PTX32-NEXT: ret;
|
|
;
|
|
; PTX64-LABEL: foo4(
|
|
; PTX64: {
|
|
; PTX64-NEXT: .local .align 4 .b8 __local_depot3[8];
|
|
; PTX64-NEXT: .reg .b64 %SP;
|
|
; PTX64-NEXT: .reg .b64 %SPL;
|
|
; PTX64-NEXT: .reg .b64 %rd<5>;
|
|
; PTX64-EMPTY:
|
|
; PTX64-NEXT: // %bb.0:
|
|
; PTX64-NEXT: mov.b64 %SPL, __local_depot3;
|
|
; PTX64-NEXT: cvta.local.u64 %SP, %SPL;
|
|
; PTX64-NEXT: add.u64 %rd1, %SP, 0;
|
|
; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
|
|
; PTX64-NEXT: add.u64 %rd3, %SP, 4;
|
|
; PTX64-NEXT: add.u64 %rd4, %SPL, 4;
|
|
; PTX64-NEXT: st.local.b32 [%rd2], 0;
|
|
; PTX64-NEXT: st.local.b32 [%rd4], 0;
|
|
; PTX64-NEXT: { // callseq 1, 0
|
|
; PTX64-NEXT: .param .b64 param0;
|
|
; PTX64-NEXT: st.param.b64 [param0], %rd1;
|
|
; PTX64-NEXT: call.uni bar, (param0);
|
|
; PTX64-NEXT: } // callseq 1
|
|
; PTX64-NEXT: { // callseq 2, 0
|
|
; PTX64-NEXT: .param .b64 param0;
|
|
; PTX64-NEXT: st.param.b64 [param0], %rd3;
|
|
; PTX64-NEXT: call.uni bar, (param0);
|
|
; PTX64-NEXT: } // callseq 2
|
|
; PTX64-NEXT: ret;
|
|
%A = alloca i32
|
|
%B = alloca i32
|
|
store i32 0, ptr %A
|
|
store i32 0, ptr %B
|
|
call void @bar(ptr %A)
|
|
call void @bar(ptr %B)
|
|
ret void
|
|
}
|