On Tue, 2016-12-06 at 11:14 -0800, [email protected] wrote: > From: Matt Arsenault <[email protected]> > > --- > .../execute/amdgcn-f32-inline-immediates.cl | 237 > +++++++++++++++++++++ > 1 file changed, 237 insertions(+) > create mode 100644 tests/cl/program/execute/amdgcn-f32-inline-immediates.cl > > diff --git a/tests/cl/program/execute/amdgcn-f32-inline-immediates.cl > b/tests/cl/program/execute/amdgcn-f32-inline-immediates.cl > new file mode 100644 > index 0000000..3541dee > --- /dev/null > +++ b/tests/cl/program/execute/amdgcn-f32-inline-immediates.cl > @@ -0,0 +1,237 @@ > +/*! > +[config] > +name: SI float inline immediates > +clc_version_min: 10 > + > +dimensions: 1 > +global_size: 10 0 0 > + > +## Addition ## > + > +[test] > +name: add 0.5 > +kernel_name: add_half > +arg_out: 0 buffer float[10] \ > + 0.5 1.0 0.0 1.5 -0.5 \ > + 2.5 -1.5 4.5 -3.5 123.5 > + > +arg_in: 1 buffer float[10] \ > + 0.0 0.5 -0.5 1.0 -1.0 \ > + 2.0 -2.0 4.0 -4.0 123.0 > + > +[test] > +name: add 1.0 > +kernel_name: add_1 > +arg_out: 0 buffer float[10] \ > + 1.0 1.5 0.5 2.0 0.0 \ > + 3.0 -1.0 5.0 -3.0 124.0 > + > +arg_in: 1 buffer float[10] \ > + 0.0 0.5 -0.5 1.0 -1.0 \ > + 2.0 -2.0 4.0 -4.0 123.0 > + > +[test] > +name: add 2.0 > +kernel_name: add_2 > +arg_out: 0 buffer float[10] \ > + 2.0 2.5 1.5 3.0 1.0 \ > + 4.0 0.0 6.0 -2.0 125.0 > + > +arg_in: 1 buffer float[10] \ > + 0.0 0.5 -0.5 1.0 -1.0 \ > + 2.0 -2.0 4.0 -4.0 123.0 > + > + > +[test] > +name: add 4.0 > +kernel_name: add_4 > +arg_out: 0 buffer float[10] \ > + 4.0 4.5 3.5 5.0 3.0 \ > + 6.0 2.0 8.0 0.0 127.0 > + > +arg_in: 1 buffer float[10] \ > + 0.0 0.5 -0.5 1.0 -1.0 \ > + 2.0 -2.0 4.0 -4.0 123.0 > + > +[test] > +name: add rcp(2pi) > +kernel_name: add_inv_2pi > +arg_out: 0 buffer float[10] \ > + 0x1.45f306p-3 0x1.517cc2p-1 -0x1.5d067cp-2 0x1.28be6p+0 -0x1.ae833ep-1 > \ > + 0x1.145f3p+1 -0x1.d741ap+0 0x1.0a2f98p+2 -0x1.eba0dp+1 0x1.eca2fap+6 > + > +arg_in: 1 buffer float[10] \ > + 0.0 0.5 -0.5 1.0 -1.0 \ > + 2.0 -2.0 4.0 -4.0 123.0 > + > +[test] > +name: sub 0.5 > +kernel_name: sub_half > +arg_out: 0 buffer float[10] \ > + -0.5 0.0 -1.0 0.5 -1.5 \ > + 1.5 -2.5 3.5 -4.5 122.5 > + > +arg_in: 1 buffer float[10] \ > + 0.0 0.5 -0.5 1.0 -1.0 \ > + 2.0 -2.0 4.0 -4.0 123.0 > + > +[test] > +name: sub 1.0 > +kernel_name: sub_1 > +arg_out: 0 buffer float[10] \ > + -1.0 -0.5 -1.5 0.0 -2.0 \ > + 1.0 -3.0 3.0 -5.0 122.0 > + > +arg_in: 1 buffer float[10] \ > + 0.0 0.5 -0.5 1.0 -1.0 \ > + 2.0 -2.0 4.0 -4.0 123.0 > + > +[test] > +name: sub 2.0 > +kernel_name: sub_2 > +arg_out: 0 buffer float[10] \ > + -2.0 -1.5 -2.5 -1.0 -3.0 \ > + 0.0 -4.0 2.0 -6.0 121.0 > + > +arg_in: 1 buffer float[10] \ > + 0.0 0.5 -0.5 1.0 -1.0 \ > + 2.0 -2.0 4.0 -4.0 123.0 > + > +[test] > +name: sub 4.0 > +kernel_name: sub_4 > +arg_out: 0 buffer float[10] \ > + -4.0 -3.5 -4.5 -3.0 -5.0 \ > + -2.0 -6.0 0.0 -8.0 119.0 > + > +arg_in: 1 buffer float[10] \ > + 0.0 0.5 -0.5 1.0 -1.0 \ > + 2.0 -2.0 4.0 -4.0 123.0 > + > +[test] > + > +# FIXME: This is a denormal and is assumed to flush to zero. This > +# should change the checked value depending on denormal support.
This one can be fixed by using -cl-denorms-are-zero build options. It
will flush when denorms are supported, and ignore the option (and
flush) when denorms are not supported. without it the test fails on
intel CPU implementation.
> +
> +name: add integer 64
> +kernel_name: add_i32_64
> +arg_out: 0 buffer float[10] \
> + 0x0 0x1p-1 -0x1p-1 0x1p+0 -0x1p+0 \
> + 0x1p+1 -0x1p+1 0x1p+2 -0x1p+2 0x1.ecp+6
> +
> +arg_in: 1 buffer float[10] \
> + 0.0 0.5 -0.5 1.0 -1.0 \
> + 2.0 -2.0 4.0 -4.0 123.0
> +
> +[test]
> +name: add other
> +kernel_name: add_other
> +arg_out: 0 buffer float[10] \
> + 10.0 10.5 9.5 11.0 9.0 \
> + 12.0 8.0 14.0 6.0 133.0
> +
> +arg_in: 1 buffer float[10] \
> + 0.0 0.5 -0.5 1.0 -1.0 \
> + 2.0 -2.0 4.0 -4.0 123.0
> +
> +[test]
> +name: sub other
> +kernel_name: sub_other
> +
> +arg_out: 0 buffer float[10] \
> +-10.0 -9.5 -10.5 -9.0 -11.0 \
> + -8.0 -12.0 -6.0 -14.0 113.0
> +
> +arg_in: 1 buffer float[10] \
> + 0.0 0.5 -0.5 1.0 -1.0 \
> + 2.0 -2.0 4.0 -4.0 123.0
> +
> +!*/
> +
> +#pragma OPENCL EXTENSION cl_khr_fp64 : enable
This should be dropped ^^
> +
> +
> +kernel void add_half(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] + 0.5;
> +}
> +
> +kernel void add_1(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] + 1.0;
> +}
> +
> +kernel void add_2(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] + 2.0;
> +}
> +
> +kernel void add_4(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] + 4.0;
> +}
> +
> +kernel void add_inv_2pi(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] + 0x1.45f306p-3;
> +}
> +
> +kernel void sub_half(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] - 0.5;
> +}
> +
> +kernel void sub_1(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] - 1.0;
> +}
> +
> +kernel void sub_2(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] - 2.0;
> +}
> +
> +kernel void sub_4(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] - 4.0;
> +}
> +
> +kernel void add_i32_neg16(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] + as_float((int)-16);
> +}
> +
> +kernel void add_i32_64(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] + as_float((int)64);
> +}
> +
> +kernel void add_i32_1(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] + as_float((int)1);
> +}
> +
> +kernel void add_other(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] + 10.0;
> +}
> +
> +kernel void sub_other(global float* out, global float* in)
> +{
> + int id = get_global_id(0);
> + out[id] = in[id] - 10.0;
> +}
> +
git complains about whitespace error here.
Reviewed-by: Jan Vesely <[email protected]>
I fixed all 3 issues locally, I can push it if you're OK with those
changes. It passes on r600 and intel CPU.
Jan
signature.asc
Description: This is a digitally signed message part
_______________________________________________ Piglit mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/piglit
