Skip to content

Commit 1c22afd

Browse files
committed
handle fdiv and other instructions where v2f32 is illegal
Requires us to lower EXTRACT_VECTOR_ELT as well.
1 parent 7de3428 commit 1c22afd

File tree

2 files changed

+15
-0
lines changed

2 files changed

+15
-0
lines changed

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -928,6 +928,7 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
928928
{ISD::FDIV, ISD::FREM, ISD::FSQRT, ISD::FSIN, ISD::FCOS}) {
929929
setOperationAction(Op, MVT::f16, Promote);
930930
setOperationAction(Op, MVT::f32, Legal);
931+
setOperationAction(Op, MVT::v2f32, Expand);
931932
setOperationAction(Op, MVT::f64, Legal);
932933
setOperationAction(Op, MVT::v2f16, Expand);
933934
setOperationAction(Op, MVT::v2bf16, Expand);

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3228,6 +3228,14 @@ let hasSideEffects = false in {
32283228
(ins Int64Regs:$s),
32293229
"{{ .reg .b32 tmp; mov.b64 {$low, tmp}, $s; }}",
32303230
[]>;
3231+
def I64toF32H : NVPTXInst<(outs Float32Regs:$high),
3232+
(ins Int64Regs:$s),
3233+
"{{ .reg .b32 tmp; mov.b64 {tmp, $high}, $s; }}",
3234+
[]>;
3235+
def I64toF32L : NVPTXInst<(outs Float32Regs:$low),
3236+
(ins Int64Regs:$s),
3237+
"{{ .reg .b32 tmp; mov.b64 {$low, tmp}, $s; }}",
3238+
[]>;
32313239

32323240
}
32333241

@@ -3251,6 +3259,12 @@ def : Pat<(extractelt vt:$src, 0),
32513259
def : Pat<(extractelt vt:$src, 1),
32523260
(I32toI16H $src)>;
32533261
}
3262+
3263+
def : Pat<(extractelt v2f32:$src, 0),
3264+
(I64toF32L $src)>;
3265+
def : Pat<(extractelt v2f32:$src, 1),
3266+
(I64toF32H $src)>;
3267+
32543268
def : Pat<(v2f16 (build_vector f16:$a, f16:$b)),
32553269
(V2I16toI32 $a, $b)>;
32563270
def : Pat<(v2bf16 (build_vector bf16:$a, bf16:$b)),

0 commit comments

Comments
 (0)