29 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --function foo --extra_scrub --default-march nvptx64 --filter-out ".*//.*" --filter-out "[\{\}\(\)]" --version 52 3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 | FileCheck %s4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 | %ptxas-verify %}5 6target triple = "nvptx-nvidia-cuda"7 8define ptx_kernel void @foo(ptr noalias readonly %ptr, ptr noalias %retval) {9; CHECK-LABEL: foo(10; CHECK: .reg .b32 %r<4>;11; CHECK: .reg .b64 %rd<5>;12; CHECK-EMPTY:13; CHECK: ld.param.b64 %rd1, [foo_param_0];14; CHECK: cvta.to.global.u64 %rd2, %rd1;15; CHECK: ld.param.b64 %rd3, [foo_param_1];16; CHECK: cvta.to.global.u64 %rd4, %rd3;17; CHECK: ld.global.nc.b8 %r1, [%rd2];18; CHECK: add.s32 %r2, %r1, 1;19; CHECK: and.b32 %r3, %r2, 1;20; CHECK: st.global.b32 [%rd4], %r3;21; CHECK: ret;22 %ld = load i1, ptr %ptr, align 123 %zext = zext i1 %ld to i3224 %add = add i32 %zext, 125 %and = and i32 %add, 126 store i32 %and, ptr %retval27 ret void28}29