UPGRADE YOUR BROWSER

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 
Search instead for 
Did you mean: 
Visitor skotti
Visitor
655 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
Xilinx Employee
Xilinx Employee
702 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

3 Replies
Xilinx Employee
Xilinx Employee
644 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
Visitor skotti
Visitor
635 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
Xilinx Employee
Xilinx Employee
703 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