Community
cancel
Showing results for 
Search instead for 
Did you mean: 
Fu_J_Intel
Employee
34 Views

vec_type_hint() does not work for me

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;
}

0 Kudos
5 Replies
Robert_I_Intel
Employee
34 Views

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)

Fu_J_Intel
Employee
34 Views

great! thanks Robert.

Fu_J_Intel
Employee
34 Views

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

 

 

Robert_I_Intel
Employee
34 Views

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

Robert_I_Intel
Employee
34 Views

See slides 18 thru 34 in the presentation above.