[OpenMP] Add flag for setting debug in the offloading device

[jhuber6]

Summary

This patch introduces the flags -fopenmp-target-debug and
-fopenmp-target-debug= to set the value of a global in the device.
This will be used to enable or disable debugging features statically in
the device runtime library.

[jdoerfert]

clang/lib/Frontend/CompilerInvocation.cpp

3900

CUDANumSMs?

Don't check for NVPTX/AMDGCN but only if the new runtime is used, if not emit a warning that the flag is useless and ignored without the new runtime.

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

254

Linkage cannot be private, I think. It might be ok if we put it in the llvm.used global otherwise we should just make it external, which is probably best.

[jhuber6]

Addressing Comments.

[jhuber6]

Forgot to format.

[jhuber6]

Adding global regex to test.

[jhuber6]

Adding constant to llvm.used. This is most likely easier than dealing with weak external linkage.

[jdoerfert]

LG

[mgorny]

This change causes clang-ast-dump to segfault for me while building clang:

[399/399] ASTNodeAPI.json
FAILED: tools/clang/lib/Tooling/ASTNodeAPI.json /home/mgorny/git/llvm-project/build/tools/clang/lib/Tooling/ASTNodeAPI.json 
cd /home/mgorny/git/llvm-project/build/tools/clang/lib/Tooling && /home/mgorny/git/llvm-project/build/bin/clang-ast-dump --skip-processing=0 -I /home/mgorny/git/llvm-project/build/lib/clang/14.0.0/include -I /home/mgorny/git/llvm-project/llvm/../clang/include -I /home/mgorny/git/llvm-project/build/tools/clang/include -I /home/mgorny/git/llvm-project/build/include -I /home/mgorny/git/llvm-project/llvm/include -I /usr/lib/gcc/x86_64-pc-linux-gnu/11.2.0/include/g++-v11 -I /usr/lib/gcc/x86_64-pc-linux-gnu/11.2.0/include/g++-v11/x86_64-pc-linux-gnu -I /usr/lib/gcc/x86_64-pc-linux-gnu/11.2.0/include/g++-v11/backward -I /usr/lib/gcc/x86_64-pc-linux-gnu/11.2.0/include -I /usr/lib/gcc/x86_64-pc-linux-gnu/11.2.0/include-fixed -I /usr/include --json-output-path /home/mgorny/git/llvm-project/build/tools/clang/lib/Tooling/ASTNodeAPI.json
ninja: build stopped: subcommand failed.

Reading symbols from /home/mgorny/git/llvm-project/build/bin/clang-ast-dump...
[New LWP 230743]
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib64/libthread_db.so.1".
Core was generated by `/home/mgorny/git/llvm-project/build/bin/clang-ast-dump --skip-processing=0 -I /'.
Program terminated with signal SIGSEGV, Segmentation fault.
#0  0x00007ff222268a87 in llvm::BumpPtrAllocatorImpl<llvm::MallocAllocator, 4096ul, 4096ul, 128ul>::~BumpPtrAllocatorImpl() ()
   from /home/mgorny/git/llvm-project/build/bin/../lib/../lib/libclangLex.so.14git
(gdb) bt
#0  0x00007ff222268a87 in llvm::BumpPtrAllocatorImpl<llvm::MallocAllocator, 4096ul, 4096ul, 128ul>::~BumpPtrAllocatorImpl() ()
   from /home/mgorny/git/llvm-project/build/bin/../lib/../lib/libclangLex.so.14git
#1  0x00007ff2222a1404 in clang::HeaderSearch::~HeaderSearch() ()
   from /home/mgorny/git/llvm-project/build/bin/../lib/../lib/libclangLex.so.14git
#2  0x00007ff2222a1859 in clang::Preprocessor::~Preprocessor() ()
   from /home/mgorny/git/llvm-project/build/bin/../lib/../lib/libclangLex.so.14git
#3  0x00007ff224d3175c in std::__shared_count<(__gnu_cxx::_Lock_policy)2>::~__shared_count() ()
   from /home/mgorny/git/llvm-project/build/bin/../lib/libclangFrontend.so.14git
#4  0x00007ff224d4c801 in clang::CompilerInstance::~CompilerInstance() ()
   from /home/mgorny/git/llvm-project/build/bin/../lib/libclangFrontend.so.14git
#5  0x0000556c688f8a95 in main ()

I'm going to try RelWithDebInfo build but I'm afraid I don't have sufficient free space for that.

[mgorny]

I'm sorry, it seems to have been caused by ccache. After clearing the cache, I can't reproduce anymore.

[jdoerfert]

In D109544#2996913, @mgorny wrote:

I'm sorry, it seems to have been caused by ccache. After clearing the cache, I can't reproduce anymore.

The problem is that clang-ast-dump has incomplete dependences and sometimes picks stuff up from the install folder.
If you delete install it works, ccache might have a similar effect.

clang/include/clang/Basic/DiagnosticDriverKinds.td

 def err_drv_emit_llvm_link : Error<
    "-emit-llvm cannot be used when linking">;
 def err_drv_optimization_remark_pattern : Error<
   "in pattern '%1': %0">;
 def err_drv_optimization_remark_format : Error<
   "unknown remark serializer format: '%0'">;
 def err_drv_no_neon_modifier : Error<"[no]neon is not accepted as modifier, please use [no]simd instead">;
 def err_drv_invalid_omp_target : Error<"OpenMP target is invalid: '%0'">;
+def err_drv_debug_no_new_runtime : Error<"OpenMP target device debugging enabled with incompatible runtime">;
 def err_drv_incompatible_omp_arch : Error<"OpenMP target architecture '%0' pointer size is incompatible with host '%1'">;
 def err_drv_omp_host_ir_file_not_found : Error<
   "provided host compiler IR file '%0' is required to generate code for OpenMP "
   "target regions but cannot be found">;
 def err_drv_omp_host_target_not_supported : Error<
   "target '%0' is not a supported OpenMP host target">;
 def err_drv_expecting_fopenmp_with_fopenmp_targets : Error<
   "'-fopenmp-targets' must be used in conjunction with a '-fopenmp' option "

clang/include/clang/Basic/LangOptions.def

 LANGOPT(OpenMPIsDevice    , 1, 0, "Generate code only for OpenMP target device")
 LANGOPT(OpenMPCUDAMode    , 1, 0, "Generate code for OpenMP pragmas in SIMT/SPMD mode")
 LANGOPT(OpenMPIRBuilder   , 1, 0, "Use the experimental OpenMP-IR-Builder codegen path.")
 LANGOPT(OpenMPCUDAForceFullRuntime , 1, 0, "Force to use full runtime in all constructs when offloading to CUDA devices")
 LANGOPT(OpenMPCUDANumSMs  , 32, 0, "Number of SMs for CUDA devices.")
 LANGOPT(OpenMPCUDABlocksPerSM  , 32, 0, "Number of blocks per SM for CUDA devices.")
 LANGOPT(OpenMPCUDAReductionBufNum , 32, 1024, "Number of the reduction records in the intermediate reduction buffer used for the teams reductions.")
 LANGOPT(OpenMPTargetNewRuntime , 1, 0, "Use the new bitcode library for OpenMP offloading")
+LANGOPT(OpenMPTargetDebug , 32, 0, "Enable debugging in the OpenMP offloading device RTL")
 LANGOPT(OpenMPOptimisticCollapse  , 1, 0, "Use at most 32 bits to represent the collapsed loop nest counter.")
 LANGOPT(RenderScript      , 1, 0, "RenderScript")
 
 LANGOPT(CUDAIsDevice      , 1, 0, "compiling for CUDA device")
 LANGOPT(CUDAAllowVariadicFunctions, 1, 0, "allowing variadic functions in CUDA device code")
 LANGOPT(CUDAHostDeviceConstexpr, 1, 1, "treating unattributed constexpr functions as __host__ __device__")
 LANGOPT(CUDADeviceApproxTranscendentals, 1, 0, "using approximate transcendental functions")
 LANGOPT(GPURelocatableDeviceCode, 1, 0, "generate relocatable device code")

clang/include/clang/Driver/Options.td

 def fno_openmp_cuda_force_full_runtime : Flag<["-"], "fno-openmp-cuda-force-full-runtime">, Group<f_Group>,
   Flags<[NoArgumentUnused, HelpHidden]>;
 def fopenmp_cuda_number_of_sm_EQ : Joined<["-"], "fopenmp-cuda-number-of-sm=">, Group<f_Group>,
   Flags<[CC1Option, NoArgumentUnused, HelpHidden]>;
 def fopenmp_cuda_blocks_per_sm_EQ : Joined<["-"], "fopenmp-cuda-blocks-per-sm=">, Group<f_Group>,
   Flags<[CC1Option, NoArgumentUnused, HelpHidden]>;
 def fopenmp_cuda_teams_reduction_recs_num_EQ : Joined<["-"], "fopenmp-cuda-teams-reduction-recs-num=">, Group<f_Group>,
   Flags<[CC1Option, NoArgumentUnused, HelpHidden]>;
+def fopenmp_target_debug : Flag<["-"], "fopenmp-target-debug">, Group<f_Group>, Flags<[CC1Option, NoArgumentUnused]>,
+  HelpText<"Enable debugging in the OpenMP offloading device RTL">;
+def fno_openmp_target_debug : Flag<["-"], "fno-openmp-target-debug">, Group<f_Group>, Flags<[NoArgumentUnused]>;
+def fopenmp_target_debug_EQ : Joined<["-"], "fopenmp-target-debug=">, Group<f_Group>, Flags<[CC1Option, NoArgumentUnused, HelpHidden]>;
 defm openmp_target_new_runtime: BoolFOption<"openmp-target-new-runtime",
   LangOpts<"OpenMPTargetNewRuntime">, DefaultFalse,
   PosFlag<SetTrue, [CC1Option], "Use the new bitcode library for OpenMP offloading">,
   NegFlag<SetFalse>>;
 defm openmp_optimistic_collapse : BoolFOption<"openmp-optimistic-collapse",
   LangOpts<"OpenMPOptimisticCollapse">, DefaultFalse,
   PosFlag<SetTrue, [CC1Option]>, NegFlag<SetFalse>, BothFlags<[NoArgumentUnused, HelpHidden]>>;
 def static_openmp: Flag<["-"], "static-openmp">,

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

unsigned CGOpenMPRuntimeGPU::getDefaultLocationReserved2Flags() const {
   }
   llvm_unreachable("Unknown flags are requested.");
 }
 
 CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
     : CGOpenMPRuntime(CGM, "_", "$") {
   if (!CGM.getLangOpts().OpenMPIsDevice)
     llvm_unreachable("OpenMP NVPTX can only handle device code.");
+
+  llvm::OpenMPIRBuilder &OMPBuilder = getOMPBuilder();
+  if (CGM.getLangOpts().OpenMPTargetNewRuntime)
+    OMPBuilder.createDebugKind(CGM.getLangOpts().OpenMPTargetDebug);
 }
 
 void CGOpenMPRuntimeGPU::emitProcBindClause(CodeGenFunction &CGF,
                                               ProcBindKind ProcBind,
                                               SourceLocation Loc) {
   // Do nothing in case of SPMD mode and L0 parallel.
   if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
     return;

clang/lib/Driver/ToolChains/Clang.cpp

case Driver::OMPRT_IOMP5:
         CmdArgs.push_back("-fopenmp-optimistic-collapse");
 
       // When in OpenMP offloading mode with NVPTX target, forward
       // cuda-mode flag
       if (Args.hasFlag(options::OPT_fopenmp_cuda_mode,
                        options::OPT_fno_openmp_cuda_mode, /*Default=*/false))
         CmdArgs.push_back("-fopenmp-cuda-mode");
 
+      // When in OpenMP offloading mode, enable or disable the new device
+      // runtime.
+      if (Args.hasFlag(options::OPT_fopenmp_target_new_runtime,
+                       options::OPT_fno_openmp_target_new_runtime,
+                       /*Default=*/false))
+        CmdArgs.push_back("-fopenmp-target-new-runtime");
+
+      // When in OpenMP offloading mode, enable debugging on the device.
+      Args.AddAllArgs(CmdArgs, options::OPT_fopenmp_target_debug_EQ);
+      if (Args.hasFlag(options::OPT_fopenmp_target_debug,
+                       options::OPT_fno_openmp_target_debug, /*Default=*/false))
+        CmdArgs.push_back("-fopenmp-target-debug");
+
       // When in OpenMP offloading mode with NVPTX target, check if full runtime
       // is required.
       if (Args.hasFlag(options::OPT_fopenmp_cuda_force_full_runtime,
                        options::OPT_fno_openmp_cuda_force_full_runtime,
                        /*Default=*/false))
         CmdArgs.push_back("-fopenmp-cuda-force-full-runtime");
       break;
     default:

clang/lib/Frontend/CompilerInvocation.cpp

#undef LANG_OPTION_WITH_MARSHALLING
 
   if (Opts.OpenMPSimd) {
     GenerateArg(Args, OPT_fopenmp_simd, SA);
 
     if (Opts.OpenMP != 50)
       GenerateArg(Args, OPT_fopenmp_version_EQ, Twine(Opts.OpenMP), SA);
   }
 
+  if (Opts.OpenMPTargetNewRuntime)
+    GenerateArg(Args, OPT_fopenmp_target_new_runtime, SA);
+
+  if (Opts.OpenMPTargetDebug != 0)
+    GenerateArg(Args, OPT_fopenmp_target_debug_EQ,
+                Twine(Opts.OpenMPTargetDebug), SA);
+
   if (Opts.OpenMPCUDANumSMs != 0)
     GenerateArg(Args, OPT_fopenmp_cuda_number_of_sm_EQ,
                 Twine(Opts.OpenMPCUDANumSMs), SA);
 
   if (Opts.OpenMPCUDABlocksPerSM != 0)
     GenerateArg(Args, OPT_fopenmp_cuda_blocks_per_sm_EQ,
                 Twine(Opts.OpenMPCUDABlocksPerSM), SA);
 
#undef LANG_OPTION_WITH_MARSHALLING
   Opts.OpenMPUseTLS =
       Opts.OpenMP && !Args.hasArg(options::OPT_fnoopenmp_use_tls);
   Opts.OpenMPIsDevice =
       Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_is_device);
   Opts.OpenMPIRBuilder =
       Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_enable_irbuilder);
   bool IsTargetSpecified =
       Opts.OpenMPIsDevice || Args.hasArg(options::OPT_fopenmp_targets_EQ);
+  Opts.OpenMPTargetNewRuntime = 
+      Opts.OpenMPIsDevice && Args.hasArg(options::OPT_fopenmp_target_new_runtime);
 
   Opts.ConvergentFunctions = Opts.ConvergentFunctions || Opts.OpenMPIsDevice;
 
   if (Opts.OpenMP || Opts.OpenMPSimd) {
     if (int Version = getLastArgIntValue(
             Args, OPT_fopenmp_version_EQ,
             (IsSimdSpecified || IsTargetSpecified) ? 50 : Opts.OpenMP, Diags))
       Opts.OpenMP = Version;
if (!Opts.OpenMPIsDevice) {
       }
     }
   }
 
   // Set the flag to prevent the implementation from emitting device exception
   // handling code for those requiring so.
   if ((Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN())) ||
       Opts.OpenCLCPlusPlus) {
+
     Opts.Exceptions = 0;
     Opts.CXXExceptions = 0;
   }
   if (Opts.OpenMPIsDevice && T.isNVPTX()) {
     Opts.OpenMPCUDANumSMs =
         getLastArgIntValue(Args, options::OPT_fopenmp_cuda_number_of_sm_EQ,
                            Opts.OpenMPCUDANumSMs, Diags);
     Opts.OpenMPCUDABlocksPerSM =
         getLastArgIntValue(Args, options::OPT_fopenmp_cuda_blocks_per_sm_EQ,
                            Opts.OpenMPCUDABlocksPerSM, Diags);
     Opts.OpenMPCUDAReductionBufNum = getLastArgIntValue(
         Args, options::OPT_fopenmp_cuda_teams_reduction_recs_num_EQ,
         Opts.OpenMPCUDAReductionBufNum, Diags);
   }
 
+  // Set the value of the debugging flag used in the new offloading device RTL.
+  // Set either by a specific value or to a default if not specified.
+  if (Opts.OpenMPIsDevice && (Args.hasArg(OPT_fopenmp_target_debug) ||
+                              Args.hasArg(OPT_fopenmp_target_debug_EQ))) {
+    if (Opts.OpenMPTargetNewRuntime) {
+      Opts.OpenMPTargetDebug = getLastArgIntValue(
+          Args, OPT_fopenmp_target_debug_EQ, Opts.OpenMPTargetDebug, Diags);
+      if (!Opts.OpenMPTargetDebug && Args.hasArg(OPT_fopenmp_target_debug))
+        Opts.OpenMPTargetDebug = 1;
+    } else {
+      Diags.Report(diag::err_drv_debug_no_new_runtime);
+    }
+  }
+
   // Get the OpenMP target triples if any.
   if (Arg *A = Args.getLastArg(options::OPT_fopenmp_targets_EQ)) {
     enum ArchPtrSize { Arch16Bit, Arch32Bit, Arch64Bit };
     auto getArchPtrSize = [](const llvm::Triple &T) {
       if (T.isArch16Bit())
         return Arch16Bit;
       if (T.isArch32Bit())
         return Arch32Bit;

clang/test/OpenMP/target_debug_codegen.cpp

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+"
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-target-debug -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-target-debug=111 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-EQ
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+//.
+// CHECK: @__omp_rtl_debug_kind = private constant i32 1
+// CHECK: @0 = private unnamed_addr constant [23 x i8] c"
+// CHECK: @1 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8
+// CHECK: @__omp_offloading_fd02_613a0c37__Z3foov_l25_exec_mode = weak constant i8 1
+// CHECK: @llvm.compiler.used = appending global [1 x i8*] [i8* @__omp_offloading_fd02_613a0c37__Z3foov_l25_exec_mode], section "llvm.metadata"
+//.
+// CHECK-EQ: @__omp_rtl_debug_kind = private constant i32 111
+// CHECK-EQ: @0 = private unnamed_addr constant [23 x i8] c"
+// CHECK-EQ: @1 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8
+// CHECK-EQ: @__omp_offloading_fd02_613a0c37__Z3foov_l25_exec_mode = weak constant i8 1
+// CHECK-EQ: @llvm.compiler.used = appending global [1 x i8*] [i8* @__omp_offloading_fd02_613a0c37__Z3foov_l25_exec_mode], section "llvm.metadata"
+//.
+void foo() {
+#pragma omp target
+  { }
+}
+
+#endif
+//
+//
+//
+//.
+// CHECK: attributes #0 = { convergent noinline norecurse nounwind optnone "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" }
+//.
+// CHECK-EQ: attributes #0 = { convergent noinline norecurse nounwind optnone "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" }
+//.
+// CHECK: !0 = !{i32 0, i32 64770, i32 1631194167, !"_Z3foov", i32 25, i32 0}
+// CHECK: !1 = !{void ()* @__omp_offloading_fd02_613a0c37__Z3foov_l25, !"kernel", i32 1}
+// CHECK: !2 = !{i32 1, !"wchar_size", i32 4}
+// CHECK: !3 = !{i32 7, !"openmp", i32 50}
+// CHECK: !4 = !{i32 7, !"openmp-device", i32 50}
+// CHECK: !5 = !{!"clang version 14.0.0"}
+//.
+// CHECK-EQ: !0 = !{i32 0, i32 64770, i32 1631194167, !"_Z3foov", i32 25, i32 0}
+// CHECK-EQ: !1 = !{void ()* @__omp_offloading_fd02_613a0c37__Z3foov_l25, !"kernel", i32 1}
+// CHECK-EQ: !2 = !{i32 1, !"wchar_size", i32 4}
+// CHECK-EQ: !3 = !{i32 7, !"openmp", i32 50}
+// CHECK-EQ: !4 = !{i32 7, !"openmp-device", i32 50}
+// CHECK-EQ: !5 = !{!"clang version 14.0.0"}
+//.

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

public:
   Constant *getOrCreateSrcLocStr(const LocationDescription &Loc);
 
   /// Return an ident_t* encoding the source location \p SrcLocStr and \p Flags.
   /// TODO: Create a enum class for the Reserve2Flags
   Value *getOrCreateIdent(Constant *SrcLocStr,
                           omp::IdentFlag Flags = omp::IdentFlag(0),
                           unsigned Reserve2Flags = 0);
 
+  /// Create a global value containing the \p DebugLevel to control debuggin in
+  /// the module.
+  GlobalValue *createDebugKind(unsigned DebugLevel);
+
   /// Generate control flow and cleanup for cancellation.
   ///
   /// \param CancelFlag Flag indicating if the cancellation is performed.
   /// \param CanceledDirective The kind of directive that is cancled.
   /// \param ExitCB Extra code to be generated in the exit block.
   void emitCancelationCheckImpl(Value *CancelFlag,
                                 omp::Directive CanceledDirective,
                                 FinalizeCallbackTy ExitCB = {});

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

void OpenMPIRBuilder::finalize(Function *Fn, bool AllowExtractorSinking) {
   // Remove work items that have been completed.
   OutlineInfos = std::move(DeferredOutlines);
 }
 
 OpenMPIRBuilder::~OpenMPIRBuilder() {
   assert(OutlineInfos.empty() && "There must be no outstanding outlinings");
 }
 
+GlobalValue *OpenMPIRBuilder::createDebugKind(unsigned DebugKind) {
+  IntegerType *I32Ty = Type::getInt32Ty(M.getContext());
+  auto *GV = new GlobalVariable(
+      M, I32Ty,
+      /* isConstant = */ true, GlobalValue::PrivateLinkage,
+      ConstantInt::get(I32Ty, DebugKind), "__omp_rtl_debug_kind");
+
+  return GV;
+}
+
 Value *OpenMPIRBuilder::getOrCreateIdent(Constant *SrcLocStr,
                                          IdentFlag LocFlags,
                                          unsigned Reserve2Flags) {
   // Enable "C-mode".
   LocFlags |= OMP_IDENT_FLAG_KMPC;
 
   Value *&Ident =
       IdentMap[{SrcLocStr, uint64_t(LocFlags) << 31 | Reserve2Flags}];