为什么这不是矢量化?
__attribute__((num_simd_work_items(4)))
__attribute__((num_compute_units(2)))
__attribute__((reqd_work_group_size(16,16,1)))
__kernel void matrix_multiplication(const int fDIM,const int gDIM, const int hDIM,
__global float* A, __global float* B, __global float* C) {
int k;
int i = get_global_id(0);
int j = get_global_id(1);
float temp_result;
if((i < gDIM) && (j<fDIM)){
temp_result= 0.0f;
for(k = 0; k<hDIM;k++) {
temp_result+= A[i*gDIM+k] * B[k*hDIM+j];
}
C[i*gDIM+j] = temp_result;
}
}
编译器警告:
内核矢量化:分支依赖于线程ID .无法矢量化。
发布于 2019-12-16 08:24:01
Q:为什么不是矢量化?
邪恶之处是“分支…无法向量化”--它与以下指令有关:
if( ( i < gDIM ) && ( j < fDIM ) ){ ... }
高效的基于指令的矢量化意味着所有的代码执行流都不是“发散的”(分支),而确实“执行”完全相同的数据/指令(即数据元素SIMD-“粘合”到数据向量中,放入足够宽的CPU、SIMD友好的寄存器中,使一次由一个SIMD友好指令计算-即对每个线程来说完全相同-a-pack D友好指令,即 not if(){...}else{...}
-分叉到不同的,不同数据的不同指令的不同序列的“发散”流.元素
基本上不可能对数据的不同部分执行不同的操作,对齐到SIMD友好的CPU寄存器----一个和唯一一个SIMD友好指令,可以同时执行存储在SIMD友好的CPU寄存器中的所有矢量组件。
关于整数和浮动SIMD矢量指令的硬件细节不同,产生的微操作延迟、SIMD处理器特定细节的编译器确实很重要,但是避免发散路径的原则对于编译器阶段的自动SIMD矢量化是很常见的。要了解更多关于SIMD-指令及其进一步的性能限制属性,可以阅读和学习阿格纳。
https://stackoverflow.com/questions/59359329
复制相似问题