Support 64-bit float

support:
  arithmetic(+ - *)
  store load
  immediate_value
  if else
  select

other change:
  add "nib control" field in machine instruction format
  support "nib control"
  fix "directly store float-64 after load float-64".
  change hard coded store size (4) to flexible size (4 or 8)
  when using float-64 load(store), change to SIMD8

example:

/* support arithmetic store load immediate_value */
kernel void f(global double *src, global double *dst) {
  int i = get_global_id(0);
  double d = 1.234567890123456789;
  dst[i] = d * (src[i] + d);
}

/* support if else */
kernel void f(global float *src, global double *dst) {
  int i = get_global_id(0);
  float d = 1.234567890123456789f;
  if (i < 14)
    dst[i] = d * (d + src[i]);
  else
    dst[i] = 14;
}

/* support select */
kernel void f(global float *src, global double *dst) {
  int i = get_global_id(0);
  float d = 1.234567890123456789f;
  dst[i] = i < 14 ? d : 14;
}

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Yang Rong <rong.r.yang@intel.com>

0 files changed, 0 insertions, 0 deletions