On Thu, 2017-02-02 at 01:24 -0800, [email protected] wrote: > From: Matt Arsenault <[email protected]> > > v2: Rename test file > --- > .../cl/program/execute/amdgcn.sign_extend_inreg.cl | 387 > +++++++++++++++++++++ > 1 file changed, 387 insertions(+) > create mode 100644 tests/cl/program/execute/amdgcn.sign_extend_inreg.cl > > diff --git a/tests/cl/program/execute/amdgcn.sign_extend_inreg.cl > b/tests/cl/program/execute/amdgcn.sign_extend_inreg.cl > new file mode 100644 > index 000000000..896747f2f > --- /dev/null > +++ b/tests/cl/program/execute/amdgcn.sign_extend_inreg.cl > @@ -0,0 +1,387 @@ > +/*! > +[config] > +name: sign extend in register > +clc_version_min: 10 > + > +dimensions: 1 > + > +## Addition ## > + > +[test] > +name: SALU i8 in i64 0 > +kernel_name: s_sext_in_reg_i8_in_i64 > + > +arg_out: 0 buffer long[14] \ > + 0x0000 \ > + 0x000f \ > + 0xfffffffffffffff0 \ > + 0xffffffffffffffff \ > + 0x0 \ > + 0x0 \ > + 0x79 \ > + 0xffffffffffffff80 \ > + 0xffffffffffffff81 \ > + 0xffffffffffffff82 \ > + 0xfffffffffffffffe \ > + 0x7f \ > + 0xffffffffffffffaa \ > + 0x55 > + > +arg_in: 1 long 0x0000 > +arg_in: 2 long 0x000f > +arg_in: 3 long 0x00f0 > +arg_in: 4 long 0x00ff > +arg_in: 5 long 0xff00 > +arg_in: 6 long 0xf000 > +arg_in: 7 long 0x0079 > +arg_in: 8 long 0x0080 > +arg_in: 9 long 0x0081 > +arg_in: 10 long 0x0082 > +arg_in: 11 long 0x00fe > +arg_in: 12 long 0x007f > +arg_in: 13 long 0x00aa > +arg_in: 14 long 0x0055 > +arg_in: 15 int 0 > + > + > +[test] > +name: SALU i16 in i64 0 > +kernel_name: s_sext_in_reg_i16_in_i64 > + > +arg_out: 0 buffer long[14] \ > + 0x0000 \ > + 0x000f \ > + 0x00f0 \ > + 0x0f00 \ > + 0xfffffffffffff000 \ > + 0xffffffffffffffff \ > + 0xfffffffffffffffe \ > + 0xffffffffffff8000 \ > + 0xffffffffffffaaaa \ > + 0xffffffffffffbbbb \ > + 0xff \ > + 0xfff \ > + 0xffffffffffff8001 \ > + 0xffffffffffff8080 > + > +arg_in: 1 long 0x0000 > +arg_in: 2 long 0x000f > +arg_in: 3 long 0x00f0 > +arg_in: 4 long 0x0f00 > +arg_in: 5 long 0xf000 > +arg_in: 6 long 0xffff > +arg_in: 7 long 0xfffe > +arg_in: 8 long 0x8000 > +arg_in: 9 long 0xaaaa > +arg_in: 10 long 0xbbbb > +arg_in: 11 long 0x00ff > +arg_in: 12 long 0x0fff > +arg_in: 13 long 0x8001 > +arg_in: 14 long 0x8080 > +arg_in: 15 int 0 > + > + > +[test] > +name: SALU i32 in i64 0 > +kernel_name: s_sext_in_reg_i32_in_i64 > +global_size: 1 0 0 > + > +arg_out: 0 buffer long[16] \ > + 0 \ > + 0xf \ > + 0xffffffffffff0000 \ > + 0xffff \ > + 0xffffffffffffffff \ > + 0xfffffffffffffffe \ > + 0x7fffffff \ > + 0xffffffff80000000 \ > + 0xffffffff80000001 \ > + 0x0 \ > + 0xffffffffaaaaaaaa \ > + 0x55555555 \ > + 0x40000000 \ > + 0xff \ > + 0x0 \ > + 0x0 > + > +arg_in: 1 long 0x0 > +arg_in: 2 long 0xf > +arg_in: 3 long 0xffff0000 > +arg_in: 4 long 0x0000ffff > +arg_in: 5 long 0xffffffffffffffff > +arg_in: 6 long 0xfffffffffffffffe > +arg_in: 7 long 0x7fffffff > +arg_in: 8 long 0x80000000 > +arg_in: 9 long 0x80000001 > +arg_in: 10 long 0x100000000 > +arg_in: 11 long 0xaaaaaaaa > +arg_in: 12 long 0x55555555 > +arg_in: 13 long 0x40000000 > +arg_in: 14 long 0xff > +arg_in: 15 long 0x0000000100000000 > +arg_in: 16 long 0x100000000000000 > + > +arg_in: 17 int 0 > + > + > +[test] > +name: VALU i8 in i64 0 > +kernel_name: v_sext_in_reg_i8_in_i64 > +global_size: 14 0 0 > + > +arg_out: 0 buffer long[14] \ > + 0x0000 \ > + 0x000f \ > + 0xfffffffffffffff0 \ > + 0xffffffffffffffff \ > + 0x0 \ > + 0x0 \ > + 0x79 \ > + 0xffffffffffffff80 \ > + 0xffffffffffffff81 \ > + 0xffffffffffffff82 \ > + 0xfffffffffffffffe \ > + 0x7f \ > + 0xffffffffffffffaa \ > + 0x55 > + > +arg_in: 1 buffer long[14] \ > + 0x0000 \ > + 0x000f \ > + 0x00f0 \ > + 0x00ff \ > + 0xff00 \ > + 0xf000 \ > + 0x0079 \ > + 0x0080 \ > + 0x0081 \ > + 0x0082 \ > + 0x00fe \ > + 0x007f \ > + 0x00aa \ > + 0x0055 > + > +arg_in: 2 int 0 > + > +[test] > +name: VALU i16 in i64 0 > +kernel_name: v_sext_in_reg_i16_in_i64 > +global_size: 14 0 0 > + > +arg_out: 0 buffer long[14] \ > + 0x0000 \ > + 0x000f \ > + 0x00f0 \ > + 0x0f00 \ > + 0xfffffffffffff000 \ > + 0xffffffffffffffff \ > + 0xfffffffffffffffe \ > + 0xffffffffffff8000 \ > + 0xffffffffffffaaaa \ > + 0xffffffffffffbbbb \ > + 0xff \ > + 0xfff \ > + 0xffffffffffff8001 \ > + 0xffffffffffff8080 > + > +arg_in: 1 buffer long[14] \ > + 0x0000 \ > + 0x000f \ > + 0x00f0 \ > + 0x0f00 \ > + 0xf000 \ > + 0xffff \ > + 0xfffe \ > + 0x8000 \ > + 0xaaaa \ > + 0xbbbb \ > + 0x00ff \ > + 0x0fff \ > + 0x8001 \ > + 0x8080 > + > +arg_in: 2 int 0 > + > +[test] > +name: VALU i32 in i64 0 > +kernel_name: v_sext_in_reg_i32_in_i64 > +global_size: 16 0 0 > + > +arg_out: 0 buffer long[16] \ > + 0 \ > + 0xf \ > + 0xffffffffffff0000 \ > + 0xffff \ > + 0xffffffffffffffff \ > + 0xfffffffffffffffe \ > + 0x7fffffff \ > + 0xffffffff80000000 \ > + 0xffffffff80000001 \ > + 0x0 \ > + 0xffffffffaaaaaaaa \ > + 0x55555555 \ > + 0x40000000 \ > + 0xff \ > + 0x0 \ > + 0x0 > + > +arg_in: 1 buffer long[16] \ > + 0x0 \ > + 0xf \ > + 0xffff0000 \ > + 0x0000ffff \ > + 0xffffffffffffffff \ > + 0xfffffffffffffffe \ > + 0x7fffffff \ > + 0x80000000 \ > + 0x80000001 \ > + 0x100000000 \ > + 0xaaaaaaaa \ > + 0x55555555 \ > + 0x40000000 \ > + 0xff \ > + 0x0000000100000000 \ > + 0x1000000000000000 > + > +arg_in: 2 int 0 > + > + > + > +!*/ > + > + > +// This test is mostly hacks to make sure we get 64-bit s_loads. > +kernel void s_sext_in_reg_i8_in_i64(global long* restrict out, > + long a0, > + long a1, > + long a2, > + long a3, > + long a4, > + long a5, > + long a6, > + long a7, > + long a8, > + long a9, > + long a10, > + long a11, > + long a12, > + long a13, > + int shift0) > +{ > + long args[] = > + { > + a0, a1, a2, a3, > + a4, a5, a6, a7, > + a8, a9, a10, a11, > + a12, a13 > + }; > + > + // Force unrolling to make sure we don't dynamically index the array. > + #pragma unroll
I think this would need cl_nv_pragma_unroll to guarantee that it's not ignored, but if it works with current clang, meh. Reviewed-by: Jan Vesely <[email protected]> Jan > + for (int i = 0 ; i < sizeof(args) / sizeof(args[0]); ++i) > + { > + // Shift by zero to make sure we load the whole 64-bit value. > + long x = args[i] << shift0; > + out[i] = (x << 56) >> 56; > + } > +} > + > +kernel void s_sext_in_reg_i16_in_i64(global long* restrict out, > + long a0, > + long a1, > + long a2, > + long a3, > + long a4, > + long a5, > + long a6, > + long a7, > + long a8, > + long a9, > + long a10, > + long a11, > + long a12, > + long a13, > + int shift0) > +{ > + long args[] = > + { > + a0, a1, a2, a3, > + a4, a5, a6, a7, > + a8, a9, a10, a11, > + a12, a13 > + }; > + > + // Force unrolling to make sure we don't dynamically index the array. > + #pragma unroll > + for (int i = 0 ; i < sizeof(args) / sizeof(args[0]); ++i) > + { > + // Shift by zero to make sure we load the whole 64-bit value. > + long x = args[i] << shift0; > + out[i] = (x << 48) >> 48; > + } > +} > + > +// This test is mostly hacks to make sure we get 64-bit s_loads. > +kernel void s_sext_in_reg_i32_in_i64(global long* restrict out, > + long a0, > + long a1, > + long a2, > + long a3, > + long a4, > + long a5, > + long a6, > + long a7, > + long a8, > + long a9, > + long a10, > + long a11, > + long a12, > + long a13, > + long a14, > + long a15, > + int shift0) > +{ > + long args[] = > + { > + a0, a1, a2, a3, > + a4, a5, a6, a7, > + a8, a9, a10, a11, > + a12, a13, a14, a15 > + }; > + > + // Force unrolling to make sure we don't dynamically index the array. > + #pragma unroll > + for (int i = 0 ; i < sizeof(args) / sizeof(args[0]); ++i) > + { > + // Shift by zero to make sure we load the whole 64-bit value. > + long x = args[i] << shift0; > + out[i] = (x << 32) >> 32; > + } > +} > + > +kernel void v_sext_in_reg_i8_in_i64(global long* restrict out, > + global long* restrict in, > + int shift0) > +{ > + int id = get_global_id(0); > + long x = in[id] << shift0; > + out[id] = (x << 56) >> 56; > +} > + > +kernel void v_sext_in_reg_i16_in_i64(global long* restrict out, > + global long* restrict in, > + int shift0) > +{ > + int id = get_global_id(0); > + long x = in[id] << shift0; > + out[id] = (x << 48) >> 48; > +} > + > +kernel void v_sext_in_reg_i32_in_i64(global long* restrict out, > + global long* restrict in, > + int shift0) > +{ > + int id = get_global_id(0); > + long x = in[id] << shift0; > + out[id] = (x << 32) >> 32; > +}
signature.asc
Description: This is a digitally signed message part
_______________________________________________ Piglit mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/piglit
