Use int for int8x4 due to performance overhead of char4

[vinx13]

Loading four int8 elements as char4 is likely to produce more integer instructions. When we use int8 intrinsics (e.g. dp4a), we need packed 32-bit data, which need extra operations for packing int8 elements.

For example, below is a ptx code snippet of
__dp4a((( char4*)(( signed char*)A_shared_local + ((k_inner_outer_outer % 2) * 32)))[0], (( char4*)(( signed char*)B_shared_local + ((k_inner_outer_outer % 2) * 32)))[0], C_local[0]);

ld.shared.v4.u8 {%rs577, %rs578, %rs579, %rs580}, [%r5+24];
ld.shared.v4.u8 {%rs625, %rs626, %rs627, %rs628}, [%r6+24];
...
cvt.u32.u16 %r2873, %rs580;
mul.wide.u16 %r2874, %rs578, 256;
cvt.u32.u16 %r2875, %rs577;
cvt.u32.u16 %r2876, %rs579;
prmt.b32 %r2877, %r2876, %r2875, 28756;
prmt.b32 %r2878, %r2873, %r2877, 1620;
or.b32 %r2010, %r2878, %r2874;

cvt.u32.u16 %r2879, %rs628;
mul.wide.u16 %r2880, %rs626, 256;
cvt.u32.u16 %r2881, %rs625;
cvt.u32.u16 %r2882, %rs627;
prmt.b32 %r2883, %r2882, %r2881, 28756;
prmt.b32 %r2884, %r2879, %r2883, 1620;
or.b32 %r1815, %r2884, %r2880;

dp4a.s32.s32 %r1785, %r2010, %r1815, %r1529;

We would like to use ld.shared.u32 in this case so that 32-bit data can be directly loaded.

This disables support for vectorized int8 arithmetic operations. Since these operations are used in few cases, we prefer better performance here.

[vinx13]

@tqchen Please review.

[tqchen]

cc @nishi-t

[tqchen]

Please add a comment block here on why are we making this choice, so people won't change it back

[nishi-t]

vectorized_add test will not work for int8 anymore. Please remove this:

check_cuda("int8", 64, 4)

https://github.com/dmlc/tvm/blob/master/tests/python/unittest/test_codegen_cuda.py#L34

[vinx13]

@tqchen Please review again.