[PATCH] D12093: [NVPTX] Support register copy from i16 to i32 register types

[Samuel Antao via llvm-commits]

sfantao created this revision.
sfantao added reviewers: jholewinski, jingyue.
sfantao added a subscriber: llvm-commits.
Herald added a subscriber: jholewinski.

In some cases register copies between i16 and i32 are created, but there is no physical register copy implementation for them. This usually happens each time a non-coherent load is emitted for 1- or 2-byte int types. This patch adds support for that.

It is possible this would have to be fixed before the copies are emitted, given that we get things like:

ld.global.nc.u8 %rs7, [%rd81]
cvt.u32.u16 %r56, %rs7;

which means we have an implicit cast from u8 to u16, which suggest that something may be wrong. Nevertheless, for the testcases I have access to, this fix works just fine.



http://reviews.llvm.org/D12093

Files:
  lib/Target/NVPTX/NVPTXInstrInfo.cpp
  test/CodeGen/NVPTX/reg-copy-int.ll

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D12093.32353.patch
Type: text/x-patch
Size: 19201 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150817/a096fe8b/attachment.bin>