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:
Observer
654 Views
Registered: ‎06-12-2008

How to make opencl_unroll attribute work in float loop ?

Hi,

We tried to use the opencl_unroll to optimize the multiply c = a * b in a loop.  If the data is int (int c = int a  * int b ),  sdaccel works fine and uses 64 x multiplies to implement the loop in the RTL. But if changed the data from int to float ( float c  = float a * float b), we were assuming sdaccel would take 64 x float multiplies for the loop, but sdaceel only used 2 x fmult in the generated RTL design.  It looks like that unroll attribute doesn't work in the float loop.

My question : Is there any missing in my opencl code ?  How to use the loop unroll optimization in the float loop  ?

BTW, we tried the same code both in sdaccel 2017.4 & 2018.2 , fpga board : KCU1500.

Attached the source code & log files .

#define LENGTH 64
__kernel void __attribute__ ((reqd_work_group_size(1, 1, 1)))
krnl_vmult(
__global float* a,
__global float* b,
__global float* c)
{
local float bufa[LENGTH];
local float bufb[LENGTH];
local float bufc[LENGTH];

for(int i = 0; i < LENGTH; i++) {
bufa[i] = a[i];
bufb[i] = b[i];
}
__attribute__((opencl_unroll_hint(64)))
for(int i = 0; i < LENGTH; i++) {
bufc[i] = bufa[i] * bufb[i];
}

for(int i = 0; i < LENGTH; i++) {
c[i] = bufc[i];
}
return;
}

1 Solution

Accepted Solutions
Xilinx Employee
705 Views
Registered: ‎07-18-2014

Re: How to make opencl_unroll attribute work in float loop ?

Hi @zxttao,

Did you also try to partition the array which are getting used inside unroll loop?

YOu can refer below example which explain about unroll using partition:

https://github.com/Xilinx/SDAccel_Examples/tree/master/getting_started/cpu_to_fpga/04_partition_ocl

-Heera

Xilinx Employee
706 Views
Registered: ‎07-18-2014

Re: How to make opencl_unroll attribute work in float loop ?

Hi @zxttao,

Did you also try to partition the array which are getting used inside unroll loop?

YOu can refer below example which explain about unroll using partition:

https://github.com/Xilinx/SDAccel_Examples/tree/master/getting_started/cpu_to_fpga/04_partition_ocl

-Heera