Emit OpenCL metadata when targeting SPIR-V

[shangwuyao]

Summary

This is required for converting function calls such as get_global_id()
into SPIR-V builtins.

[yaxunl]

Is this because your HIP threadIdx etc are implemented using OpenCL builtins so that the emitted LLVM IR contains calls of OpenCL builtins?

[shangwuyao]

In D123049#3426849, @yaxunl wrote:

Is this because your HIP threadIdx etc are implemented using OpenCL builtins so that the emitted LLVM IR contains calls of OpenCL builtins?

Yes, that's correct.

[yaxunl]

LGTM since currently there is only one HIP/SPIRV implementation. If in the future there is another HIP/SPIRV implementation that does not need this, it could disable it by triple.

clang/lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::Release() {
   // Indicate whether this Module was compiled with -fopenmp
   if (getLangOpts().OpenMP && !getLangOpts().OpenMPSimd)
     getModule().addModuleFlag(llvm::Module::Max, "openmp", LangOpts.OpenMP);
   if (getLangOpts().OpenMPIsDevice)
     getModule().addModuleFlag(llvm::Module::Max, "openmp-device",
                               LangOpts.OpenMP);
 
   // Emit OpenCL specific module metadata: OpenCL/SPIR version.
-  if (LangOpts.OpenCL) {
+  if (LangOpts.OpenCL || (LangOpts.CUDAIsDevice && getTriple().isSPIRV())) {
     EmitOpenCLMetadata();
     // Emit SPIR version.
     if (getTriple().isSPIR()) {
       // SPIR v2.0 s2.12 - The SPIR version used by the module is stored in the
       // opencl.spir.version named metadata.
       // C++ for OpenCL has a distinct mapping for version compatibility with
       // OpenCL.
       auto Version = LangOpts.getOpenCLCompatibleVersion();

clang/lib/Frontend/CompilerInvocation.cpp

if (Opts.HIP) {
     // device library functions in LLVM bitcode, which causes accuracy loss in
     // certain math functions, e.g. tan(-1e20) becomes -0.933 instead of 0.8446.
     // For device library functions in bitcode to work, 'Strict' or 'Standard'
     // FPOpFusion options in backends is needed. Therefore 'fast-honor-pragmas'
     // FP contract option is used to allow fuse across statements in frontend
     // whereas respecting contract flag in backend.
     Opts.setDefaultFPContractMode(LangOptions::FPM_FastHonorPragmas);
   } else if (Opts.CUDA) {
+    if (T.isSPIRV()) {
+      // Emit OpenCL version metadata in LLVM IR when targeting SPIR-V.
+      Opts.OpenCLVersion = 200;
+    }
     // Allow fuse across statements disregarding pragmas.
     Opts.setDefaultFPContractMode(LangOptions::FPM_Fast);
   }
 
   Opts.RenderScript = IK.getLanguage() == Language::RenderScript;
 
   // OpenCL, C++ and C2x have bool, true, false keywords.
   Opts.Bool = Opts.OpenCL || Opts.CPlusPlus || Opts.C2x;

clang/test/CodeGenCUDASPIRV/kernel-cc.cu

 // RUN: %clang_cc1 -fcuda-is-device -triple spirv32 -o - -emit-llvm -x cuda %s  | FileCheck %s
 // RUN: %clang_cc1 -fcuda-is-device -triple spirv64 -o - -emit-llvm -x cuda %s  | FileCheck %s
 
 // Verifies that building CUDA targeting SPIR-V {32,64} generates LLVM IR with
 // spir_kernel attributes for kernel functions.
 
 // CHECK: define spir_kernel void @_Z6kernelv()
 
 __attribute__((global)) void kernel() { return; }
+
+// CHECK: !opencl.ocl.version = !{[[OCL:![0-9]+]]}
+// CHECK: [[OCL]] = !{i32 2, i32 0}