performance - 编译警告OpenCL矩阵乘法

为什么这不是矢量化？

__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 ...无法向量化。

最佳答案

问：为什么不进行向量化？


邪恶之处在于“分支……无法向量化”-它与以下指令有关：

if( ( i < gDIM ) && ( j < fDIM ) ){ ... }

高效的基于SIMD指令的矢量化意味着所有代码执行流都不会“分叉”（分叉），并且会“执行”完全相同的数据/指令（即将数据元素SIMD-“粘合”到DATA矢量中，放入足够宽的范围内） ，CPU和SIMD友好的寄存器，它们通过一条SIMD友好的指令立即进行计算-也就是说，对于每一个线程中的SIMD友好指令，它们都是完全相同的，即不会if(){...}else{...}分成不同的，针对不同数据元素的不同指令的不同序列的“分歧”流

原则上，不可能对要对齐SIMD友好型CPU寄存器的数据的不同部分执行不同的操作-对于存储在SIMD友好型CPU中​​的所有矢量分量，一次只能执行一条SIMD友好型指令-寄存器。

整数和浮点SIMD矢量指令的硬件细节各不相同，由此产生的微操作延迟也很重要，编译器针对SIMD处理器的具体细节确实很重要，但是避免分散路径的原理对于自动SIMD矢量化是很普遍的。编译器阶段。有关SIMD指令及其更多性能限制属性的更多详细信息，请阅读Agner

关于performance - 编译警告OpenCL矩阵乘法，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/59359329/