// Copyright 2021 The IREE Authors
//
// Licensed under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

#include "compiler/plugins/target/LLVMCPU/LibraryBuilder.h"

#include "llvm/IR/DerivedTypes.h"
#include "llvm/IR/IRBuilder.h"

// =============================================================================
//
// NOTE: these structures model 1:1 those in iree/hal/local/executable_library.h
//
// This file must always track the latest version. Backwards compatibility with
// existing runtimes using older versions of the header is maintained by
// emitting variants of the structs matching those in the older headers and
// selecting between them in the query function based on the requested version.
//
// =============================================================================

namespace mlir::iree_compiler::IREE::HAL {

static inline int64_t roundUpToAlignment(int64_t value, int64_t alignment) {
  return (value + (alignment - 1)) & ~(alignment - 1);
}

//===----------------------------------------------------------------------===//
// iree/hal/local/executable_library.h structure types
//===----------------------------------------------------------------------===//
// The IR snippets below were pulled from clang running with `-S -emit-llvm`
// on the executable_library.h header: https://godbolt.org/z/6bMv5jfvf

// %struct.iree_hal_executable_import_table_v0_t = type {
//   i32,
//   i8**
// }
static llvm::StructType *makeImportTableType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_import_table_v0_t")) {
    return existingType;
  }
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *type = llvm::StructType::create(context,
                                        {
                                            i32Type,
                                            llvm::PointerType::get(context, 0),
                                        },
                                        "iree_hal_executable_import_table_v0_t",
                                        /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_environment_v0_t = type {
//   ...
// }
static llvm::StructType *makeEnvironmentType(llvm::LLVMContext &context) {
  auto *type = llvm::StructType::getTypeByName(
      context, "iree_hal_executable_environment_v0_t");
  if (!type) {
    type = llvm::StructType::create(context,
                                    "iree_hal_executable_environment_v0_t");
  }
  return type;
}

// %struct.iree_hal_executable_dispatch_state_v0_t = type {
//   ...
// }
static llvm::StructType *makeDispatchStateType(llvm::LLVMContext &context) {
  auto *type = llvm::StructType::getTypeByName(
      context, "iree_hal_executable_dispatch_state_v0_t");
  if (!type) {
    type = llvm::StructType::create(context,
                                    "iree_hal_executable_dispatch_state_v0_t");
  }
  return type;
}

// %struct.iree_hal_executable_workgroup_state_v0_t = type {
//   ...
// }
static llvm::StructType *makeWorkgroupStateType(llvm::LLVMContext &context) {
  auto *type = llvm::StructType::getTypeByName(
      context, "iree_hal_executable_workgroup_state_v0_t");
  if (!type) {
    type = llvm::StructType::create(context,
                                    "iree_hal_executable_workgroup_state_v0_t");
  }
  return type;
}

// i32 (%struct.iree_hal_executable_environment_v0_t*,
//      %struct.iree_hal_executable_dispatch_state_v0_t*,
//      i8*)
static llvm::FunctionType *
makeDispatchFunctionType(llvm::LLVMContext &context) {
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *ptrType = llvm::PointerType::get(context, 0);
  return llvm::FunctionType::get(i32Type,
                                 {
                                     ptrType,
                                     ptrType,
                                     ptrType,
                                 },
                                 /*isVarArg=*/false);
}

// %struct.iree_hal_executable_dispatch_attrs_v0_t = type {
//   i64,     // flags
//   i16,     // local_memory_pages
//   i8,      // constant_count
//   i8,      // binding_count
//   i32,     // workgroup_size_x
//   i32,     // workgroup_size_y
//   i16,     // workgroup_size_z
//   i16,     // reserved_0
//   i64[5]   // reserved_1
// }
static llvm::StructType *makeDispatchAttrsType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_dispatch_attrs_v0_t")) {
    return existingType;
  }
  auto *i8Type = llvm::IntegerType::getInt8Ty(context);
  auto *i16Type = llvm::IntegerType::getInt16Ty(context);
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *i64Type = llvm::IntegerType::getInt64Ty(context);
  auto *type =
      llvm::StructType::create(context,
                               {
                                   i64Type, // flags
                                   i16Type, // local_memory_pages
                                   i8Type,  // constant_count
                                   i8Type,  // binding_count
                                   i32Type, // workgroup_size_x
                                   i32Type, // workgroup_size_y
                                   i16Type, // workgroup_size_z
                                   i16Type, // parameter_count
                                   i64Type, // reserved_1[0]
                                   i64Type, // reserved_1[1]
                                   i64Type, // reserved_1[2]
                                   i64Type, // reserved_1[3]
                                   i64Type, // reserved_1[4]
                               },
                               "iree_hal_executable_dispatch_attrs_v0_t",
                               /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_source_location_v0_t = type {
//   i32,
//   i32,
//   i8*
// }
static llvm::StructType *makeSourceLocationType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_source_location_v0_t")) {
    return existingType;
  }
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *i8PtrType = llvm::PointerType::getUnqual(context);
  auto *type =
      llvm::StructType::create(context,
                               {
                                   i32Type,
                                   i32Type,
                                   i8PtrType,
                               },
                               "iree_hal_executable_source_location_v0_t",
                               /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_stage_location_table_v0_t = type {
//   i32,
//   i8**,
//   %struct.iree_hal_executable_source_location_v0_t*,
// }
static llvm::StructType *
makeStageLocationTableType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_stage_location_table_v0_t")) {
    return existingType;
  }
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *type =
      llvm::StructType::create(context,
                               {
                                   i32Type,
                                   llvm::PointerType::get(context, 0),
                                   llvm::PointerType::get(context, 0),
                               },
                               "iree_hal_executable_stage_location_table_v0_t",
                               /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_dispatch_parameter_v0_t = type {
//   i8,
//   i8,
//   i16,
//   i16,
//   i16
// }
static llvm::StructType *makeDispatchParameterType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_dispatch_parameter_v0_t")) {
    return existingType;
  }
  auto *i8Type = llvm::IntegerType::getInt8Ty(context);
  auto *i16Type = llvm::IntegerType::getInt16Ty(context);
  auto *type =
      llvm::StructType::create(context,
                               {
                                   i8Type,  // type
                                   i8Type,  // size
                                   i16Type, // flags
                                   i16Type, // name
                                   i16Type, // offset
                               },
                               "iree_hal_executable_dispatch_parameter_v0_t",
                               /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_dispatch_occupancy_v0_t = type {
//   i32
// }
static llvm::StructType *makeDispatchOccupancyType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_dispatch_occupancy_v0_t")) {
    return existingType;
  }
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *type =
      llvm::StructType::create(context,
                               {
                                   i32Type, // reserved
                               },
                               "iree_hal_executable_dispatch_occupancy_v0_t",
                               /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_export_table_v0_t = type {
//   i32,
//   i32*,
//   %struct.iree_hal_executable_dispatch_attrs_v0_t*,
//   %struct.iree_hal_executable_dispatch_parameter_v0_t**,
//   %struct.iree_hal_executable_dispatch_occupancy_v0_t*,
//   i8**,
//   i8**,
//   i8**,
//   %struct.iree_hal_executable_source_location_v0_t*,
//   %struct.iree_hal_executable_stage_location_table_v0_t*,
// }
static llvm::StructType *makeExportTableType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_export_table_v0_t")) {
    return existingType;
  }
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *ptrType = llvm::PointerType::get(context, 0);
  auto *type = llvm::StructType::create(context,
                                        {
                                            i32Type, // count
                                            ptrType, // ptrs
                                            ptrType, // attrs
                                            ptrType, // params
                                            ptrType, // occupancy
                                            ptrType, // names
                                            ptrType, // tags
                                            ptrType, // parameter_names
                                            ptrType, // source_locations
                                            ptrType, // stage_locations
                                        },
                                        "iree_hal_executable_export_table_v0_t",
                                        /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_constant_table_v0_t = type {
//   i32
// }
static llvm::StructType *makeConstantTableType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_constant_table_v0_t")) {
    return existingType;
  }
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *type =
      llvm::StructType::create(context,
                               {
                                   i32Type,
                               },
                               "iree_hal_executable_constant_table_v0_t",
                               /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_source_file_v0_t = type {
//   i32,
//   i8*,
//   i32,
//   i8*
// }
static llvm::StructType *makeSourceFileType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_source_file_v0_t")) {
    return existingType;
  }
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *i8PtrType = llvm::PointerType::getUnqual(context);
  auto *type = llvm::StructType::create(context,
                                        {
                                            i32Type,
                                            i8PtrType,
                                            i32Type,
                                            i8PtrType,
                                        },
                                        "iree_hal_executable_source_file_v0_t",
                                        /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_source_file_table_v0_t = type {
//   i32,
//   %struct.iree_hal_executable_source_file_v0_t*,
// }
static llvm::StructType *makeSourceTableType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_source_file_table_v0_t")) {
    return existingType;
  }
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *type =
      llvm::StructType::create(context,
                               {
                                   i32Type,
                                   llvm::PointerType::get(context, 0),
                               },
                               "iree_hal_executable_source_file_table_v0_t",
                               /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_library_header_t = type {
//   i32,
//   i8*,
//   i32,
//   i32
// }
static llvm::StructType *makeLibraryHeaderType(llvm::LLVMContext &context) {
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_library_header_t")) {
    return existingType;
  }
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *i8PtrType = llvm::PointerType::getUnqual(context);
  auto *type = llvm::StructType::create(context,
                                        {
                                            i32Type,
                                            i8PtrType,
                                            i32Type,
                                            i32Type,
                                        },
                                        "iree_hal_executable_library_header_t",
                                        /*isPacked=*/false);
  return type;
}

// %struct.iree_hal_executable_library_v0_t = type {
//   %struct.iree_hal_executable_library_header_t*,
//   %struct.iree_hal_executable_import_table_v0_t,
//   %struct.iree_hal_executable_export_table_v0_t,
// }
static llvm::StructType *makeLibraryType(llvm::StructType *libraryHeaderType) {
  auto &context = libraryHeaderType->getContext();
  if (auto *existingType = llvm::StructType::getTypeByName(
          context, "iree_hal_executable_library_v0_t")) {
    return existingType;
  }
  auto *importTableType = makeImportTableType(context);
  auto *exportTableType = makeExportTableType(context);
  auto *constantTableType = makeConstantTableType(context);
  auto *sourceTableType = makeSourceTableType(context);
  auto *type = llvm::StructType::create(context,
                                        {
                                            llvm::PointerType::get(context, 0),
                                            importTableType,
                                            exportTableType,
                                            constantTableType,
                                            sourceTableType,
                                        },
                                        "iree_hal_executable_library_v0_t",
                                        /*isPacked=*/false);
  return type;
}

//===----------------------------------------------------------------------===//
// IR construction utilities
//===----------------------------------------------------------------------===//

// Creates a global NUL-terminated string constant.
//
// Example:
//   @.str.2 = private unnamed_addr constant [6 x i8] c"lib_a\00", align 1
static llvm::Constant *createStringConstant(StringRef value,
                                            llvm::Module *module) {
  auto i8Type = llvm::IntegerType::getInt8Ty(module->getContext());
  auto i32Type = llvm::IntegerType::getInt32Ty(module->getContext());
  auto *stringType = llvm::ArrayType::get(i8Type, value.size() + /*NUL*/ 1);
  auto *literal =
      llvm::ConstantDataArray::getString(module->getContext(), value);
  auto *global = new llvm::GlobalVariable(*module, stringType,
                                          /*isConstant=*/true,
                                          llvm::GlobalVariable::PrivateLinkage,
                                          literal, /*Name=*/"");
  global->setAlignment(llvm::MaybeAlign(1));
  llvm::Constant *zero = llvm::ConstantInt::get(i32Type, 0);
  return llvm::ConstantExpr::getInBoundsGetElementPtr(
      stringType, global, ArrayRef<llvm::Constant *>{zero, zero});
}

// Creates a global NUL-terminated string constant or NULL if the string is
// empty.
static llvm::Constant *createStringConstantOrNull(StringRef value,
                                                  llvm::Module *module) {
  if (value.empty()) {
    return llvm::ConstantPointerNull::get(
        llvm::PointerType::get(module->getContext(), 0));
  }
  return createStringConstant(value, module);
}

// Creates a global serialized buffer constant (or string without NUL).
//
// Example:
//   @.data = private unnamed_addr constant [5 x i8] c"lib_a", align 1
static llvm::Constant *createBufferConstant(StringRef name,
                                            ArrayRef<char> value,
                                            llvm::Module *module) {
  auto i8Type = llvm::IntegerType::getInt8Ty(module->getContext());
  auto i32Type = llvm::IntegerType::getInt32Ty(module->getContext());
  auto *bufferType = llvm::ArrayType::get(i8Type, value.size());
  auto *literal = llvm::ConstantDataArray::get(module->getContext(), value);
  auto *global = new llvm::GlobalVariable(
      *module, bufferType,
      /*isConstant=*/true, llvm::GlobalVariable::PrivateLinkage, literal, name);
  global->setAlignment(llvm::MaybeAlign(1));
  llvm::Constant *zero = llvm::ConstantInt::get(i32Type, 0);
  return llvm::ConstantExpr::getInBoundsGetElementPtr(
      bufferType, global, ArrayRef<llvm::Constant *>{zero, zero});
}

// Creates a global constant with the given elements.
static llvm::Constant *createArrayConstant(StringRef name,
                                           llvm::Type *elementType,
                                           ArrayRef<llvm::Constant *> elements,
                                           llvm::Module *module) {
  auto *i32Type = llvm::IntegerType::getInt32Ty(module->getContext());
  llvm::Constant *zero = llvm::ConstantInt::get(i32Type, 0);
  auto *arrayType = llvm::ArrayType::get(elementType, elements.size());
  auto *global = new llvm::GlobalVariable(
      *module, arrayType, /*isConstant=*/true,
      llvm::GlobalVariable::PrivateLinkage,
      llvm::ConstantArray::get(arrayType, elements), name);
  return llvm::ConstantExpr::getInBoundsGetElementPtr(
      arrayType, global, ArrayRef<llvm::Constant *>{zero, zero});
}

//===----------------------------------------------------------------------===//
// Builder interface
//===----------------------------------------------------------------------===//

llvm::Function *LibraryBuilder::build(StringRef queryFuncName) {
  auto &context = module->getContext();
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *ptrType = llvm::PointerType::get(context, 0);

  // %struct.iree_hal_executable_library_header_t**
  // @iree_hal_library_query(i32, %struct.iree_hal_executable_environment_v0_t*)
  auto *queryFuncType = llvm::FunctionType::get(ptrType,
                                                {
                                                    i32Type,
                                                    ptrType,
                                                },
                                                /*isVarArg=*/false);
  auto *func =
      llvm::Function::Create(queryFuncType, llvm::GlobalValue::InternalLinkage,
                             queryFuncName, *module);

  auto *entryBlock = llvm::BasicBlock::Create(context, "entry", func);
  llvm::IRBuilder<> builder(entryBlock);

  // Build out the header for each version and select it at runtime.
  // NOTE: today there is just one version so this is rather simple:
  //   return max_version == 0 ? &library : NULL;
  auto *v0 = buildLibraryV0((queryFuncName + "_v0").str());
  builder.CreateRet(builder.CreateSelect(
      builder.CreateICmpEQ(func->getArg(0),
                           llvm::ConstantInt::get(
                               i32Type, static_cast<int64_t>(Version::LATEST))),
      builder.CreatePointerCast(v0, ptrType),
      llvm::ConstantPointerNull::get(ptrType)));

  return func;
}

llvm::Constant *
LibraryBuilder::buildLibraryV0ImportTable(std::string libraryName) {
  auto &context = module->getContext();
  auto *importTableType = makeImportTableType(context);
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  llvm::Constant *symbolNames =
      llvm::Constant::getNullValue(llvm::PointerType::get(context, 0));
  if (!imports.empty()) {
    SmallVector<llvm::Constant *> symbolNameValues;
    for (auto &import : imports) {
      auto symbolName = import.symbol_name;
      if (import.weak)
        symbolName = "?" + symbolName;
      symbolNameValues.push_back(createStringConstant(symbolName, module));
    }
    symbolNames = createArrayConstant(libraryName + "_import_names",
                                      llvm::PointerType::get(context, 0),
                                      symbolNameValues, module);
  }
  return llvm::ConstantStruct::get(
      importTableType, {
                           // count=
                           llvm::ConstantInt::get(i32Type, imports.size()),
                           // symbols=
                           symbolNames,
                       });
}

llvm::Constant *
LibraryBuilder::buildLibraryV0ExportTable(std::string libraryName) {
  auto &context = module->getContext();
  auto *exportTableType = makeExportTableType(context);
  auto *dispatchAttrsType = makeDispatchAttrsType(context);
  auto *sourceLocationType = makeSourceLocationType(context);
  auto *stageLocationTableType = makeStageLocationTableType(context);
  auto *i8Type = llvm::IntegerType::getInt8Ty(context);
  auto *i16Type = llvm::IntegerType::getInt16Ty(context);
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  auto *i64Type = llvm::IntegerType::getInt64Ty(context);
  auto *ptrType = llvm::PointerType::get(context, 0);

  // iree_hal_executable_export_table_v0_t::ptrs
  SmallVector<llvm::Constant *> exportPtrValues;
  for (auto dispatch : exports)
    exportPtrValues.push_back(dispatch.func);
  llvm::Constant *exportPtrs = createArrayConstant(
      libraryName + "_funcs", ptrType, exportPtrValues, module);

  // iree_hal_executable_export_table_v0_t::attrs
  // Always populate the attrs table as it contains required dispatch metadata.
  SmallVector<llvm::Constant *> exportAttrValues;
  for (auto dispatch : exports) {
    exportAttrValues.push_back(llvm::ConstantStruct::get(
        dispatchAttrsType,
        {
            // flags=
            llvm::ConstantInt::get(i64Type,
                                   static_cast<uint64_t>(dispatch.attrs.flags)),
            // local_memory_pages=
            llvm::ConstantInt::get(
                i16Type, roundUpToAlignment(dispatch.attrs.localMemorySize,
                                            kWorkgroupLocalMemoryPageSize) /
                             kWorkgroupLocalMemoryPageSize),
            // constant_count=
            llvm::ConstantInt::get(i8Type, dispatch.attrs.constantCount),
            // binding_count=
            llvm::ConstantInt::get(i8Type, dispatch.attrs.bindingCount),
            // workgroup_size_x=
            llvm::ConstantInt::get(i32Type, dispatch.attrs.workgroupSize[0]),
            // workgroup_size_y=
            llvm::ConstantInt::get(i32Type, dispatch.attrs.workgroupSize[1]),
            // workgroup_size_z=
            llvm::ConstantInt::get(i16Type, dispatch.attrs.workgroupSize[2]),
            // parameter_count=
            llvm::ConstantInt::get(i16Type, dispatch.params.size()),
            // reserved_1[0]=
            llvm::ConstantInt::get(i64Type, 0),
            // reserved_1[1]=
            llvm::ConstantInt::get(i64Type, 0),
            // reserved_1[2]=
            llvm::ConstantInt::get(i64Type, 0),
            // reserved_1[3]=
            llvm::ConstantInt::get(i64Type, 0),
            // reserved_1[4]=
            llvm::ConstantInt::get(i64Type, 0),
        }));
  }
  llvm::Constant *exportAttrs = createArrayConstant(
      libraryName + "_attrs", dispatchAttrsType, exportAttrValues, module);

  // iree_hal_executable_export_table_v0_t::names
  llvm::Constant *exportNames = llvm::Constant::getNullValue(ptrType);
  if (mode == Mode::INCLUDE_REFLECTION_ATTRS) {
    SmallVector<llvm::Constant *> exportNameValues;
    for (auto dispatch : exports)
      exportNameValues.push_back(createStringConstant(dispatch.name, module));
    exportNames = createArrayConstant(libraryName + "_names", ptrType,
                                      exportNameValues, module);
  }

  // iree_hal_executable_export_table_v0_t::tags
  llvm::Constant *exportTags = llvm::Constant::getNullValue(ptrType);
  bool hasAnyTags = llvm::any_of(
      exports, [](auto &dispatch) { return !dispatch.tag.empty(); });
  if (mode == Mode::INCLUDE_REFLECTION_ATTRS && hasAnyTags) {
    SmallVector<llvm::Constant *> exportTagValues;
    for (auto dispatch : exports)
      exportTagValues.push_back(
          createStringConstantOrNull(dispatch.tag, module));
    exportTags = createArrayConstant(libraryName + "_tags", ptrType,
                                     exportTagValues, module);
  }

  // iree_hal_executable_export_table_v0_t::source_locations
  llvm::Constant *exportSourceLocations = llvm::Constant::getNullValue(ptrType);
  if (mode == Mode::INCLUDE_REFLECTION_ATTRS) {
    SmallVector<llvm::Constant *> exportSourceLocationValues;
    for (auto dispatch : exports) {
      exportSourceLocationValues.push_back(llvm::ConstantStruct::get(
          sourceLocationType,
          {
              // line=
              llvm::ConstantInt::get(i32Type, dispatch.sourceLocation.line),
              // path_length=
              llvm::ConstantInt::get(i32Type,
                                     dispatch.sourceLocation.path.size()),
              // path=
              createStringConstant(dispatch.sourceLocation.path, module),
          }));
    }
    exportSourceLocations = createArrayConstant(
        libraryName + "_source_locations", sourceLocationType,
        exportSourceLocationValues, module);
  }

  // iree_hal_executable_export_table_v0_t::params
  // Build unique string table for parameter names.
  llvm::DenseSet<StringRef> uniqueParamNames;
  for (const auto &dispatch : exports) {
    for (const auto &param : dispatch.params) {
      if (!param.name.empty()) {
        uniqueParamNames.insert(param.name);
      }
    }
  }

  // Create sorted list of parameter names for consistent indexing.
  SmallVector<StringRef> paramNameList(uniqueParamNames.begin(),
                                       uniqueParamNames.end());
  llvm::sort(paramNameList);

  // Build name to index mapping.
  llvm::DenseMap<StringRef, uint16_t> paramNameIndices;
  for (size_t i = 0; i < paramNameList.size(); ++i) {
    paramNameIndices[paramNameList[i]] = static_cast<uint16_t>(i);
  }

  // Build params array (2D - one array per export).
  auto *dispatchParamType = makeDispatchParameterType(context);
  llvm::Constant *exportParams = llvm::Constant::getNullValue(ptrType);
  bool hasAnyParams = llvm::any_of(
      exports, [](const auto &dispatch) { return !dispatch.params.empty(); });
  if (hasAnyParams) {
    SmallVector<llvm::Constant *> exportParamArrays;
    for (const auto &dispatch : exports) {
      if (dispatch.params.empty()) {
        exportParamArrays.push_back(llvm::Constant::getNullValue(ptrType));
      } else {
        SmallVector<llvm::Constant *> paramValues;
        for (const auto &param : dispatch.params) {
          uint16_t nameIndex = param.name.empty()
                                   ? static_cast<uint16_t>(-1)
                                   : paramNameIndices[param.name];
          paramValues.push_back(llvm::ConstantStruct::get(
              dispatchParamType,
              {
                  // type=
                  llvm::ConstantInt::get(i8Type,
                                         static_cast<uint8_t>(param.type)),
                  // size=
                  llvm::ConstantInt::get(i8Type, param.size),
                  // flags=
                  llvm::ConstantInt::get(i16Type, param.flags),
                  // name=
                  llvm::ConstantInt::get(i16Type, nameIndex),
                  // offset=
                  llvm::ConstantInt::get(i16Type, param.offset),
              }));
        }
        auto *paramArray =
            createArrayConstant(libraryName + "_" + dispatch.name + "_params",
                                dispatchParamType, paramValues, module);
        exportParamArrays.push_back(paramArray);
      }
    }
    exportParams = createArrayConstant(libraryName + "_params_table", ptrType,
                                       exportParamArrays, module);
  }

  // iree_hal_executable_export_table_v0_t::occupancy
  // Currently just placeholder nulls.
  llvm::Constant *exportOccupancy = llvm::Constant::getNullValue(ptrType);

  // iree_hal_executable_export_table_v0_t::parameter_names
  llvm::Constant *exportParamNames = llvm::Constant::getNullValue(ptrType);
  if (!paramNameList.empty()) {
    SmallVector<llvm::Constant *> paramNameValues;
    for (const auto &name : paramNameList) {
      paramNameValues.push_back(createStringConstant(name, module));
    }
    exportParamNames = createArrayConstant(libraryName + "_parameter_names",
                                           ptrType, paramNameValues, module);
  }

  // iree_hal_executable_export_table_v0_t::stage_locations
  llvm::Constant *exportStageLocations = llvm::Constant::getNullValue(ptrType);
  if (mode == Mode::INCLUDE_REFLECTION_ATTRS) {
    SmallVector<llvm::Constant *> exportStageTableValues;
    for (auto dispatch : exports) {
      SmallVector<llvm::Constant *> exportStageNameValues;
      SmallVector<llvm::Constant *> exportSourceLocationValues;
      for (auto &stageLocation : dispatch.stageLocations) {
        exportStageNameValues.push_back(
            createStringConstant(stageLocation.stage, module));
        exportSourceLocationValues.push_back(llvm::ConstantStruct::get(
            sourceLocationType,
            {
                // line=
                llvm::ConstantInt::get(i32Type, stageLocation.line),
                // path_length=
                llvm::ConstantInt::get(i32Type, stageLocation.path.size()),
                // path=
                createStringConstant(stageLocation.path, module),
            }));
      }
      llvm::Constant *stageNamesPtr = createArrayConstant(
          libraryName + "_" + dispatch.name + "_stage_names", ptrType,
          exportStageNameValues, module);
      llvm::Constant *sourceLocationsPtr = createArrayConstant(
          libraryName + "_" + dispatch.name + "_stage_source_locations",
          sourceLocationType, exportSourceLocationValues, module);
      exportStageTableValues.push_back(llvm::ConstantStruct::get(
          stageLocationTableType,
          {
              // count=
              llvm::ConstantInt::get(i32Type, exportStageNameValues.size()),
              // names=
              stageNamesPtr,
              // locations=
              sourceLocationsPtr,
          }));
    }
    if (!exportStageTableValues.empty()) {
      exportStageLocations = createArrayConstant(
          libraryName + "_stage_location_tables", stageLocationTableType,
          exportStageTableValues, module);
    }
  }

  return llvm::ConstantStruct::get(
      exportTableType, {
                           // count=
                           llvm::ConstantInt::get(i32Type, exports.size()),
                           // ptrs=
                           exportPtrs,
                           // attrs=
                           exportAttrs,
                           // params=
                           exportParams,
                           // occupancy=
                           exportOccupancy,
                           // names=
                           exportNames,
                           // tags=
                           exportTags,
                           // parameter_names=
                           exportParamNames,
                           // source_locations=
                           exportSourceLocations,
                           // stage_locations=
                           exportStageLocations,
                       });
}

llvm::Constant *
LibraryBuilder::buildLibraryV0ConstantTable(std::string libraryName) {
  auto &context = module->getContext();
  auto *constantTableType = makeConstantTableType(context);
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  return llvm::ConstantStruct::get(
      constantTableType, {
                             // count=
                             llvm::ConstantInt::get(i32Type, constantCount),
                         });
}

llvm::Constant *
LibraryBuilder::buildLibraryV0SourceTable(std::string libraryName) {
  auto &context = module->getContext();
  auto *sourceFileType = makeSourceFileType(context);
  auto *sourceTableType = makeSourceTableType(context);
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);
  llvm::Constant *sourceFilesValue =
      llvm::Constant::getNullValue(llvm::PointerType::get(context, 0));
  if (!sourceFiles.empty()) {
    SmallVector<llvm::Constant *> sourceFileValues;
    for (auto &sourceFile : sourceFiles) {
      sourceFileValues.push_back(llvm::ConstantStruct::get(
          sourceFileType,
          {
              // path_length=
              llvm::ConstantInt::get(i32Type, sourceFile.path.size()),
              // path=
              createStringConstant(sourceFile.path, module),
              // content_length=
              llvm::ConstantInt::get(i32Type, sourceFile.contents.size()),
              // content=
              createBufferConstant(sourceFile.path, sourceFile.contents,
                                   module),
          }));
    }
    sourceFilesValue = createArrayConstant(
        libraryName + "_sources", sourceFileType, sourceFileValues, module);
  }
  return llvm::ConstantStruct::get(
      sourceTableType, {
                           // count=
                           llvm::ConstantInt::get(i32Type, sourceFiles.size()),
                           // files=
                           sourceFilesValue,
                       });
}

llvm::Constant *LibraryBuilder::buildLibraryV0(std::string libraryName) {
  auto &context = module->getContext();
  auto *libraryHeaderType = makeLibraryHeaderType(context);
  auto *libraryType = makeLibraryType(libraryHeaderType);
  auto *i32Type = llvm::IntegerType::getInt32Ty(context);

  // ----- Header -----

  auto *libraryHeader = new llvm::GlobalVariable(
      *module, libraryHeaderType, /*isConstant=*/true,
      llvm::GlobalVariable::PrivateLinkage,
      llvm::ConstantStruct::get(
          libraryHeaderType,
          {
              // version=
              llvm::ConstantInt::get(i32Type,
                                     static_cast<int64_t>(Version::LATEST)),
              // name=
              createStringConstant(module->getName(), module),
              // features=
              llvm::ConstantInt::get(i32Type, static_cast<int64_t>(features)),
              // sanitizer=
              llvm::ConstantInt::get(i32Type,
                                     static_cast<int64_t>(sanitizerKind)),
          }),
      /*Name=*/libraryName + "_header");
  // TODO(benvanik): force alignment (8? natural pointer width?)

  // ----- Library -----

  auto *library = new llvm::GlobalVariable(
      *module, libraryType, /*isConstant=*/true,
      llvm::GlobalVariable::PrivateLinkage,
      llvm::ConstantStruct::get(libraryType,
                                {
                                    // header=
                                    libraryHeader,
                                    // imports=
                                    buildLibraryV0ImportTable(libraryName),
                                    // exports=
                                    buildLibraryV0ExportTable(libraryName),
                                    // constants=
                                    buildLibraryV0ConstantTable(libraryName),
                                    // sources=
                                    buildLibraryV0SourceTable(libraryName),
                                }),
      /*Name=*/libraryName);
  // TODO(benvanik): force alignment (8? natural pointer width?)

  return library;
}

} // namespace mlir::iree_compiler::IREE::HAL