cancel
Showing results for 
Search instead for 
Did you mean: 
Highlighted
Observer
Observer
810 Views
Registered: ‎06-18-2018

Loop bound is not divisible by vector type uint16 length

Jump to solution

Hello everyone!

I have a problem with vector types.  I don't know the right way to address an individual element.

For example:

I have a loop which is:

for (int i = 0; i < 99; i++) {

 <some array computations>

}

 

and I want to use int16. Obviously the loop bound is not divisible on 16. I don't want to execute one more operation which will exceed 99 and poison  other values in array. How can I make computations on the other 3 array elements (99/16 = 6, 1875, 6*16 = 96)? Should I use any shift registers or so?

0 Kudos
1 Solution

Accepted Solutions
Highlighted
Xilinx Employee
Xilinx Employee
857 Views
Registered: ‎01-12-2017

Re: Loop bound is not divisible by vector type uint16 length

Jump to solution

Hi @skotti ,

 

Please have a look at following examples which will help you to understand vector data type usage in SDAccel OpenCL kernel better.

 

https://github.com/Xilinx/SDAccel_Examples/tree/master/vision/watermarking

https://github.com/Xilinx/SDAccel_Examples/tree/master/vision/median_filter

 

Khronos OpenCL Reference Guide:

https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf [Chapter: 6.1.6, 6.1.7]

 

Thanks

Kali

View solution in original post

3 Replies
Highlighted
Xilinx Employee
Xilinx Employee
799 Views
Registered: ‎01-12-2017

Re: Loop bound is not divisible by vector type uint16 length

Jump to solution

Hi @skotti

 

Elements of vector data types can be accessed using numerical indexes as below

 

Vector data type width, 

 

2 - 0, 1

4 - 0, 1, 2, 3

8 - 0, 1, 2, 3, 4, 5, 6, 7

16 - 0, 1, 2, 3, 4, 5, 6, 7,8, 9, a, A, b, B, c, C, d, D, e, E,f, F

These indexes must be preceded by s or S.

 

For example:

 

int16 a ; // If you want to access a 3rd element you can do this way --> a.s2 

 

Thanks

Kali

0 Kudos
Highlighted
Observer
Observer
790 Views
Registered: ‎06-18-2018

Re: Loop bound is not divisible by vector type uint16 length

Jump to solution

@kalib, so will this be kind of correct pattern ?

 

__kernel void kernel0(__global int16 *A)

{

local int16 local_A[7];


for (int i = 0; i < 6; i++){

local_A[i] = A[i];

}


local_A[7].s1 = A[7].s1;

local_A[7].s2 = A[7].s2;

local_A[7].s3 = A[7].s3;


<...>
}

 

 

I mean if it is correct to address a loop with i iterating from new indexes with new step, or step should be 16? I just can't find appropriate examples.

 

I also noticed that such types as double8 are not accessible yet in sdx environment, I had to declare it by myself as typedef double double8 __attribute__((ext_vector_type(8)));

 

Are all the access patterns applicable to such kind of type?

 

And should I declare array the same time in the host code, or it will be mapped automatically to this uint16 when passing to kernel?

 

0 Kudos
Highlighted
Xilinx Employee
Xilinx Employee
858 Views
Registered: ‎01-12-2017

Re: Loop bound is not divisible by vector type uint16 length

Jump to solution

Hi @skotti ,

 

Please have a look at following examples which will help you to understand vector data type usage in SDAccel OpenCL kernel better.

 

https://github.com/Xilinx/SDAccel_Examples/tree/master/vision/watermarking

https://github.com/Xilinx/SDAccel_Examples/tree/master/vision/median_filter

 

Khronos OpenCL Reference Guide:

https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf [Chapter: 6.1.6, 6.1.7]

 

Thanks

Kali

View solution in original post