Add AddVISADumpDeclarationsToEnd flag

Add a IGC flag that adds a comment with
.decl section at the end of VISAASM console dump.

Author: fda0

IGC/Compiler/CISACodeGen/CISABuilder.cpp

@@ -3819,6 +3819,9 @@ void CEncoder::InitVISABuilderOptions(TARGET_PLATFORM VISAPlatform, bool canAbor
   if (IGC_IS_FLAG_ENABLED(DumpASMToConsole)) {
     SaveOption(vISA_asmToConsole, true);
   }
+  if (IGC_IS_FLAG_ENABLED(AddVISADumpDeclarationsToEnd)) {
+    SaveOption(vISA_AddISAASMDeclarationsToEnd, true);
+  }
   if (IGC_IS_FLAG_ENABLED(AddExtraIntfInfo)) {
     SaveOption(vISA_AddExtraIntfInfo, true);
   }

IGC/common/igc_flags.h

@@ -56,6 +56,8 @@ DECLARE_IGC_REGKEY(bool, EnableIGAEncoder, false, "Enable VISA IGA encoder", fal
 DECLARE_IGC_REGKEY(bool, EnableVISADumpCommonISA, false, "Enable VISA Dump Common ISA", true)
 DECLARE_IGC_REGKEY(bool, DumpVISAASMToConsole, false, "Dump VISAASM to console and do early exit", true)
 DECLARE_IGC_REGKEY(bool, DumpASMToConsole, false, "Dump ASM to console and do early exit", true)
+DECLARE_IGC_REGKEY(bool, AddVISADumpDeclarationsToEnd, false,
+                   "Add a comment with .decl section to the end of VISA console dump. Used in tests.", true)
 DECLARE_IGC_REGKEY(bool, EnableVISABinary, false, "Enable VISA Binary", true)
 DECLARE_IGC_REGKEY(bool, EnableVISAOutput, false, "Enable VISA GenISA output", true)
 DECLARE_IGC_REGKEY(bool, EnableVISASlowpath, false, "Enable VISA Slowpath. Needed to dump .visaasm", true)

IGC/ocloc_tests/Options/AddVISADumpDeclarationsToEnd.cl

@@ -0,0 +1,22 @@
+/*========================== begin_copyright_notice ============================
+
+Copyright (C) 2025 Intel Corporation
+
+SPDX-License-Identifier: MIT
+
+============================= end_copyright_notice ===========================*/
+
+// REQUIRES: regkeys
+
+// RUN: ocloc compile -file %s -device dg2 -options "-igc_opts 'DumpVISAASMToConsole=1,AddVISADumpDeclarationsToEnd=1,ForceOCLSIMDWidth=8'" | FileCheck %s
+
+// CHECK-LABEL: .kernel "foo"
+// CHECK: add (M1_NM, 1) [[DST:.*]](0,0)<1> [[SRC_A:.*]](0,0)<0;1,0> [[SRC_B:.*]](0,0)<0;1,0>
+// CHECK: lsc_store{{.*}} (M1_NM, 1) {{.*}} [[DST]]
+// CHECK-DAG: // .decl [[SRC_A]] v_type=G type=d num_elts=1
+// CHECK-DAG: // .decl [[SRC_B]] v_type=G type=d num_elts=1
+// CHECK-DAG: // .decl [[DST]] v_type=G type=d num_elts=1
+
+__kernel void foo(int a, int b, __global int *res) {
+    *res = a + b;
+}

visa/BuildCISAIR.h

@@ -127,7 +127,8 @@ class CISA_IR_Builder : public VISABuilder {
 
   std::string isaDump(const VISAKernelImpl *kernel,
                       const VISAKernelImpl *mainKernel,
-                      bool printVersion = true) const;
+                      bool printVersion = true,
+                      bool addDeclCommentAtEnd = false) const;
 
   static void cat(std::stringstream &ss) {}
   template <typename T, typename... Ts>

visa/Common_ISA_framework.cpp

@@ -152,7 +152,8 @@ bool allowDump(const Options &options, const std::string &fullPath) {
 // repKernel has the actual shader body.
 std::string CISA_IR_Builder::isaDump(const VISAKernelImpl *kernel,
                                      const VISAKernelImpl *repKernel,
-                                     bool printVersion) const {
+                                     bool printVersion,
+                                     bool addDeclCommentAtEnd) const {
   std::stringstream sstr;
   VISAKernel_format_provider header(repKernel);
 
@@ -170,6 +171,12 @@ std::string CISA_IR_Builder::isaDump(const VISAKernelImpl *kernel,
     sstr << header.printInstruction(inst, kernel->getOptions()) << "\n";
   }
 
+  if (addDeclCommentAtEnd) {
+    // This path is enabled vISA_AddISAASMDeclarationsToEnd
+    // and is used in tests to verify the .decl section.
+    sstr << header.printDeclSection(true) << "\n";
+  }
+
 #ifdef DLL_MODE
   // Print the options used to compile this vISA object to assist debugging.
   if (kernel->getCISABuilder()->getBuilderOption() == VISA_BUILDER_BOTH) {
@@ -189,6 +196,7 @@ int CISA_IR_Builder::isaDump(const char *combinedIsaasmName) const {
   return VISA_SUCCESS;
 #else
   const bool isaasmToConsole = m_options.getOption(vISA_ISAASMToConsole);
+  const bool isaasmaddDeclAtEnd = m_options.getOption(vISA_AddISAASMDeclarationsToEnd);
   const bool genIsaasm = m_options.getOption(vISA_GenerateISAASM);
   const bool allowIsaasmDump = combinedIsaasmName && combinedIsaasmName[0] != '\0' &&
       CisaFramework::allowDump(m_options, combinedIsaasmName);
@@ -239,7 +247,7 @@ int CISA_IR_Builder::isaDump(const char *combinedIsaasmName) const {
     }
     if (genIsaasm) {
       if (isaasmToConsole) {
-        std::cout << isaDump(kTemp, repKernel);
+        std::cout << isaDump(kTemp, repKernel, true, isaasmaddDeclAtEnd);
       } else {
         std::ofstream out(asmFileName);
         out << isaDump(kTemp, repKernel);

visa/IsaDisassembly.cpp

@@ -3435,34 +3435,7 @@ std::string VISAKernel_format_provider::printKernelHeader(bool printVersion) {
   }
   sstr << "\n";
 
-  // emit var decls
-  //.decl  V<#> name=<name> type=<type> num_elts=<num_elements> [align=<align>]
-  //[alias=(<alias_index>,<alias_offset>)]
-  for (unsigned i = 0; i < getVarCount(); i++) {
-    sstr << "\n" << printVariableDecl(this, i, options);
-  }
-  // address decls
-  for (unsigned i = 0; i < getAddrCount(); i++) {
-    sstr << "\n" << printAddressDecl(this, i);
-  }
-  // pred decls
-  for (unsigned i = 0; i < getPredCount(); i++) {
-    // P0 is reserved; starting from P1 if there is predicate decl
-    sstr << "\n" << printPredicateDecl(this, i);
-  }
-  // sampler
-  for (unsigned i = 0; i < getSamplerCount(); i++) {
-    sstr << "\n" << printSamplerDecl(this, i);
-  }
-  // surface
-  unsigned numPreDefinedSurfs = Get_CISA_PreDefined_Surf_Count();
-  for (unsigned i = 0; i < getSurfaceCount(); i++) {
-    sstr << "\n" << printSurfaceDecl(this, i, numPreDefinedSurfs);
-  }
-  // inputs to kernel
-  for (unsigned i = 0; i < getInputCount(); i++) {
-    sstr << "\n" << printFuncInput(this, i, options);
-  }
+  sstr << printDeclSection(false);
 
   bool isTargetSet = false;
   for (unsigned i = 0; i < getAttrCount(); i++) {
@@ -3500,6 +3473,46 @@ std::string VISAKernel_format_provider::printKernelHeader(bool printVersion) {
   return sstr.str();
 }
 
+std::string VISAKernel_format_provider::printDeclSection(bool printAsComment) {
+  // printAsComment adds "// " at the beginning of each line
+  // This is used for adding duplicated decl section at the end
+  // of the dump when vISA_AddISAASMDeclarationsToEnd is enabled.
+
+  std::stringstream sstr;
+  const Options *options = m_kernel->getOptions();
+  const char *newLine = printAsComment ? "\n// " : "\n";
+
+  // emit var decls
+  //.decl  V<#> name=<name> type=<type> num_elts=<num_elements> [align=<align>]
+  //[alias=(<alias_index>,<alias_offset>)]
+  for (unsigned i = 0; i < getVarCount(); i++) {
+    sstr << newLine << printVariableDecl(this, i, options);
+  }
+  // address decls
+  for (unsigned i = 0; i < getAddrCount(); i++) {
+    sstr << newLine << printAddressDecl(this, i);
+  }
+  // pred decls
+  for (unsigned i = 0; i < getPredCount(); i++) {
+    // P0 is reserved; starting from P1 if there is predicate decl
+    sstr << newLine << printPredicateDecl(this, i);
+  }
+  // sampler
+  for (unsigned i = 0; i < getSamplerCount(); i++) {
+    sstr << newLine << printSamplerDecl(this, i);
+  }
+  // surface
+  unsigned numPreDefinedSurfs = Get_CISA_PreDefined_Surf_Count();
+  for (unsigned i = 0; i < getSurfaceCount(); i++) {
+    sstr << newLine << printSurfaceDecl(this, i, numPreDefinedSurfs);
+  }
+  // inputs to kernel
+  for (unsigned i = 0; i < getInputCount(); i++) {
+    sstr << newLine << printFuncInput(this, i, options);
+  }
+  return sstr.str();
+}
+
 std::string printFunctionDecl(const print_format_provider_t *header,
                               bool isKernel) {
   std::stringstream sstr;

visa/VISAKernel.h

@@ -1503,6 +1503,7 @@ class VISAKernel_format_provider : public print_format_provider_t {
   uint32_t getInputCount() const { return m_kernel->m_input_count; }
 
   std::string printKernelHeader(bool printVersion);
+  std::string printDeclSection(bool printAsComment);
 
   std::string printInstruction(const CISA_INST *instruction,
                                const Options *opt) const;

visa/include/VISAOptionsDefs.h

@@ -622,6 +622,8 @@ DEF_VISA_OPTION(vISA_GenerateCombinedISAASM, ET_BOOL, "-dumpcombinedcisa",
 DEF_VISA_OPTION(vISA_ISAASMToConsole, ET_BOOL, "-isaasmToConsole",
                 "The option can be used with -dumpcommonisa to make finalizer"
                 "emit isaasm to stdout instead of file and do early exit", false)
+DEF_VISA_OPTION(vISA_AddISAASMDeclarationsToEnd, ET_BOOL, "-isaasmAddDeclarationsAtEnd",
+                "Add a comment with .decl section to the end of isaasm console dump. Used in tests.", false)
 DEF_VISA_OPTION(vISA_DumpIsaVarNames, ET_BOOL, "-dumpisavarnames", UNUSED, true)
 DEF_VISA_OPTION(vISA_UniqueLabels, ET_BOOL, "-uniqueLabel", UNUSED, false)
 DEF_VISA_OPTION(vISA_ShaderDumpRegexFilter, ET_CSTR, "-shaderDumpRegexFilter",