Forum Discussion
Kernel Vectorization: branching is thread ID dependent ... cannot vectorize
I have the following code, and in the kernel the norm is get from the rowdata and row data is dependent on the thread ID.
In the myfastsquaredDistance function, the if else statement is dependent on the norm, thus when I compile, it gives warning that it cannot vectorize. Does anyone know how to get around this? Thanks!
double myfastSquaredDistance(__constant double *rowdata, __constant double *rowcenter, double norm1, double norm2, int n, int y)
{
double precision = 1e-6;
double sumSquaredNorm = norm1 * norm1 + norm2 * norm2;
double normDiff = norm1 - norm2;
double sqDist = 0.0;
double precisionBound1 = 2.0 * epsilon() * sumSquaredNorm / (normDiff * normDiff + epsilon());
if (precisionBound1 < precision) {
sqDist = sumSquaredNorm - 2.0 * mydot(rowdata, rowcenter, n, y);
} else {
sqDist = sqdist(rowdata, rowcenter, n, y);
}
return sqDist;
}
__attribute__((num_simd_work_items(4)))
__attribute__((reqd_work_group_size(32,1,1)))
__kernel
void kmeansTest(
int num_vectors,
int vector_length,
int num_clusters,
__constant double* restrict data,
__constant double* restrict center,
__global double* restrict result
)
{
const uint y = get_global_id(0);
int bestindex;
double norm1, norm2;
double min_distance = FLT_MAX;
__constant double* rowdata = &data [y * (vector_length + 1)];
norm1 = rowdata[vector_length];
for (int i = 0; i < num_clusters; ++i){
double dotProduct = 0;
double precision = 1e-6;
__constant double *rowcenter = center + i * (vector_length + 1);
norm2 = rowcenter[vector_length];
dotProduct = myfastSquaredDistance(data, rowcenter, norm1, norm2, vector_length, y);
if(dotProduct < min_distance){
min_distance = dotProduct;
bestindex = i;
}
}
result[y * 2] = bestindex;
result[y * 2 + 1] = sqrt(min_distance);
}
6 Replies
- Nooraini_Y_Intel
Frequent Contributor
Hi,
Currently I am reviewing the forum for any open questions and found this thread. I apologize that no one seems to answer this question that you posted. Since it has been a while you posted this question, I'm wondering if you have found the answer? If not, please let me know, I will try to assign someone to assist you. Thank you.
Regards,
Nooraini
- RJin1
New Contributor
Hi,
I haven't figured out the answer yet. If you can find someone to help me that would be great. Thanks!
- Nooraini_Y_Intel
Frequent Contributor
Hi RJin1,
Sure, I'm checking here to find someone that may able to assist you on this thread.
Regards,
Nooraini
- MuhammadAr_U_Intel
Contributor
- RJin1
New Contributor
16.0 sdk
- HRZ
Frequent Contributor
Your code does not compile as it is due to missing definition for many of the functions. I believe it requires an external header. Please provide a code snippet that can be compiled (preferably using the code insertion mechanism provided in the forum so that indentation is preserved) so that we can take a look at the report and compiler messages. I fail to see why this cannot be vectorized, though; the statement in the condition does not depend on thread-ID, only its value does. You might be able to get around it using the (condition) ? (value1) : (value2) notation instead of if/else.