编译警告 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 指令的矢量化意味着所有代码执行流程都不会 "divergent"（分支）并且 做 "execute" 完全相同 data/instruction（即数据元素 SIMD-"glued" 放入数据向量中，放入足够宽的、CPU、SIMD 友好的寄存器中，由单个 SIMD 立即计算 -friendly instruction - 即每个 thread-in-a-pack SIMD-friendly 指令都相同，即 not if(){...}else{...}-分为不同的、"divergent" 不同数据元素的不同指令序列

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

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

编译警告 OpenCL 矩阵乘法

Compilation warning OpenCL Matrix Multplication

parallel-processing

performance

opencl

low-latency

pyopencl