We have detected your current browser version is not the latest one. Xilinx.com uses the latest web technologies to bring you the best online experience possible. Please upgrade to a Xilinx.com supported browser:Chrome, Firefox, Internet Explorer 11, Safari. Thank you!

cancel
Showing results for
Did you mean:
Visitor
655 Views
Registered: ‎06-18-2018

## Loop bound is not divisible by vector type uint16 length

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?

Tags (3)
1 Solution

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

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

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

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

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

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

Visitor
635 Views
Registered: ‎06-18-2018

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

@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?

Xilinx Employee
703 Views
Registered: ‎01-12-2017

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

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