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

[Samuel Antao via llvm-commits]

sfantao commandeered this revision.
sfantao edited reviewers, added: jingyue; removed: sfantao.
sfantao added a comment.

Thanks for the review. I am sending a new diff with a more concise regression test. See comments inlined.


================
Comment at: lib/Target/NVPTX/NVPTXInstrInfo.cpp:40
@@ -39,2 +39,3 @@
 
-  if (DestRC->getSize() != SrcRC->getSize())
+  if (DestRC->getSize() != SrcRC->getSize()) {
+    // If the sizes differ it may be possible we are copying a i16 to a i32
----------------
jholewinski wrote:
> jingyue wrote:
> > Out of curiosity, which code pattern leads to 16-to-32 copy? Normally, COPY traditionally copies the same type. Later, I added int-to-float and float-to-int because bitcast does that. 
> As far as I know, a register copy from i16 to i32 should not occur.  Trying to emit such a copy is usually an indication of something wrong earlier on.  Can you please explain precisely the case you're hitting?
The machine instructions that are getting produced that lead to that physical register copy are:

```
%vreg159<def> = INT_PTX_LDG_GLOBAL_i8areg64 %vreg204; mem:LD1[%lsr.iv810(addrspace=1)](tbaa=!3) Int16Regs:%vreg159 Int64Regs:%vreg204
%vreg161<def> = COPY %vreg159; Int32Regs:%vreg161 Int16Regs:%vreg159
```
and the IR that generates this is any load of readonly function argument with no aliases.

================
Comment at: lib/Target/NVPTX/NVPTXInstrInfo.cpp:40
@@ -39,2 +39,3 @@
 
-  if (DestRC->getSize() != SrcRC->getSize())
+  if (DestRC->getSize() != SrcRC->getSize()) {
+    // If the sizes differ it may be possible we are copying a i16 to a i32
----------------
sfantao wrote:
> jholewinski wrote:
> > jingyue wrote:
> > > Out of curiosity, which code pattern leads to 16-to-32 copy? Normally, COPY traditionally copies the same type. Later, I added int-to-float and float-to-int because bitcast does that. 
> > As far as I know, a register copy from i16 to i32 should not occur.  Trying to emit such a copy is usually an indication of something wrong earlier on.  Can you please explain precisely the case you're hitting?
> The machine instructions that are getting produced that lead to that physical register copy are:
> 
> ```
> %vreg159<def> = INT_PTX_LDG_GLOBAL_i8areg64 %vreg204; mem:LD1[%lsr.iv810(addrspace=1)](tbaa=!3) Int16Regs:%vreg159 Int64Regs:%vreg204
> %vreg161<def> = COPY %vreg159; Int32Regs:%vreg161 Int16Regs:%vreg159
> ```
> and the IR that generates this is any load of readonly function argument with no aliases.
I traced the problem down to SelectLDGLDU. INT_PTX_LDG_GLOBAL_i8areg64 is generated in there and loads an i8 to an i16 register.

================
Comment at: lib/Target/NVPTX/NVPTXInstrInfo.cpp:51
@@ -41,1 +50,3 @@
+    }
     report_fatal_error("Copy one register into another with a different width");
+  }
----------------
jingyue wrote:
> The error message here needs to be updated. 
Done!

================
Comment at: test/CodeGen/NVPTX/reg-copy-int.ll:14
@@ +13,3 @@
+
+; CHECK-LABEL __omptgt__0_1045bf9_30_(
+; CHECK: ld.global.nc.u8 {{.*}}[[r1:%.+]], [%r{{.+}}];
----------------
jingyue wrote:
> I'd reduce this test case. The current one is too large to demonstrate what you're really testing. 
I'm was very happy with the regression test either... I was having an hard time replicating the problem with a small example and just realized that this issue only happens for readonly function arguments.

The new regression is more concise.


http://reviews.llvm.org/D12093