diff options
Diffstat (limited to 'llvm/lib/Target/NVPTX/NVPTXIntrinsics.td')
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 12 |
1 files changed, 12 insertions, 0 deletions
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index 66419f034f6..31bed350f38 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -1937,6 +1937,12 @@ multiclass NG_TO_G<string Str, Intrinsic Intrin> { def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src), !strconcat("cvta.", Str, ".u64 \t$result, $src;"), [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>; + def _yes_6432 : NVPTXInst<(outs Int64Regs:$result), (ins Int32Regs:$src), + "{{ .reg .b64 %tmp;\n\t" + #" cvt.u64.u32 \t%tmp, $src;\n\t" + #" cvta." # Str # ".u64 \t$result, %tmp; }}", + [(set Int64Regs:$result, (Intrin Int32Regs:$src))]>, + Requires<[useShortPtr]>; } multiclass G_TO_NG<string Str, Intrinsic Intrin> { @@ -1946,6 +1952,12 @@ multiclass G_TO_NG<string Str, Intrinsic Intrin> { def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src), !strconcat("cvta.to.", Str, ".u64 \t$result, $src;"), [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>; + def _yes_3264 : NVPTXInst<(outs Int32Regs:$result), (ins Int64Regs:$src), + "{{ .reg .b64 %tmp;\n\t" + #" cvta.to." # Str # ".u64 \t%tmp, $src;\n\t" + #" cvt.u32.u64 \t$result, %tmp; }}", + [(set Int32Regs:$result, (Intrin Int64Regs:$src))]>, + Requires<[useShortPtr]>; } defm cvta_local : NG_TO_G<"local", int_nvvm_ptr_local_to_gen>; |
