On 11/04/2014 05:51 PM, Bernd Schmidt wrote:
On 11/04/2014 05:48 PM, Richard Henderson wrote:
On 10/28/2014 03:56 PM, Bernd Schmidt wrote:
+nvptx_ptx_type_from_mode (enum machine_mode mode, bool promote)
+{
+ switch (mode)
+ {
+ case BLKmode:
+ return ".b8";
+ case BImode:
+ return ".pred";
+ case QImode:
+ if (promote)
+ return ".u32";
+ else
+ return ".u8";
+ case HImode:
+ return ".u16";
Promote here too? Or does this have nothing to do with
+static enum machine_mode
+arg_promotion (enum machine_mode mode)
+{
+ if (mode == QImode || mode == HImode)
+ return SImode;
+ return mode;
+}
No, these are different problems - the one in arg promotion is purely
about K&R C and trying to match untyped function decls with calls, while
the type_from_mode bit was about some ptx ideosyncracy. Although I
forget what the problem was, that code is more than a year old - I'll
see if I can get rid of this.
Err, no, it's quite necessary. From the manual "The .u8, .s8 and .b8
instruction types are restricted to ld, st and cvt instructions." This
means that if the compiler generates reasonable-looking code along the
lines of
.reg .u8 %r70;
mov.u8 %r70,48;
you get
ptxas 20000211-1.o, line 191; error : Arguments mismatch for
instruction 'mov'
Now, one _could_ write .cvt.u8.u32 for the load immediate, but then one
would also have to write .cvt.u8.u8 for register-register moves, and
that's starting to look iffy. I don't really want to rely on the ptx
assembler to do the right thing for "conversions" from one type to itself.
Bernd
