a9ab8a019e
This change uses the untyped variant of `selp.` in all integer cases to simplify the ISel TableGen logic. It is not expected to have any impact on the final SASS.
65 lines
2.2 KiB
LLVM
65 lines
2.2 KiB
LLVM
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
|
|
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 | FileCheck %s
|
|
; RUN: %if ptxas-12.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 | %ptxas-verify -arch=sm_90 %}
|
|
|
|
target triple = "nvptx64-nvidia-cuda"
|
|
|
|
declare {i32, i1} @llvm.nvvm.elect.sync(i32)
|
|
|
|
define {i32, i1} @elect_sync(i32 %mask) {
|
|
; CHECK-LABEL: elect_sync(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .pred %p<2>;
|
|
; CHECK-NEXT: .reg .b16 %rs<2>;
|
|
; CHECK-NEXT: .reg .b32 %r<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.u32 %r1, [elect_sync_param_0];
|
|
; CHECK-NEXT: elect.sync %r2|%p1, %r1;
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;
|
|
; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;
|
|
; CHECK-NEXT: ret;
|
|
%val = call {i32, i1} @llvm.nvvm.elect.sync(i32 %mask)
|
|
ret {i32, i1} %val
|
|
}
|
|
|
|
define {i32, i1} @elect_sync_imm() {
|
|
; CHECK-LABEL: elect_sync_imm(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .pred %p<2>;
|
|
; CHECK-NEXT: .reg .b16 %rs<2>;
|
|
; CHECK-NEXT: .reg .b32 %r<2>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: elect.sync %r1|%p1, -1;
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
|
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;
|
|
; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;
|
|
; CHECK-NEXT: ret;
|
|
%val = call {i32, i1} @llvm.nvvm.elect.sync(i32 u0xffffffff)
|
|
ret {i32, i1} %val
|
|
}
|
|
|
|
; When there are two elect.sync's make sure that
|
|
; the second one is not optimized away.
|
|
define {i32, i1} @elect_sync_twice(i32 %mask) {
|
|
; CHECK-LABEL: elect_sync_twice(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .pred %p<3>;
|
|
; CHECK-NEXT: .reg .b16 %rs<2>;
|
|
; CHECK-NEXT: .reg .b32 %r<4>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.u32 %r1, [elect_sync_twice_param_0];
|
|
; CHECK-NEXT: elect.sync %r2|%p1, %r1;
|
|
; CHECK-NEXT: elect.sync %r3|%p2, %r1;
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;
|
|
; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;
|
|
; CHECK-NEXT: ret;
|
|
%val = call {i32, i1} @llvm.nvvm.elect.sync(i32 %mask)
|
|
%val2 = call {i32, i1} @llvm.nvvm.elect.sync(i32 %mask)
|
|
ret {i32, i1} %val
|
|
}
|