performance - 编译警告 OpenCL 矩阵乘法

Question

为什么这不是矢量化？

__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;

    }
}

编译器警告:

Kernel Vectorization: branching is thread ID dependent ... cannot vectorize.

Answer 1

Q : Why is this not vectorizing?

问题是“分支…无法向量化”——它与这条指令有关:

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

高效的基于 SIMD 指令的矢量化意味着所有代码执行流都不是“发散的”(分支的)并且“执行”完全相同的数据/指令(即数据元素 SIMD-“粘合”放入数据向量中，放入足够宽的、CPU、SIMD 友好的寄存器中，这些寄存器通过单个 SIMD 友好指令立即计算 - 即每个线程都相同-pack SIMD 友好指令，即不 if(){...}else{...} -发散成不同的，不同数据元素的不同指令的不同序列的“发散”流

原则上不可能对数据的不同部分进行不同的操作，对齐到 SIMD 友好的 CPU 寄存器中 - 一次且只有一条 SIMD 友好指令可以针对所有向量组件执行 存储到 SIMD 友好的 CPU 寄存器中。

关于整数和 float SIMD 向量指令的硬件细节各不相同，由此产生的微操作延迟也是如此，SIMD 处理器特定细节形式的编译器非常重要，但避免发散路径的原则对于自动化 SIMD 来说很常见 -编译器阶段的矢量化。有关 SIMD 指令及其进一步性能限制属性的更多详细信息，请阅读并了解 Agner