[HIP] Emit amdgpu_code_object_version module flag

[yaxunl]

Summary

code object version determines ABI, therefore should not be mixed.

This patch emits amdgpu_code_object_version module flag in LLVM IR
based on code object version (default 4).

The amdgpu_code_object_version value is code object version times 100.

LLVM IR with different amdgpu_code_object_version module flag cannot
be linked.

The -cc1 option -mcode-object-version=none is for ROCm device library use
only, which supports multiple ABI.

[tra]

clang/lib/CodeGen/CodeGenModule.cpp

575

When will it ever be set to none? Does the new option parser enforce the default & version values specified in the tablegen?
If so, then it should never be none.
If the Values specified for the option are not enforced, then the condition will be true if user specifies any invalid value other than none.

[yaxunl]

clang/lib/CodeGen/CodeGenModule.cpp

575

Normal HIP programs should only use -mcode-object-version={2|3|4|5}. clang driver enforces that.

ROCm device library need to be compiled with -Xclang -mcode-object-version=none so that the module flag is not emitted. Since this use case is not for common users, -mcode-object-version=none can only be used with -cc1.

[tra]

clang/lib/CodeGen/CodeGenModule.cpp

575

I'm surprised that cc1 does not check option validity. It does check them at the driver level: https://godbolt.org/z/9T6n47es9

Assuming that's intentional, and there are no checks in cc1, then we should probably remove the now-useless Values from the option tablegen and add a test to verify what happens with invalid values.

[yaxunl]

clang/lib/CodeGen/CodeGenModule.cpp

575

I will use MarshallingInfoEnum, which will check the values. Will add a test for invalid values.

[yaxunl]

marshalling the arg as enum. fix test failures for -cc1as. temporarily disable it except v5.

[tra]

clang/include/clang/Driver/Options.td

3453–3458

none is missing. I'm not sure if we need to enumerate all values -- the list will eventually grow unacceptably long.
Does the option parser print allowed values on error? That would be sufficient, IMO.
Up to you.

clang/lib/Driver/ToolChains/Clang.cpp

1166

Do we really need a special case for -cc1as? What happens if we do pass the option to it?

If cc1as does need a special case handling, we may want to add a test case for that.

[yaxunl]

clang/include/clang/Driver/Options.td

3453–3458

The help text is for clang driver, where 'none' is not allowed. 'none' is only allowed with clang -cc1.

clang/lib/Driver/ToolChains/Clang.cpp

1166

TargetOptions are for -cc1 only. when passed to -cc1as, they are treated as unknown options and cause error. will add a test for that.

[yaxunl]

add a test for -cc1as

[tra]

clang/lib/Driver/ToolChains/Clang.cpp

1166

does not need -mcode-object-version --> does not accept...

clang/test/CodeGenCUDA/amdgpu-code-object-version.cu

25–26

I think what we want to check is that if top-level driver decides t invoke cc1as, that we do not pass the -mcode-object-version to it.

That cc1as would not accept that option is not very interesting -- it's a CC1-only option, after all.
I guess that makes my test idea also not very interesting -- if we were to accidentally pass that option, it would cause an error.

OK, let's just update the comment above cc1as check and remove this test.

[yaxunl]

clang/lib/Driver/ToolChains/Clang.cpp

1166

will do

clang/test/CodeGenCUDA/amdgpu-code-object-version.cu

25–26

will remove this test and add a driver test for -cc1as

[yaxunl]

fix comments and add a driver test for -cc1as

clang/include/clang/Basic/TargetOptions.h

public:
 
   /// \brief If enabled, use 32-bit pointers for accessing const/local/shared
   /// address space.
   bool NVPTXUseShortPointers = false;
 
   /// \brief If enabled, allow AMDGPU unsafe floating point atomics.
   bool AllowAMDGPUUnsafeFPAtomics = false;
 
+  /// \brief Enumeration value for AMDGPU code object version, which is the
+  /// code object version times 100.
+  enum CodeObjectVersionKind {
+    COV_None,
+    COV_2 = 200,
+    COV_3 = 300,
+    COV_4 = 400,
+    COV_5 = 500,
+  };
+  /// \brief Code object version for AMDGPU.
+  CodeObjectVersionKind CodeObjectVersion;
+
   // The code model to be used as specified by the user. Corresponds to
   // CodeModel::Model enum defined in include/llvm/Support/CodeGen.h, plus
   // "default" for the case when the user has not explicitly specified a
   // code model.
   std::string CodeModel;
 
   /// The version of the SDK which was used during the compilation.
   /// The option is used for two different purposes:

clang/include/clang/Driver/Options.td

defm amdgpu_ieee : BoolOption<"m", "amdgpu-ieee",
   CodeGenOpts<"EmitIEEENaNCompliantInsts">, DefaultTrue,
   PosFlag<SetTrue, [], "Sets the IEEE bit in the expected default floating point "
   " mode register. Floating point opcodes that support exception flag "
   "gathering quiet and propagate signaling NaN inputs per IEEE 754-2008. "
   "This option changes the ABI. (AMDGPU only)">,
   NegFlag<SetFalse, [CC1Option]>>, Group<m_Group>;
 
 def mcode_object_version_EQ : Joined<["-"], "mcode-object-version=">, Group<m_Group>,
-  HelpText<"Specify code object ABI version. Defaults to 4. (AMDGPU only)">,
-  MetaVarName<"<version>">, Values<"2,3,4,5">;
+  HelpText<"Specify code object ABI version. Allowed values are 2, 3, 4, and 5. Defaults to 4. (AMDGPU only)">,
+  Flags<[CC1Option]>,
+  Values<"none,2,3,4,5">,
+  NormalizedValuesScope<"TargetOptions">,
+  NormalizedValues<["COV_None", "COV_2", "COV_3", "COV_4", "COV_5"]>,
+  MarshallingInfoEnum<TargetOpts<"CodeObjectVersion">, "COV_4">;
 
 defm code_object_v3_legacy : SimpleMFlag<"code-object-v3",
   "Legacy option to specify code object ABI V3",
   "Legacy option to specify code object ABI V2",
   " (AMDGPU only)">;
 defm cumode : SimpleMFlag<"cumode",
   "Specify CU wavefront", "Specify WGP wavefront",
   " execution mode (AMDGPU only)", m_amdgpu_Features_Group>;

clang/lib/CodeGen/CodeGenModule.cpp

if (CodeGenOpts.SanitizeCfiCrossDso) {
     CodeGenFunction(*this).EmitCfiCheckStub();
   }
   emitAtAvailableLinkGuard();
   if (Context.getTargetInfo().getTriple().isWasm() &&
       !Context.getTargetInfo().getTriple().isOSEmscripten()) {
     EmitMainVoidAlias();
   }
 
+  if (getTriple().isAMDGPU()) {
     // Emit reference of __amdgpu_device_library_preserve_asan_functions to
     // preserve ASAN functions in bitcode libraries.
-  if (LangOpts.Sanitize.has(SanitizerKind::Address) && getTriple().isAMDGPU()) {
+    if (LangOpts.Sanitize.has(SanitizerKind::Address)) {
       auto *FT = llvm::FunctionType::get(VoidTy, {});
       auto *F = llvm::Function::Create(
           FT, llvm::GlobalValue::ExternalLinkage,
           "__amdgpu_device_library_preserve_asan_functions", &getModule());
       auto *Var = new llvm::GlobalVariable(
           getModule(), FT->getPointerTo(),
           /*isConstant=*/true, llvm::GlobalValue::WeakAnyLinkage, F,
           "__amdgpu_device_library_preserve_asan_functions_ptr", nullptr,
           llvm::GlobalVariable::NotThreadLocal);
       addCompilerUsedGlobal(Var);
       if (!getModule().getModuleFlag("amdgpu_hostcall")) {
         getModule().addModuleFlag(llvm::Module::Override, "amdgpu_hostcall", 1);
       }
     }
+    // Emit amdgpu_code_object_version module flag, which is code object version
+    // times 100.
+    // ToDo: Enable module flag for all code object version when ROCm device
+    // library is ready.
+    if (getTarget().getTargetOpts().CodeObjectVersion == TargetOptions::COV_5) {
+      getModule().addModuleFlag(llvm::Module::Error,
+                                "amdgpu_code_object_version",
+                                getTarget().getTargetOpts().CodeObjectVersion);
+    }
+  }
 
   emitLLVMUsed();
   if (SanStats)
     SanStats->finish();
 
   if (CodeGenOpts.Autolink &&
       (Context.getLangOpts().Modules || !LinkerOptionsMetadata.empty())) {
     EmitModuleLinkOptions();

clang/lib/Driver/ToolChains/Clang.cpp

case llvm::Reloc::RWPI:
     return "rwpi";
   case llvm::Reloc::ROPI_RWPI:
     return "ropi-rwpi";
   }
   llvm_unreachable("Unknown Reloc::Model kind");
 }
 static void handleAMDGPUCodeObjectVersionOptions(const Driver &D,
                                                  const ArgList &Args,
-                                                 ArgStringList &CmdArgs) {
+                                                 ArgStringList &CmdArgs,
+                                                 bool IsCC1As = false) {
   // If no version was requested by the user, use the default value from the
   // back end. This is consistent with the value returned from
   // getAMDGPUCodeObjectVersion. This lets clang emit IR for amdgpu without
   // requiring the corresponding llvm to have the AMDGPU target enabled,
   // provided the user (e.g. front end tests) can use the default.
   if (haveAMDGPUCodeObjectVersionArgument(D, Args)) {
     unsigned CodeObjVer = getAMDGPUCodeObjectVersion(D, Args);
     CmdArgs.insert(CmdArgs.begin() + 1,
                    Args.MakeArgString(Twine("--amdhsa-code-object-version=") +
                                       Twine(CodeObjVer)));
     CmdArgs.insert(CmdArgs.begin() + 1, "-mllvm");
+    // -cc1as does not accept -mcode-object-version option.
+    if (!IsCC1As)
+      CmdArgs.insert(CmdArgs.begin() + 1,
+                     Args.MakeArgString(Twine("-mcode-object-version=") +
+                                        Twine(CodeObjVer)));
   }
 }
 
 void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
                                     const Driver &D, const ArgList &Args,
                                     ArgStringList &CmdArgs,
                                     const InputInfo &Output,
                                     const InputInfoList &Inputs) const {
void ClangAs::ConstructJob(Compilation &C, const JobAction &JA,
   Arg *A;
   if (getDebugFissionKind(D, Args, A) == DwarfFissionKind::Split &&
       T.isOSBinFormatELF()) {
     CmdArgs.push_back("-split-dwarf-output");
     CmdArgs.push_back(SplitDebugName(JA, Args, Input, Output));
   }
 
   if (Triple.isAMDGPU())
-    handleAMDGPUCodeObjectVersionOptions(D, Args, CmdArgs);
+    handleAMDGPUCodeObjectVersionOptions(D, Args, CmdArgs, /*IsCC1As=*/true);
 
   assert(Input.isFilename() && "Invalid input.");
   CmdArgs.push_back(Input.getFilename());
 
   const char *Exec = getToolChain().getDriver().getClangProgramPath();
   if (D.CC1Main && !D.CCGenDiagnostics) {
     // Invoke cc1as directly in this process.
     C.addCommand(std::make_unique<CC1Command>(JA, *this,

clang/test/CodeGenCUDA/amdgpu-asan-printf.cu

 // RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -target-cpu gfx906 -fsanitize=address \
 // RUN:   -O3 -x hip | FileCheck -check-prefixes=MFCHECK %s
 
-// MFCHECK: !llvm.module.flags = !{![[FLAG1:[0-9]+]], ![[FLAG2:[0-9]+]]}
-// MFCHECK: ![[FLAG1]] = !{i32 4, !"amdgpu_hostcall", i32 1}
+// MFCHECK: !{{.*}} = !{i32 4, !"amdgpu_hostcall", i32 1}
 
 // Test to check hostcall module flag metadata is generated correctly
 // when a program has printf call and compiled with -fsanitize=address.
 #include "Inputs/cuda.h"
 __device__ void non_kernel() {
   printf("sanitized device function");
 }
 
 __global__ void kernel() {
   non_kernel();
 }
 

clang/test/CodeGenCUDA/amdgpu-asan.cu

 
 // REQUIRES: amdgpu-registered-target
 
 // ASAN-DAG: define weak void @__amdgpu_device_library_preserve_asan_functions()
 // ASAN-DAG: @__amdgpu_device_library_preserve_asan_functions_ptr = weak addrspace(1) constant void ()* @__amdgpu_device_library_preserve_asan_functions
 // ASAN-DAG: @llvm.compiler.used = {{.*}}@__amdgpu_device_library_preserve_asan_functions_ptr
 // ASAN-DAG: define weak void @__asan_report_load1(i64 %{{.*}})
 
-// MFCHECK: !llvm.module.flags = !{![[FLAG1:[0-9]+]], ![[FLAG2:[0-9]+]]}
-// MFCHECK: ![[FLAG1]] = !{i32 4, !"amdgpu_hostcall", i32 1}
+// MFCHECK: !{{.*}} = !{i32 4, !"amdgpu_hostcall", i32 1}
 
 // CHECK-NOT: @__amdgpu_device_library_preserve_asan_functions
 // CHECK-NOT: @__asan_report_load1

clang/test/CodeGenCUDA/amdgpu-code-object-version.cu

+// Create module flag for code object version.
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -o - %s | FileCheck %s -check-prefix=NONE
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -mcode-object-version=2 -o - %s | FileCheck -check-prefix=NONE %s
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -mcode-object-version=3 -o - %s | FileCheck -check-prefix=NONE %s
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -mcode-object-version=4 -o - %s | FileCheck -check-prefix=NONE %s
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -mcode-object-version=5 -o - %s | FileCheck -check-prefix=V5 %s
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -mcode-object-version=none -o - %s | FileCheck %s -check-prefix=NONE
+
+// RUN: not %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -mcode-object-version=4.1 -o - %s 2>&1| FileCheck %s -check-prefix=INV
+
+// V5: !{{.*}} = !{i32 1, !"amdgpu_code_object_version", i32 500}
+// NONE-NOT: !{{.*}} = !{i32 1, !"amdgpu_code_object_version",
+// INV: error: invalid value '4.1' in '-mcode-object-version=4.1'

clang/test/Driver/hip-code-object-version.hip

 // RUN:   %s 2>&1 | FileCheck -check-prefix=V3 %s
 
 // RUN: %clang -### -target x86_64-linux-gnu \
 // RUN:   -mcode-object-version=4 -mcode-object-version=3 \
 // RUN:   --offload-arch=gfx906 -nogpulib \
 // RUN:   %s 2>&1 | FileCheck -check-prefix=V3 %s
 
 // V3-WARN: warning: argument '-mcode-object-v3' is deprecated, use '-mcode-object-version=3' instead [-Wdeprecated]
+// V3: "-mcode-object-version=3"
 // V3: "-mllvm" "--amdhsa-code-object-version=3"
 // V3: "-targets=host-x86_64-unknown-linux,hip-amdgcn-amd-amdhsa--gfx906"
 
 // Check bundle ID for code object version 4.
 
 // RUN: %clang -### -target x86_64-linux-gnu \
 // RUN:   -mcode-object-version=4 \
 // RUN:   --offload-arch=gfx906 -nogpulib \
 // RUN:   %s 2>&1 | FileCheck -check-prefix=V4 %s
 
+// V4: "-mcode-object-version=4"
 // V4: "-mllvm" "--amdhsa-code-object-version=4"
 // V4: "-targets=host-x86_64-unknown-linux,hipv4-amdgcn-amd-amdhsa--gfx906"
 
 // Check bundle ID for code object version 5.
 
 // RUN: %clang -### -target x86_64-linux-gnu \
 // RUN:   -mcode-object-version=5 \
 // RUN:   --offload-arch=gfx906 -nogpulib \
 // RUN:   %s 2>&1 | FileCheck -check-prefix=V5 %s
 
+// V5: "-mcode-object-version=5"
 // V5: "-mllvm" "--amdhsa-code-object-version=5"
 // V5: "-targets=host-x86_64-unknown-linux,hipv4-amdgcn-amd-amdhsa--gfx906"
 
 // Check bundle ID for code object version default
 
 // RUN: %clang -### -target x86_64-linux-gnu \
 // RUN:   --offload-arch=gfx906 -nogpulib \
 // RUN:   %s 2>&1 | FileCheck -check-prefix=VD %s
 
 // VD: "-targets=host-x86_64-unknown-linux,hipv4-amdgcn-amd-amdhsa--gfx906"
 
 // Check invalid code object version option.
 
 // RUN: %clang -### -target x86_64-linux-gnu \
 // RUN:   -mcode-object-version=1 \
 // RUN:   --offload-arch=gfx906 -nogpulib \
 // RUN:   %s 2>&1 | FileCheck -check-prefix=INVALID %s
 // INVALID: error: invalid integral value '1' in '-mcode-object-version=1'
 // INVALID-NOT: error: invalid integral value
 
+// Check LLVM code object version option --amdhsa-code-object-version
+// is passed to -cc1 and -cc1as, and -mcode-object-version is passed
+// to -cc1 but not -cc1as.
+
+// RUN: %clang -### -target x86_64-linux-gnu \
+// RUN:   -mcode-object-version=5 \
+// RUN:   --offload-arch=gfx906 -nogpulib -save-temps \
+// RUN:   %s 2>&1 | FileCheck -check-prefix=CC1 %s
+
+// CC1: "-cc1" {{.*}}"-mcode-object-version=5" {{.*}}"-mllvm" "--amdhsa-code-object-version=5"
+// CC1: "-cc1as" {{.*}}"-mllvm" "--amdhsa-code-object-version=5"
+
+// RUN: %clang -### -target x86_64-linux-gnu \
+// RUN:   -mcode-object-version=5 \
+// RUN:   --offload-arch=gfx906 -nogpulib -save-temps \
+// RUN:   %s 2>&1 | FileCheck -check-prefix=CC1NEG %s
+
+// CC1NEG-NOT: "-cc1as" {{.*}}"-mcode-object-version=5"
+
 // Check warnings are emitted for legacy options before -mcode-object-version options.
 // Check warnings are emitted only once.
 
 // RUN: %clang -### -target x86_64-linux-gnu \
 // RUN:   -mno-code-object-v3 -mcode-object-v3 -mcode-object-version=4 \
 // RUN:   --offload-arch=gfx906 -nogpulib \
 // RUN:   %s 2>&1 | FileCheck -check-prefixes=WARN %s
 // WARN: warning: argument '-mno-code-object-v3' is deprecated, use '-mcode-object-version=2' instead [-Wdeprecated]
 // WARN: warning: argument '-mcode-object-v3' is deprecated, use '-mcode-object-version=3' instead [-Wdeprecated]
 // WARN-NOT: warning: argument {{.*}} is deprecated