Honored ContributorHello all,
I have finally gotten some code I have been working on to compile, specifically after installing the April update to the OpenCL compiler. However, the Compile has been going for almost a week now, is there something wrong or is a long compile time possible for complex code? Sincerely, Aaron S.
Honored ContributorWhen exactly, which task (analyse, fitter...) does it spend very much time ?
You have also Quartus Advisor in menu Quartus > Tools > .... > compile time advisor where you get advices and where you can "fix" settings.Honored ContributorIt is during the OpenCL compile that I have the problem, I am only running aoc -c -v etc.
Sincerely, Aaron S.Honored ContributorDouble check your kernel to make sure there are not any coding structures that could cause the compiler problems (like recursion). If you have any of the performance attributes or the unroll pragma defined in your kernel try scaling those back. The only time I have run into the issue you are seeing is when had an incredibly deep dependency tree in my kernel that became much worst when I unrolled a long loop.
Honored ContributorThank you for your reply, I guess I have 2 questions.
1) Is the compile in an infinite loop? 2) Without unrolling all of the loops the logic utilization was above 100% so is there a way to balance the two issues, loop unrolls and logic utilization? Sincerely, Aaron S.Honored ContributorThe compiler will probably finish eventually but I always end up killing it because if it takes that long to generate the logic then I Quartus will take a while to synthesize the logic.
I don't really have an answer for# 2 because it's very kernel specific and each way you tackle the problem is different. If you can share kernel through a service request or FAE I'm sure Altera can take a look and make suggestions. One thing to look into is whether your kernel is sequential in nature. For example the vector_add kernel posted on the design examples page on www.altera.com (http://www.altera.com) could easily be written as a task kernel like this:__attribute ((task))
__kernel void vectorAdd(__global const float * restrict x,
__global const float * restrict y,
__global float * restrict z,
int N)
{
// iterate over elements of the arrays instead of indexing using work-item ID
for(int index = 0; index < N; index++)
{
// add the vector elements
z = x + y;
}
}
I do not know if this is the exact issue here, but the previous versions of the aoc compiler was having trouble trying to schedule instructions when there are large number of load and store instructions (hundreds or thousands or loads/stores). There could be multiple reasons for having excessive number of loads/stores; 1) loop unrolling, 2) a large private array being accessed with dynamic addresses.
The compile-time issue is fixed in 14.0 update 1 release. However, in general, users should try to avoid having excessive number of load/store instructions for performance reasons.Honored ContributorI think I'm running into the "long compile times due to large loops being unrolled issue." I have a task kernel and synthesis does not complete in a reasonable amount of time if I unroll over 320. When can we expect 14.0 update 1?
"However, in general, users should try to avoid having excessive number of load/store instructions for performance reasons." If I'm doing a large parallel load, what are my other options? Should I create an I/O kernel that writes to a channel instead?Honored ContributorHonored ContributorAccording to my calendar, the update release should be available now.
I am still assuming you are talking about "aoc -c" compile time and your parallel load is from global memory. If so, make sure your access pattern is simple. If you are accessing consecutive elements in global buffer and unroll the loop, the compiler will merge your accesses to form wider accesses which reduces the number of loads. You can also try sending the data to the kernel via channel as you suggested. This may increase resource usage due to channels, but may make compile-time faster because you are splitting your memory accesses in two kernels. --- Quote Start --- I think I'm running into the "long compile times due to large loops being unrolled issue." I have a task kernel and synthesis does not complete in a reasonable amount of time if I unroll over 320. When can we expect 14.0 update 1? "However, in general, users should try to avoid having excessive number of load/store instructions for performance reasons." If I'm doing a large parallel load, what are my other options? Should I create an I/O kernel that writes to a channel instead? --- Quote End ---Honored Contributor
