topic vec_type_hint() does not work for me in OpenCL* for CPU

vec_type_hint() does not work for me

Fu_J_Intel — Thu, 12 May 2016 21:08:44 GMT

trying to disable auto-vectorization, I used vec_type_hint(char) as below. But checking the built assembly code, I see the kernel is still compiled as SIMD32. Any advice?

__kernel __attribute__((vec_type_hint(uchar))) void modulate_v1_uchar(global const uchar *pSrc, global uchar *pDst)
//__kernel void modulate_v1_uchar(global const uchar *pSrc, global uchar *pDst)
{

uint idx = get_global_id(0);

   uchar src = pSrc[idx];
   src >>= MODULATE_SHIFT_FACTOR;
   pDst[idx] = src;
}

You won't be able to disable

Robert_I_Intel — Fri, 13 May 2016 01:46:29 GMT

You won't be able to disable vectorization on the GPU. That hint is for the CPU. You could use _attribute__((reqd_work_gr oup_size(X, Y, Z))) could to change SIMD width to 8 (8, 1, 1), 16 (16, 1, 1) or 32 (32, 1, 1)

great! thanks Robert.

Fu_J_Intel — Fri, 13 May 2016 05:46:42 GMT

great! thanks Robert.

I actually have a further

Fu_J_Intel — Fri, 13 May 2016 18:36:44 GMT

I actually have a further question: how is workgroup size (eg, _attribute__((reqd_work_gr oup_size(X, Y, Z)))) related to SIMD width? does it mean work-item of different workgroups can not be handled on the same thread? Thanks

Yes, typically a work-group

Robert_I_Intel — Sat, 14 May 2016 00:12:13 GMT

Yes, typically a work-group maps to a hardware thread or can span multiple threads or even multiple EUs.505229

See slides 18 thru 34 in the

Robert_I_Intel — Sat, 14 May 2016 00:13:11 GMT

See slides 18 thru 34 in the presentation above.