Remove preserve_fp32_precision until it is replaced by unpack_to_dest…

…_mode (issue #879)
tenstorrent · Oct 9, 2024 · 977591c · 977591c
1 parent 94d52b5
commit 977591c
Show file tree

Hide file tree

Showing 5 changed files with 4 additions and 9 deletions.
diff --git a/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td b/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td
@@ -32,13 +32,12 @@ def TTKernel_TensixConfigAttr: TTKernel_Attr<"TensixConfig", "tensix_config", [T
   }];
   let parameters = (ins "MathFidelity":$math_fidelity,
                         "bool":$fp32_dest_acc_en,
-                        "bool":$preserve_fp32_precision,
                         "bool":$math_approx_mode);
-  let assemblyFormat =  "`<` $math_fidelity`,` $fp32_dest_acc_en`,` $preserve_fp32_precision`,` $math_approx_mode `>`";
+  let assemblyFormat =  "`<` $math_fidelity`,` $fp32_dest_acc_en`,` $math_approx_mode `>`";
 
   let extraClassDeclaration = [{
     static TensixConfigAttr get(::mlir::MLIRContext *context){
-      return TensixConfigAttr::get(context, MathFidelity::HiFi4, false, false, false);
+      return TensixConfigAttr::get(context, MathFidelity::HiFi4, false, false);
     };
 
     ThreadType getThreadType() const { return ThreadType::Tensix; }

diff --git a/include/ttmlir/Target/TTMetal/program.fbs b/include/ttmlir/Target/TTMetal/program.fbs
@@ -19,7 +19,6 @@ table NocConfig {
 table TensixConfig {
   math_fidelity: MathFidelity;
   fp32_dest_acc_en: bool;
-  preserve_fp32_precision: bool;
   math_approx_mode: bool;
 }
 

diff --git a/lib/Conversion/TTIRToTTMetal/TTIRToTTMetal.cpp b/lib/Conversion/TTIRToTTMetal/TTIRToTTMetal.cpp
@@ -299,7 +299,7 @@ class TTIRToTTMetalLayoutRewriter : public OpRewritePattern<ttir::ToLayoutOp> {
     assert(inputLayout.getGrid() == outputLayout.getGrid());
 
     auto tensixAttr = rewriter.getAttr<ttkernel::TensixConfigAttr>(
-        ttkernel::MathFidelity::HiFi4, false, false, false);
+        ttkernel::MathFidelity::HiFi4, false, false);
     SmallVector<Attribute> kernelConfigs = {tensixAttr};
     SmallVector<Attribute> coreRanges = {
         rewriter.getAttr<ttmetal::CoreRangeAttr>(inputLayout.getGrid()),
@@ -1007,7 +1007,7 @@ class TTIRToTTMetalDispatchRewriter : public OpRewritePattern<ttir::GenericOp> {
     }
 
     auto tensixAttr = rewriter.getAttr<ttkernel::TensixConfigAttr>(
-        ttkernel::MathFidelity::HiFi4, false, false, false);
+        ttkernel::MathFidelity::HiFi4, false, false);
     SmallVector<Attribute> kernelConfigs = {tensixAttr};
     SmallVector<Attribute> coreRanges = {
         rewriter.getAttr<ttmetal::CoreRangeAttr>(op.getGrid()),

diff --git a/lib/Target/TTMetal/TTMetalToFlatbuffer.cpp b/lib/Target/TTMetal/TTMetalToFlatbuffer.cpp
@@ -116,7 +116,6 @@ toFlatbuffer(::flatbuffers::FlatBufferBuilder &fbb,
     auto config = ::tt::target::metal::CreateTensixConfig(
         fbb, toFlatbuffer(tensixConfigAttr.getMathFidelity()),
         tensixConfigAttr.getFp32DestAccEn(),
-        tensixConfigAttr.getPreserveFp32Precision(),
         tensixConfigAttr.getMathApproxMode());
     return std::make_pair(configType, config.Union());
   }

diff --git a/runtime/lib/ttmetal/command_queue.cpp b/runtime/lib/ttmetal/command_queue.cpp
@@ -297,8 +297,6 @@ createKernelConfig(::tt::target::metal::KernelSource const *kernelSource) {
 
     computeConfig.fp32_dest_acc_en =
         kernelSource->config_as_TensixConfig()->fp32_dest_acc_en();
-    computeConfig.preserve_fp32_precision =
-        kernelSource->config_as_TensixConfig()->preserve_fp32_precision();
     computeConfig.math_approx_mode =
         kernelSource->config_as_TensixConfig()->math_approx_mode();
     return computeConfig;