Honored Contributor
Honored ContributorI think it may be best to describe what you've observed with a simple example:
bool comparison; comparison = get_global_id(0) % 2 == 0 || get_global_id(0) % 2 == 1; // this is always true if (comparison) {...} // conditional 1 and statement 1 if (get_global_id(0) % 2 == 0 || get_global_id(0) % 2 == 1) {...} // conditional 2 and statement 2 The two conditional tests above are not equivalent because of possible early exits, i.e., if the global id is even then you can ignore the second logical test and immediately execute statement 2, while the odd global ids have to perform both logical tests before executing statement 2. This may cause a branch divergence for GPUs given their SIMD architecture and reduce performance by half in the above example. The two statements seem like they would perform equally well for FPGAs. Are you using any kernel attributes, like num_simd_work_items?
