编译警告 OpenCL 矩阵乘法答案

【问题标题】：Compilation warning OpenCL Matrix Multplication编译警告 OpenCL 矩阵乘法
【发布时间】：2020-04-09 01:49:46
【问题描述】：

为什么这不是矢量化的？

__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 ...不能向量化。

【问题讨论】：

标签： performance parallel-processing opencl low-latency pyopencl

【解决方案1】：

问：为什么这不矢量化？

邪恶是“branching…无法矢量化”——它与这条指令有关：

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

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

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

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

【讨论】：