如何对Clang编译的OpenMP卸载代码进行性能分析？

问题描述 #

使用从GitHub构建的LLVM/clang-16进行OpenMP卸载开发，已通过LIBOMPTARGET_PROFILE=profile.json、LIBOMPTARGET_INFO等环境变量确认代码在GPU上执行，但使用NVIDIA Nsight工具集中的nvprof或ncu分析时，出现如下警告：

> ncu ./saxpy
Time of kernel: 0.000004
==WARNING== No kernels were profiled.
==WARNING== Profiling kernels launched by child processes requires the --target-processes all option.

测试代码：

#include <iostream>
#include <omp.h>
#include <cstdlib>

void saxpy(float a, float* x, float* y, int sz) {
    double t = 0.0;
    double tb, te;
    tb = omp_get_wtime();
#pragma omp target teams distribute parallel for map(to:x[0:sz]) map(tofrom:y[0:sz])
{
    for (int i = 0; i < sz; i++) {
        y[i] = a * x[i] + y[i];
    }
}
    te = omp_get_wtime();
    t = te - tb;
    printf("Time of kernel: %lf\n", t);
}

int main() {
    auto x = (float*) malloc(1000 * sizeof(float));
    auto y = (float*) calloc(1000, sizeof(float));
    
    for (int i = 0; i < 1000; i++) {
        x[i] = i;
    }
    
    saxpy(42, x, y, 1000);
    
    return 0;
}

编译命令：

> clang++ -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda main.cpp -o saxpy --cuda-path=/opt/nvidia/hpc_sdk/Linux_x86_64/22.11/cuda/10.2 --offload-arch=sm_61 -fopenmp-offload-mandatory

已知他人无需额外步骤即可用ncu分析Clang编译的OpenMP卸载代码，需解决配置问题以启用性能分析。

解决方案 #

• 添加子进程追踪参数：Clang的OpenMP运行时可能通过子进程启动GPU内核，按照警告提示添加--target-processes all参数，执行命令改为：

  ncu --target-processes all ./saxpy
  

• 强制主进程内启动内核：设置环境变量LIBOMPTARGET_USE_HSA=0，强制OpenMP运行时在主进程内启动GPU内核，避免子进程导致的追踪问题：

  export LIBOMPTARGET_USE_HSA=0
  ncu ./saxpy
  

• 调整编译选项提升可分析性：编译时添加-g（生成调试信息）和-O0（禁用优化），帮助Nsight工具更准确捕获内核信息：

  clang++ -g -O0 -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda main.cpp -o saxpy --cuda-path=/opt/nvidia/hpc_sdk/Linux_x86_64/22.11/cuda/10.2 --offload-arch=sm_61 -fopenmp-offload-mandatory
  

• 检查版本兼容性：clang-16推荐搭配CUDA 11.x及以上版本，当前使用的CUDA 10.2可能存在兼容性问题，尝试升级CUDA到11.x版本后重新编译测试。

内容的提问来源于stack exchange，提问作者Dogyman