SYCL中sycl::vec传值/传引用选择及目标代码查看方法问询

问题1：sycl::vec<double, 3>的传参方式选择 #

优先选择传值方式，原因如下：

• sycl::vec是可平凡拷贝的轻量值类型，3个double元素仅占24字节，拷贝开销远低于传引用带来的间接寻址开销，尤其是在GPU等加速器设备上执行时，传值可以让编译器直接将数据存入寄存器，无需额外访存操作，性能表现更好
• 若选择传const引用，在设备端编译时虽然大概率会被编译器优化为传值，但在主机-设备参数传递的场景下，传引用可能引发不必要的内存地址捕获，反而增加内核启动开销
• 仅当使用长度大于16元素的大尺寸vec时，才需要评估传const引用的收益，3元素的场景直接传值即可

问题2：SYCL CUDA目标代码的查看方法 #

你使用Intel DPC++/Clang的场景下，可以通过两类方式查看目标代码：

• 在线编译查看工具：Compiler Explorer本身已完整支持SYCL代码编译与目标代码展示，只需选择对应版本的Intel DPC++编译器，添加编译参数-fsycl -fsycl-targets=nvptx64-nvidia-cuda，即可直接查看生成的PTX代码、SPIR-V中间表示、主机端代码等输出
• 本地工具链查看：直接使用DPC++自带的工具即可完成查看：
  • 生成CUDA PTX代码：编译时添加参数-fsycl -fsycl-targets=nvptx64-nvidia-cuda -save-temps，编译完成后会在当前目录生成后缀为.ptx的目标代码文件，可直接打开查看
  • 查看LLVM中间表示：编译时添加参数-fsycl -emit-llvm -S，会生成可读的LLVM IR格式代码
  • 拆解二进制文件的设备汇编：使用llvm-objdump工具，执行命令llvm-objdump -d 你的二进制文件名称即可输出对应设备架构的汇编代码

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