Fix incomplete object size (=4n + 3) support of amd_wave_read_first_lane()

[poyenc]

This is a complement to PR 711

Though it's unlikely to use (4n + 3) byte class types in real-world scenarios. I think it's still necessary to fix missing size support of ck::amd_wave_read_first_lane().

We can create such type by declaring a struct S which has only single byte-array data member:

struct S {
   unsigned char uca[3];
};
static_assert(sizeof(S) == 3);

const auto s = ck::amd_wave_read_first_lane(S{1, 2, 3}); // compilation error before this PR
assert(s.uca[0] == 1);
assert(s.uca[1] == 2);
assert(s.uca[2] == 3);

[poyenc]

Fixed the compilation error (using __device__ function from host method)