Cleanup the way config values are retrieved. (#21610)

The main motivation here is to remove `getConfigStringAttr`,
`getConfigIntegerAttr` and `getConfigBoolAttr` methods from `Utils.h` to
avoid the leaking of names used for attributes stored in the
`DictionaryAttr:configuration` field of
`IREE::HAL::ExecutableTargetAttr`. Instead

- Have specific methods to retrieve known properties from the
`configuration` and also have these methods return the value directly
instead of the attribute.

- Most of these methods know just use `DictionaryAttr`. They could use
`IREE::HAL::ExecutableAttr` directly. This was initially tried but some
cases (related to MaterializeEncoding) use the `DictionaryAttr` without
using the `IREE::HAL::ExecutableAttr`. So it was more uniform to just
use the `DictionaryAttr` everywhere (only exception being
`getGPUTargetAttr` due to the complication of target being specified
from command line instead of `IREE::HAL::ExecutableTargetAttr`)

- The names used to access the various target information encoding in
the `configuration` field of the `hal.executable.target` are now made
private and these fields can be queried through new utility methods
added.

- Rename `iree.gpu.target` -> `iree_codegen.target_info`.

Most of the checking of null values needed were mostly to make lit-tests
happy. The end-to-end compilation flow always ensure the attributes are
never empty.

---------

Signed-off-by: MaheshRavishankar <[email protected]>

Author: MaheshRavishankar

compiler/plugins/target/CUDA/CUDATarget.cpp

@@ -424,7 +424,7 @@ class CUDATargetBackend final : public TargetBackend {
 
     if (auto target = GPU::getCUDATargetDetails(
             options.clTarget, options.clTargetFeatures, context)) {
-      configItems.emplace_back(kGPUTargetAttrName, target);
+      addConfigGPUTarget(context, target, configItems);
     }
 
     return b.getAttr<IREE::HAL::ExecutableTargetAttr>(
@@ -467,7 +467,8 @@ class CUDATargetBackend final : public TargetBackend {
     auto targetAttr = variantOp.getTargetAttr();
     StringRef targetArch = options.clTarget;
     StringRef targetFeatures = options.clTargetFeatures;
-    if (auto attr = getGPUTargetAttr(targetAttr)) {
+    if (auto attr =
+            getGPUTargetAttr(executableBuilder.getContext(), targetAttr)) {
       targetArch = attr.getArch();
       targetFeatures = attr.getFeatures();
     }

compiler/plugins/target/LLVMCPU/BUILD.bazel

@@ -128,6 +128,7 @@ iree_compiler_cc_library(
     ],
     deps = [
         ":ResolveCPUAndCPUFeatures",
+        "//compiler/src/iree/compiler/Codegen/LLVMCPU",
         "//compiler/src/iree/compiler/Utils",
         "@llvm-project//llvm:Analysis",
         "@llvm-project//llvm:Core",

compiler/plugins/target/LLVMCPU/CMakeLists.txt

@@ -123,6 +123,7 @@ iree_cc_library(
     LLVMTarget
     LLVMTargetParser
     MLIRIR
+    iree::compiler::Codegen::LLVMCPU
     iree::compiler::Utils
   PUBLIC
 )

compiler/plugins/target/LLVMCPU/LLVMCPUTarget.cpp

@@ -245,7 +245,8 @@ class LLVMCPUTargetBackend final : public TargetBackend {
 
   void buildTranslationPassPipeline(IREE::HAL::ExecutableTargetAttr targetAttr,
                                     OpPassManager &passManager) override {
-    bool enableAArch64SME = isAArch64(targetAttr) && hasSMEFeature(targetAttr);
+    bool enableAArch64SME = isAArch64(targetAttr.getConfiguration()) &&
+                            hasSMEFeature(targetAttr.getConfiguration());
     buildLLVMCPUCodegenPassPipeline(passManager, enableAArch64SME);
   }
 
@@ -556,7 +557,7 @@ class LLVMCPUTargetBackend final : public TargetBackend {
 
     if (target.linkUkernelBitcode) {
       // Link in ukernel bitcode.
-      if (hasUkernel(variantOp.getTarget())) {
+      if (hasUkernel(variantOp.getTarget().getConfiguration())) {
         llvm::Expected<std::unique_ptr<llvm::Module>> bitcode =
             loadUKernelBitcode(targetMachine.get(), context);
         if (!bitcode) {

compiler/plugins/target/LLVMCPU/LLVMTargetOptions.cpp

@@ -7,6 +7,7 @@
 #include "compiler/plugins/target/LLVMCPU/LLVMTargetOptions.h"
 
 #include "compiler/plugins/target/LLVMCPU/ResolveCPUAndCPUFeatures.h"
+#include "iree/compiler/Codegen/LLVMCPU/Utils.h"
 #include "llvm/ADT/APFloat.h"
 #include "llvm/ADT/StringRef.h"
 #include "llvm/Analysis/TargetTransformInfo.h"
@@ -126,20 +127,17 @@ void LLVMTarget::storeToConfigAttrs(MLIRContext *context,
   auto addBool = [&](StringRef name, bool value) {
     config.emplace_back(b.getStringAttr(name), b.getBoolAttr(value));
   };
-  auto addInt64 = [&](StringRef name, int64_t value) {
-    config.emplace_back(b.getStringAttr(name), b.getI64IntegerAttr(value));
-  };
 
-  addString("target_triple", triple);
+  addConfigTargetTriple(context, triple, config);
   addString("cpu", cpu);
-  addString("cpu_features", cpuFeatures);
+  addConfigCpuFeatures(context, cpuFeatures, config);
   if (!dataLayout.empty()) {
-    addString("data_layout", dataLayout);
+    addConfigDataLayout(context, dataLayout, config);
   }
   if (vectorWidthInBytes != DEFAULT_VECTOR_WIDTH_IN_BYTES) {
-    addInt64("native_vector_size", vectorWidthInBytes);
+    addConfigNativeVectorSize(context, vectorWidthInBytes, config);
   }
-  addInt64("max_stack_allocation_size", maxStackAllocSizeInBytes);
+  addConfigMaxStackAllocationSize(context, maxStackAllocSizeInBytes, config);
   if (linkEmbedded != DEFAULT_LINK_EMBEDDED) {
     addBool("link_embedded", linkEmbedded);
   }
@@ -234,24 +232,13 @@ LLVMTarget::loadFromConfigAttr(Location loc, DictionaryAttr config,
     }
     return fallback;
   };
-  auto getInt64 = [&](StringRef name, int64_t fallback) -> int64_t {
-    Attribute attr = config.get(name);
-    if (auto iattr = llvm::dyn_cast_if_present<IntegerAttr>(attr)) {
-      return iattr.getValue().getSExtValue();
-    } else if (attr) {
-      hasFailures = true;
-      emitError(loc) << "executable config '" << name
-                     << "' requires i64 but got " << attr;
-    }
-    return fallback;
-  };
 
   LLVMTarget target;
 
   // Constructor arguments.
-  auto triple = getOptionalString("target_triple");
+  auto triple = getConfigTargetTriple(config);
   auto cpu = getOptionalString("cpu");
-  auto cpuFeatures = getOptionalString("cpu_features");
+  auto cpuFeatures = getConfigCpuFeatures(config);
   bool linkEmbedded = getBool("link_embedded", DEFAULT_LINK_EMBEDDED);
   if (triple || cpu || cpuFeatures) {
     if (!triple) {
@@ -279,9 +266,9 @@ LLVMTarget::loadFromConfigAttr(Location loc, DictionaryAttr config,
     target.copy(defaultTarget);
   }
 
-  target.dataLayout = getString("data_layout", DEFAULT_DATA_LAYOUT, false);
+  target.dataLayout = getConfigDataLayout(config).value_or(DEFAULT_DATA_LAYOUT);
   target.vectorWidthInBytes =
-      getInt64("native_vector_size", DEFAULT_VECTOR_WIDTH_IN_BYTES);
+      getConfigNativeVectorSize(config).value_or(DEFAULT_VECTOR_WIDTH_IN_BYTES);
 
   target.debugSymbols = getBool("debug_symbols", DEFAULT_DEBUG_SYMBOLS);
   target.linkStatic = getBool("link_static", DEFAULT_LINK_STATIC);

compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp

@@ -98,7 +98,7 @@ class MetalSPIRVTargetBackend : public TargetBackend {
     Builder b(context);
     SmallVector<NamedAttribute, 1> configItems;
     if (auto target = GPU::getMetalTargetDetails(context)) {
-      configItems.emplace_back(kGPUTargetAttrName, target);
+      addConfigGPUTarget(context, target, configItems);
     }
 
     return b.getAttr<IREE::HAL::ExecutableTargetAttr>(

compiler/plugins/target/MetalSPIRV/test/smoketest.mlir

@@ -4,7 +4,7 @@ module attributes {
   hal.device.targets = [
     #hal.device.target<"metal", [
       #hal.executable.target<"metal-spirv", "metal-msl-fb", {
-        iree.gpu.target = #iree_gpu.target<arch = "", features = "spirv:v1.3,cap:Shader", wgp = <
+        iree_codegen.target_info = #iree_gpu.target<arch = "", features = "spirv:v1.3,cap:Shader", wgp = <
           compute = fp32|int32, storage = b32, subgroup = none, dot = none, mma = [], scaled_mma = [], subgroup_size_choices = [32],
           max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384,
           max_workgroup_counts = [65535, 65535, 65535]>>

compiler/plugins/target/ROCM/ROCMTarget.cpp

@@ -285,7 +285,7 @@ class ROCMTargetBackend final : public TargetBackend {
 
     if (auto target = GPU::getHIPTargetDetails(
             options.target, options.targetFeatures, context)) {
-      addConfig(kGPUTargetAttrName, target);
+      addConfigGPUTarget(context, target, configItems);
       if (options.encodingLayoutResolver !=
           GPU::kNoEncodingLayoutResolverName) {
         if (Attribute encoding = GPU::getHIPTargetEncodingLayoutAttr(
@@ -326,7 +326,7 @@ class ROCMTargetBackend final : public TargetBackend {
 
     addConfig("ukernels", b.getStringAttr(options.enableROCMUkernels));
     if (options.wavesPerEu > 0) {
-      addConfig("waves_per_eu", b.getI64IntegerAttr(options.wavesPerEu));
+      addConfigWavesPerEu(b.getContext(), options.wavesPerEu, configItems);
     }
 
     return b.getAttr<IREE::HAL::ExecutableTargetAttr>(
@@ -350,7 +350,7 @@ class ROCMTargetBackend final : public TargetBackend {
   buildConfigurationPassPipeline(IREE::HAL::ExecutableTargetAttr targetAttr,
                                  OpPassManager &passManager) override {
     if (options.specializeDispatches) {
-      if (auto attr = getGPUTargetAttr(targetAttr)) {
+      if (auto attr = getGPUTargetAttr(targetAttr.getContext(), targetAttr)) {
         ROCM::ApplyBuiltinPDLPatternsPassOptions options;
         options.enableSpecialization = true;
         if (IREE::GPU::TargetChipAttr chip = attr.getChip()) {
@@ -443,7 +443,7 @@ class ROCMTargetBackend final : public TargetBackend {
     auto targetAttr = variantOp.getTargetAttr();
     StringRef targetArch = options.target;
     StringRef targetFeatures = options.targetFeatures;
-    if (auto attr = getGPUTargetAttr(targetAttr)) {
+    if (auto attr = getGPUTargetAttr(variantOp.getContext(), targetAttr)) {
       targetArch = attr.getArch();
       targetFeatures = attr.getFeatures();
     }

compiler/plugins/target/ROCM/test/external_function_validation.mlir

@@ -5,7 +5,7 @@
 // remain in the final bitcode (post device bitcode linking).
 
 #executable_target_rocm_hsaco_fb = #hal.executable.target<"rocm", "rocm-hsaco-fb",
-  {iree.gpu.target = #iree_gpu.target<arch = "gfx942", features = "",
+  {iree_codegen.target_info = #iree_gpu.target<arch = "gfx942", features = "",
                                       wgp = <compute = fp16, storage =  b16,
                                       subgroup =  none, dot =  none, mma = [], scaled_mma = [],
                                       subgroup_size_choices = [64],

compiler/plugins/target/ROCM/test/smoketest_hsaco.mlir

@@ -17,7 +17,7 @@
 ]>
 #executable_target = #hal.executable.target<"rocm", "rocm-hsaco-fb", {
   abi = "hip",
-  iree.gpu.target = #iree_gpu.target<arch = "gfx942", features = "", wgp = <compute =  fp64|fp32|fp16|int64|int32|int16|int8, storage =  b64|b32|b16|b8, subgroup =  shuffle|arithmetic, dot =  dp4xi8toi32, mma = [<MFMA_F32_16x16x16_BF16>, <MFMA_F32_32x32x8_BF16>, <MFMA_F32_16x16x32_F8E5M2FNUZ>, <MFMA_F32_16x16x32_F8E5M2FNUZ_F8E4M3FNUZ>, <MFMA_F32_16x16x32_F8E4M3FNUZ>, <MFMA_F32_16x16x32_F8E4M3FNUZ_F8E5M2FNUZ>, <MFMA_F32_32x32x16_F8E5M2FNUZ>, <MFMA_F32_32x32x16_F8E5M2FNUZ_F8E4M3FNUZ>, <MFMA_F32_32x32x16_F8E4M3FNUZ>, <MFMA_F32_32x32x16_F8E4M3FNUZ_F8E5M2FNUZ>, <MFMA_I32_16x16x32_I8>, <MFMA_I32_32x32x16_I8>, <MFMA_F64_16x16x4_F64>, <MFMA_F32_16x16x4_F32>, <MFMA_F32_16x16x16_F16>, <MFMA_F32_32x32x8_F16>], scaled_mma = [], subgroup_size_choices = [64], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 65536, max_workgroup_counts = [2147483647, 2147483647, 2147483647], max_load_instruction_bits = 128, simds_per_wgp = 4, vgpr_space_bits = 16384>>
+  iree_codegen.target_info = #iree_gpu.target<arch = "gfx942", features = "", wgp = <compute =  fp64|fp32|fp16|int64|int32|int16|int8, storage =  b64|b32|b16|b8, subgroup =  shuffle|arithmetic, dot =  dp4xi8toi32, mma = [<MFMA_F32_16x16x16_BF16>, <MFMA_F32_32x32x8_BF16>, <MFMA_F32_16x16x32_F8E5M2FNUZ>, <MFMA_F32_16x16x32_F8E5M2FNUZ_F8E4M3FNUZ>, <MFMA_F32_16x16x32_F8E4M3FNUZ>, <MFMA_F32_16x16x32_F8E4M3FNUZ_F8E5M2FNUZ>, <MFMA_F32_32x32x16_F8E5M2FNUZ>, <MFMA_F32_32x32x16_F8E5M2FNUZ_F8E4M3FNUZ>, <MFMA_F32_32x32x16_F8E4M3FNUZ>, <MFMA_F32_32x32x16_F8E4M3FNUZ_F8E5M2FNUZ>, <MFMA_I32_16x16x32_I8>, <MFMA_I32_32x32x16_I8>, <MFMA_F64_16x16x4_F64>, <MFMA_F32_16x16x4_F32>, <MFMA_F32_16x16x16_F16>, <MFMA_F32_32x32x8_F16>], scaled_mma = [], subgroup_size_choices = [64], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 65536, max_workgroup_counts = [2147483647, 2147483647, 2147483647], max_load_instruction_bits = 128, simds_per_wgp = 4, vgpr_space_bits = 16384>>
 }>
 hal.executable public @executable {
   hal.executable.variant public @rocm_hsaco_fb target(#executable_target) {