Fail to insert a call to function in opt pass

Beginners
#1

I am writing a custom pass to insert a function recording some arguments of specific functions, but I fail to insert a call of self-record function.

  1. Following is my custom pass
#include<string>
#include<vector>
#include<iostream>
#include<fstream>
#include "llvm/Pass.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/Module.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/Support/FileSystem.h"
#include "llvm/IR/LegacyPassManager.h"
#include "llvm/Transforms/IPO/PassManagerBuilder.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Mangler.h"
#include "llvm/IR/Argument.h"

using namespace llvm;

//This Pass is to recognize every GemmEx for fp32 and output the M,N,K into one file as well as
//Replace them with wrapper function with an additional ID
//then output IDs of gemm wrapped functions which are called into ont file

/*
Module *ParseBitcodeFile(MemoryBuffer *Buffer, LLVMContext& Context,
                         std::string *ErrMsg = 0);
*/

namespace {
    struct MDFID : public ModulePass {
        static char ID;

        MDFID();
        ~MDFID() {};

        bool runOnModule (Module & M)override;
        bool set_dimensionfile_path(std::string module_id);
        bool set_funcIDfile_path(std::string module_id);

        //self data
        int Func_Counter;
        std::string dimension_file_path;
        std::string FuncID_run_file_path; 
    };

    MDFID::MDFID() : ModulePass(ID)
    {
        Func_Counter = 0;
    }

    bool MDFID::set_dimensionfile_path(std::string id)          //Module ID would be the full abs path of input ll file 
    {
        dimension_file_path = id+"_dimension.txt";
        errs()<<"The target dimension file path is "<<dimension_file_path<<"\n";
        std::fstream df;
        df.open(dimension_file_path, std::ios::out);
        if(!df)
        {
            errs()<<"Dimension file not exist, create one\n";
            df<<"\n";
            df.close();
            return true;
        }
        else
        {
            //TO.DO.: Need a check to ensure it can reflush the file with original context
            errs()<<"Dimension file exists, but need to reflush it\n";
            df<<"\n";
            df.close();
            return true;
        }
    }

    bool MDFID::set_funcIDfile_path(std::string id)
    {
        FuncID_run_file_path = id+"_calledFuncID.txt";      //TO.DO.: the id appears to be <stdin> in result. fix it 
        errs()<<"The target CalledFuncID file path is "<<FuncID_run_file_path<<"\n";
        std::fstream df;
        df.open(FuncID_run_file_path, std::ios::out);
        if(!df)
        {
            errs()<<"CalledFuncID file not exist, create one\n";
            df<<"\n";
            df.close();
            return true;
        }
        else
        {
            //TO.DO.: Need a check to ensure it can reflush the file with original context
            errs()<<"CalledFuncID file exists, but need to reflush it\n";
            df<<"\n";
            df.close();
            return true;
        }
    }

    bool MDFID::runOnModule(Module & M)
    {
        //Set up the output file
        std::string ModuleName = M.getModuleIdentifier();
        std::cout<<"Module name is "<<ModuleName<<std::endl;                //TO.DO.: The modulename is <stdin>. why?       ANS.: Dont use < in command line in opt. just pass the filename and the Module name will be the abs path of input file module
        if(!set_dimensionfile_path(ModuleName))
        {
            errs()<<"Wrong when setting dimension file path \n";
            exit(1);
        }
        if(!set_funcIDfile_path(ModuleName))
        {
            errs()<<"Wrong when setting funcID file path \n";
            exit(1);
        }

        //Loop over all instructions, if one is call GemmEx, add one IR right behind it to output its dimension
        //QUES.: Can I use getFunction to locate all GemmEx functions?
        //NOTE: We cannot use the result of M.getFunctionList() outside the current line
        //const Module::FunctionListType & func_list = M.getFunctionList();

        //The tool to generate IR
        //IRBuilder<> builder(M.getContext());

        //get the inserted function ptr
        //TO.DO.: Maybe we can use Module::getOrInsertFunction() to get function ptr;
        //Function * output_dimension_func_ptr = M.getFunction("output_dimension_file");            //This cannot get the function
        /*
        std::string mangled_output_dimension_func_name;
        raw_string_ostream mangled_output_dimension_func_stream(mangled_output_dimension_func_name);
        Mangler::getNameWithPrefix(mangled_output_dimension_func_stream,"output_dimension_file",M.getDataLayout());
        std::string new_name = mangled_output_dimension_func_stream.str();
        */
        Function * output_dimension_func_ptr = M.getFunction(StringRef("_Z21output_dimension_fileiii"));
        if(output_dimension_func_ptr == nullptr)
        {
            errs()<<"Cannot get the odf func\n";
            exit(1);
        }
        Function * output_CalledFuncID_func_ptr = M.getFunction("output_CalledFuncID_file");
        //Value * arg_CalledFuncID_file_path = builder.CreateGlobalStringPtr(StringRef(FuncID_run_file_path));
        FunctionType * output_func_type = FunctionType::get(FunctionType::getVoidTy(M.getContext()),false);     //getVoidTy is a method in Class Type
        

        //SymbolTableList<Function> == FunctionListType
        for(Module::FunctionListType::iterator func = M.getFunctionList().begin(),
        end_Func = M.getFunctionList().end(); func != end_Func; func++)
        {
            //Function::iterator == BasicBlockListType::iterator
            errs()<<"Now we are facing declare of function "<<func->getName()<<"\n";
            for(Function::iterator bb = func->begin(); bb != func->end(); bb++)
            {
                for(BasicBlock::iterator inst = bb->begin(); inst != bb->end(); inst++)
                {
                    if(isa<CallInst>(inst))
                    {
                        //QUES.: What about InvokeInst which means indirect call?
                        Function * called_func = dyn_cast<CallInst>(inst)->getCalledFunction();
                        if(called_func && called_func->getName() == "_Z6GemmExiiiiii")
                        {
                            errs()<<"we now get the gemm func\n";
                            /*
                            int counter = 0;
                            Argument * arg_M, * arg_N, * arg_K;
                            for(auto arg = called_func->arg_begin(), arg_end = called_func->arg_end(); arg != arg_end; arg++)           //QUES.: Is this arg the one in def or in calling?
                            {
                                if(counter == 3) arg_M = arg;
                                else if(counter == 4) arg_N = arg;
                                else if(counter == 5) arg_K = arg;
                                counter++;
                            }
                            */
                            CallInst * call_inst = dyn_cast<CallInst>(inst);            //Only in this way can we get the true arguments used in callinst
                            Value * arg_M = call_inst->getArgOperand(3);
                            Value * arg_N = call_inst->getArgOperand(4);
                            Value * arg_K = call_inst->getArgOperand(5);
                            errs()<<"Get all args we need\n";

                            //Dont directly use errs() to print, use a print_func in self-made library(which can use errs() inside)
                            //builder.SetInsertPoint(inst);
                            IRBuilder<> builder((Instruction*)&*inst);
                            char file_path[100];
                            strcpy(file_path,dimension_file_path.c_str());
                            //Value * arg_dimension_file_path = builder.CreateGlobalStringPtr(file_path);
                            Value * argv_M = dynamic_cast<Value*>(arg_M);
                            Value * argv_N = dynamic_cast<Value*>(arg_N);
                            Value * argv_K = dynamic_cast<Value*>(arg_K);
                            std::vector<Value*> args = {argv_M,argv_N,argv_K};
                            CallInst * new_inst = builder.CreateCall(output_func_type,output_dimension_func_ptr,makeArrayRef(args));
                            
                            
                        }
                        else
                        {
                            errs()<<"This is not a call to function or a GemmEx function\n";
                        }
                    }
                }
            }
        }
        return true;
    }

}


char MDFID::ID = 0;

static RegisterPass<MDFID> X("mdfid","A Pass to Mark GEMM Dimension and IDs of Called GEMM Functions",
                                true, false);

static RegisterStandardPasses Y(
    PassManagerBuilder::EP_EarlyAsPossible,
    [] (const PassManagerBuilder & Builder, 
    legacy::PassManagerBase & PM) {PM.add(new MDFID());});
  1. Following is error report
Called function is not the same type as the call!
  call void @_Z21output_dimension_fileiii(i32 4, i32 5, i32 6)
in function main
LLVM ERROR: Broken function found, compilation aborted!
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace.
Stack dump:
0.      Program arguments: opt -load /mnt/data/home/mzw/hipcc_llvm_4.3.0/lib/libMDFID.so -mdfid -enable-new-pm=0 /mnt/data/home/mzw/workspace/test_space/llvm_test/mark_test/mark_test-c5c146.ll -o /mnt/data/home/mzw/workspace/test_space/llvm_test/mark_test/mark_test-new.bc
1.      Running pass 'Function Pass Manager' on module '/mnt/data/home/mzw/workspace/test_space/llvm_test/mark_test/mark_test-c5c146.ll'.
2.      Running pass 'Module Verifier' on function '@main'
 #0 0x000055a70e3ae64c llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/mnt/data/home/mzw/hipcc_llvm_4.3.0/bin/opt+0x1c5364c)
 #1 0x000055a70e3ac3f4 llvm::sys::RunSignalHandlers() (/mnt/data/home/mzw/hipcc_llvm_4.3.0/bin/opt+0x1c513f4)
 #2 0x000055a70e3ac563 SignalHandler(int) Signals.cpp:0:0
 #3 0x00007efcf994d3c0 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x153c0)
 #4 0x00007efcf93ee18b raise (/lib/x86_64-linux-gnu/libc.so.6+0x4618b)
 #5 0x00007efcf93cd859 abort (/lib/x86_64-linux-gnu/libc.so.6+0x25859)
 #6 0x000055a70e316fe6 llvm::report_fatal_error(llvm::Twine const&, bool) (/mnt/data/home/mzw/hipcc_llvm_4.3.0/bin/opt+0x1bbbfe6)
 #7 0x000055a70e317118 (/mnt/data/home/mzw/hipcc_llvm_4.3.0/bin/opt+0x1bbc118)
 #8 0x000055a70dc4b01f (anonymous namespace)::VerifierLegacyPass::runOnFunction(llvm::Function&) Verifier.cpp:0:0
 #9 0x000055a70dbc95d7 llvm::FPPassManager::runOnFunction(llvm::Function&) (/mnt/data/home/mzw/hipcc_llvm_4.3.0/bin/opt+0x146e5d7)
#10 0x000055a70dbc9cf1 llvm::FPPassManager::runOnModule(llvm::Module&) (/mnt/data/home/mzw/hipcc_llvm_4.3.0/bin/opt+0x146ecf1)
#11 0x000055a70dbc894f llvm::legacy::PassManagerImpl::run(llvm::Module&) (/mnt/data/home/mzw/hipcc_llvm_4.3.0/bin/opt+0x146d94f)
#12 0x000055a70cc866eb main (/mnt/data/home/mzw/hipcc_llvm_4.3.0/bin/opt+0x52b6eb)
#13 0x00007efcf93cf0b3 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x270b3)
#14 0x000055a70cd0d51a _start (/mnt/data/home/mzw/hipcc_llvm_4.3.0/bin/opt+0x5b251a)

Any help will be highly appreciated, thank you.

#2
  1. following is my source llvm IR code
; ModuleID = '/mnt/data/home/mzw/workspace/test_space/llvm_test/mark_test//mark_test.hip'
source_filename = "/mnt/data/home/mzw/workspace/test_space/llvm_test/mark_test//mark_test.hip"
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

%"class.std::ios_base::Init" = type { i8 }
%"class.std::basic_ostream" = type { i32 (...)**, %"class.std::basic_ios" }
%"class.std::basic_ios" = type { %"class.std::ios_base", %"class.std::basic_ostream"*, i8, i8, %"class.std::basic_streambuf"*, %"class.std::ctype"*, %"class.std::num_put"*, %"class.std::num_get"* }
%"class.std::ios_base" = type { i32 (...)**, i64, i64, i32, i32, i32, %"struct.std::ios_base::_Callback_list"*, %"struct.std::ios_base::_Words", [8 x %"struct.std::ios_base::_Words"], i32, %"struct.std::ios_base::_Words"*, %"class.std::locale" }
%"struct.std::ios_base::_Callback_list" = type { %"struct.std::ios_base::_Callback_list"*, void (i32, %"class.std::ios_base"*, i32)*, i32, i32 }
%"struct.std::ios_base::_Words" = type { i8*, i64 }
%"class.std::locale" = type { %"class.std::locale::_Impl"* }
%"class.std::locale::_Impl" = type { i32, %"class.std::locale::facet"**, i64, %"class.std::locale::facet"**, i8** }
%"class.std::locale::facet" = type <{ i32 (...)**, i32, [4 x i8] }>
%"class.std::basic_streambuf" = type { i32 (...)**, i8*, i8*, i8*, i8*, i8*, i8*, %"class.std::locale" }
%"class.std::ctype" = type <{ %"class.std::locale::facet.base", [4 x i8], %struct.__locale_struct*, i8, [7 x i8], i32*, i32*, i16*, i8, [256 x i8], [256 x i8], i8, [6 x i8] }>
%"class.std::locale::facet.base" = type <{ i32 (...)**, i32 }>
%struct.__locale_struct = type { [13 x %struct.__locale_data*], i16*, i32*, i32*, [13 x i8*] }
%struct.__locale_data = type opaque
%"class.std::num_put" = type { %"class.std::locale::facet.base", [4 x i8] }
%"class.std::num_get" = type { %"class.std::locale::facet.base", [4 x i8] }
%"class.std::__cxx11::basic_string" = type { %"struct.std::__cxx11::basic_string<char>::_Alloc_hider", i64, %union.anon }
%"struct.std::__cxx11::basic_string<char>::_Alloc_hider" = type { i8* }
%union.anon = type { i64, [8 x i8] }
%"class.std::basic_fstream" = type { %"class.std::basic_iostream.base", %"class.std::basic_filebuf", %"class.std::basic_ios" }
%"class.std::basic_iostream.base" = type { %"class.std::basic_istream.base", %"class.std::basic_ostream.base" }
%"class.std::basic_istream.base" = type { i32 (...)**, i64 }
%"class.std::basic_ostream.base" = type { i32 (...)** }
%"class.std::basic_filebuf" = type { %"class.std::basic_streambuf", %union.pthread_mutex_t, %"class.std::__basic_file", i32, %struct.__mbstate_t, %struct.__mbstate_t, %struct.__mbstate_t, i8*, i64, i8, i8, i8, i8, i8*, i8*, i8, %"class.std::codecvt"*, i8*, i64, i8*, i8* }
%union.pthread_mutex_t = type { %struct.__pthread_mutex_s }
%struct.__pthread_mutex_s = type { i32, i32, i32, i32, i32, i16, i16, %struct.__pthread_internal_list }
%struct.__pthread_internal_list = type { %struct.__pthread_internal_list*, %struct.__pthread_internal_list* }
%"class.std::__basic_file" = type <{ %struct._IO_FILE*, i8, [7 x i8] }>
%struct._IO_FILE = type { i32, i8*, i8*, i8*, i8*, i8*, i8*, i8*, i8*, i8*, i8*, i8*, %struct._IO_marker*, %struct._IO_FILE*, i32, i32, i64, i16, i8, [1 x i8], i8*, i64, %struct._IO_codecvt*, %struct._IO_wide_data*, %struct._IO_FILE*, i8*, i64, i32, [20 x i8] }
%struct._IO_marker = type opaque
%struct._IO_codecvt = type opaque
%struct._IO_wide_data = type opaque
%struct.__mbstate_t = type { i32, %union.anon.0 }
%union.anon.0 = type { i32 }
%"class.std::codecvt" = type { %"class.std::__codecvt_abstract_base.base", %struct.__locale_struct* }
%"class.std::__codecvt_abstract_base.base" = type { %"class.std::locale::facet.base" }

@_ZStL8__ioinit = internal global %"class.std::ios_base::Init" zeroinitializer, align 1
@__dso_handle = external hidden global i8
@.str = private unnamed_addr constant [79 x i8] c"/mnt/data/home/mzw/workspace/test_space/llvm_test/mark_test/dimension_file.txt\00", align 1
@_ZSt4cout = external dso_local global %"class.std::basic_ostream", align 8
@.str.1 = private unnamed_addr constant [29 x i8] c"Cant open the dimension file\00", align 1
@.str.2 = private unnamed_addr constant [2 x i8] c" \00", align 1
@_ZTTSt13basic_fstreamIcSt11char_traitsIcEE = external unnamed_addr constant [10 x i8*], align 8
@llvm.global_ctors = appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 65535, void ()* @_GLOBAL__sub_I_mark_test.hip, i8* null }]

declare dso_local void @_ZNSt8ios_base4InitC1Ev(%"class.std::ios_base::Init"* nonnull dereferenceable(1)) unnamed_addr #0

; Function Attrs: nounwind
declare dso_local void @_ZNSt8ios_base4InitD1Ev(%"class.std::ios_base::Init"* nonnull dereferenceable(1)) unnamed_addr #1

; Function Attrs: nofree nounwind
declare dso_local i32 @__cxa_atexit(void (i8*)*, i8*, i8*) local_unnamed_addr #2

; Function Attrs: noinline uwtable
define dso_local void @_Z21output_dimension_fileiii(i32 %0, i32 %1, i32 %2) local_unnamed_addr #3 personality i8* bitcast (i32 (...)* @__gxx_personality_v0 to i8*) {
  %4 = alloca i64, align 8
  %5 = alloca %"class.std::__cxx11::basic_string", align 8
  %6 = alloca %"class.std::basic_fstream", align 8
  %7 = bitcast %"class.std::__cxx11::basic_string"* %5 to i8*
  call void @llvm.lifetime.start.p0i8(i64 32, i8* nonnull %7) #14
  %8 = getelementptr inbounds %"class.std::__cxx11::basic_string", %"class.std::__cxx11::basic_string"* %5, i64 0, i32 2
  %9 = bitcast %"class.std::__cxx11::basic_string"* %5 to %union.anon**
  store %union.anon* %8, %union.anon** %9, align 8, !tbaa !2
  %10 = bitcast %union.anon* %8 to i8*
  %11 = bitcast i64* %4 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %11) #14
  store i64 78, i64* %4, align 8, !tbaa !7
  %12 = call i8* @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE9_M_createERmm(%"class.std::__cxx11::basic_string"* nonnull dereferenceable(32) %5, i64* nonnull align 8 dereferenceable(8) %4, i64 0)
  %13 = getelementptr inbounds %"class.std::__cxx11::basic_string", %"class.std::__cxx11::basic_string"* %5, i64 0, i32 0, i32 0
  store i8* %12, i8** %13, align 8, !tbaa !9
  %14 = load i64, i64* %4, align 8, !tbaa !7
  %15 = getelementptr inbounds %"class.std::__cxx11::basic_string", %"class.std::__cxx11::basic_string"* %5, i64 0, i32 2, i32 0
  store i64 %14, i64* %15, align 8, !tbaa !11
  call void @llvm.memcpy.p0i8.p0i8.i64(i8* noundef nonnull align 1 dereferenceable(78) %12, i8* noundef nonnull align 1 dereferenceable(78) getelementptr inbounds ([79 x i8], [79 x i8]* @.str, i64 0, i64 0), i64 78, i1 false) #14
  %16 = getelementptr inbounds %"class.std::__cxx11::basic_string", %"class.std::__cxx11::basic_string"* %5, i64 0, i32 1
  store i64 %14, i64* %16, align 8, !tbaa !12
  %17 = getelementptr inbounds i8, i8* %12, i64 %14
  store i8 0, i8* %17, align 1, !tbaa !11
  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %11) #14
  %18 = bitcast %"class.std::basic_fstream"* %6 to i8*
  call void @llvm.lifetime.start.p0i8(i64 528, i8* nonnull %18) #14
  invoke void @_ZNSt13basic_fstreamIcSt11char_traitsIcEEC1ERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode(%"class.std::basic_fstream"* nonnull dereferenceable(264) %6, %"class.std::__cxx11::basic_string"* nonnull align 8 dereferenceable(32) %5, i32 1)
          to label %19 unwind label %36

19:                                               ; preds = %3
  %20 = bitcast %"class.std::basic_fstream"* %6 to i8**
  %21 = load i8*, i8** %20, align 8, !tbaa !13
  %22 = getelementptr i8, i8* %21, i64 -24
  %23 = bitcast i8* %22 to i64*
  %24 = load i64, i64* %23, align 8
  %25 = getelementptr inbounds i8, i8* %18, i64 %24
  %26 = getelementptr inbounds i8, i8* %25, i64 32
  %27 = bitcast i8* %26 to i32*
  %28 = load i32, i32* %27, align 8, !tbaa !15
  %29 = and i32 %28, 5
  %30 = icmp eq i32 %29, 0
  br i1 %30, label %40, label %31

31:                                               ; preds = %19
  %32 = invoke nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(%"class.std::basic_ostream"* nonnull align 8 dereferenceable(8) @_ZSt4cout, i8* nonnull getelementptr inbounds ([29 x i8], [29 x i8]* @.str.1, i64 0, i64 0), i64 28)
          to label %33 unwind label %38

33:                                               ; preds = %31
  %34 = invoke nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(%"class.std::basic_ostream"* nonnull align 8 dereferenceable(8) @_ZSt4cout)
          to label %35 unwind label %38

35:                                               ; preds = %33
  call void @exit(i32 1) #15
  unreachable

36:                                               ; preds = %3
  %37 = landingpad { i8*, i32 }
          cleanup
  br label %125

38:                                               ; preds = %83, %80, %74, %73, %64, %90, %85, %48, %44, %33, %31, %50, %46, %40
  %39 = landingpad { i8*, i32 }
          cleanup
  call void @_ZNSt13basic_fstreamIcSt11char_traitsIcEED1Ev(%"class.std::basic_fstream"* nonnull dereferenceable(264) %6) #14
  br label %125

40:                                               ; preds = %19
  %41 = getelementptr inbounds %"class.std::basic_fstream", %"class.std::basic_fstream"* %6, i64 0, i32 0, i32 1
  %42 = bitcast %"class.std::basic_ostream.base"* %41 to %"class.std::basic_ostream"*
  %43 = invoke nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZNSolsEi(%"class.std::basic_ostream"* nonnull dereferenceable(8) %42, i32 %0)
          to label %44 unwind label %38

44:                                               ; preds = %40
  %45 = invoke nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(%"class.std::basic_ostream"* nonnull align 8 dereferenceable(8) %43, i8* nonnull getelementptr inbounds ([2 x i8], [2 x i8]* @.str.2, i64 0, i64 0), i64 1)
          to label %46 unwind label %38

46:                                               ; preds = %44
  %47 = invoke nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZNSolsEi(%"class.std::basic_ostream"* nonnull dereferenceable(8) %43, i32 %1)
          to label %48 unwind label %38

48:                                               ; preds = %46
  %49 = invoke nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(%"class.std::basic_ostream"* nonnull align 8 dereferenceable(8) %47, i8* nonnull getelementptr inbounds ([2 x i8], [2 x i8]* @.str.2, i64 0, i64 0), i64 1)
          to label %50 unwind label %38

50:                                               ; preds = %48
  %51 = invoke nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZNSolsEi(%"class.std::basic_ostream"* nonnull dereferenceable(8) %47, i32 %2)
          to label %52 unwind label %38

52:                                               ; preds = %50
  %53 = bitcast %"class.std::basic_ostream"* %51 to i8**
  %54 = load i8*, i8** %53, align 8, !tbaa !13
  %55 = getelementptr i8, i8* %54, i64 -24
  %56 = bitcast i8* %55 to i64*
  %57 = load i64, i64* %56, align 8
  %58 = bitcast %"class.std::basic_ostream"* %51 to i8*
  %59 = getelementptr inbounds i8, i8* %58, i64 %57
  %60 = getelementptr inbounds i8, i8* %59, i64 240
  %61 = bitcast i8* %60 to %"class.std::ctype"**
  %62 = load %"class.std::ctype"*, %"class.std::ctype"** %61, align 8, !tbaa !22
  %63 = icmp eq %"class.std::ctype"* %62, null
  br i1 %63, label %64, label %66

64:                                               ; preds = %52
  invoke void @_ZSt16__throw_bad_castv() #16
          to label %65 unwind label %38

65:                                               ; preds = %64
  unreachable

66:                                               ; preds = %52
  %67 = getelementptr inbounds %"class.std::ctype", %"class.std::ctype"* %62, i64 0, i32 8
  %68 = load i8, i8* %67, align 8, !tbaa !25
  %69 = icmp eq i8 %68, 0
  br i1 %69, label %73, label %70

70:                                               ; preds = %66
  %71 = getelementptr inbounds %"class.std::ctype", %"class.std::ctype"* %62, i64 0, i32 9, i64 10
  %72 = load i8, i8* %71, align 1, !tbaa !11
  br label %80

73:                                               ; preds = %66
  invoke void @_ZNKSt5ctypeIcE13_M_widen_initEv(%"class.std::ctype"* nonnull dereferenceable(570) %62)
          to label %74 unwind label %38

74:                                               ; preds = %73
  %75 = bitcast %"class.std::ctype"* %62 to i8 (%"class.std::ctype"*, i8)***
  %76 = load i8 (%"class.std::ctype"*, i8)**, i8 (%"class.std::ctype"*, i8)*** %75, align 8, !tbaa !13
  %77 = getelementptr inbounds i8 (%"class.std::ctype"*, i8)*, i8 (%"class.std::ctype"*, i8)** %76, i64 6
  %78 = load i8 (%"class.std::ctype"*, i8)*, i8 (%"class.std::ctype"*, i8)** %77, align 8
  %79 = invoke signext i8 %78(%"class.std::ctype"* nonnull dereferenceable(570) %62, i8 signext 10)
          to label %80 unwind label %38

80:                                               ; preds = %74, %70
  %81 = phi i8 [ %72, %70 ], [ %79, %74 ]
  %82 = invoke nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZNSo3putEc(%"class.std::basic_ostream"* nonnull dereferenceable(8) %51, i8 signext %81)
          to label %83 unwind label %38

83:                                               ; preds = %80
  %84 = invoke nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZNSo5flushEv(%"class.std::basic_ostream"* nonnull dereferenceable(8) %82)
          to label %85 unwind label %38

85:                                               ; preds = %83
  %86 = getelementptr inbounds %"class.std::basic_fstream", %"class.std::basic_fstream"* %6, i64 0, i32 1
  %87 = invoke %"class.std::basic_filebuf"* @_ZNSt13basic_filebufIcSt11char_traitsIcEE5closeEv(%"class.std::basic_filebuf"* nonnull dereferenceable(240) %86)
          to label %88 unwind label %38

88:                                               ; preds = %85
  %89 = icmp eq %"class.std::basic_filebuf"* %87, null
  br i1 %89, label %90, label %101

90:                                               ; preds = %88
  %91 = load i8*, i8** %20, align 8, !tbaa !13
  %92 = getelementptr i8, i8* %91, i64 -24
  %93 = bitcast i8* %92 to i64*
  %94 = load i64, i64* %93, align 8
  %95 = getelementptr inbounds i8, i8* %18, i64 %94
  %96 = bitcast i8* %95 to %"class.std::basic_ios"*
  %97 = getelementptr inbounds i8, i8* %95, i64 32
  %98 = bitcast i8* %97 to i32*
  %99 = load i32, i32* %98, align 8, !tbaa !15
  %100 = or i32 %99, 4
  invoke void @_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate(%"class.std::basic_ios"* nonnull dereferenceable(264) %96, i32 %100)
          to label %101 unwind label %38

101:                                              ; preds = %88, %90
  %102 = load i32 (...)**, i32 (...)*** bitcast ([10 x i8*]* @_ZTTSt13basic_fstreamIcSt11char_traitsIcEE to i32 (...)***), align 8
  %103 = getelementptr inbounds %"class.std::basic_fstream", %"class.std::basic_fstream"* %6, i64 0, i32 0, i32 0, i32 0
  store i32 (...)** %102, i32 (...)*** %103, align 8, !tbaa !13
  %104 = load i32 (...)**, i32 (...)*** bitcast (i8** getelementptr inbounds ([10 x i8*], [10 x i8*]* @_ZTTSt13basic_fstreamIcSt11char_traitsIcEE, i64 0, i64 8) to i32 (...)***), align 8
  %105 = getelementptr i32 (...)*, i32 (...)** %102, i64 -3
  %106 = bitcast i32 (...)** %105 to i64*
  %107 = load i64, i64* %106, align 8
  %108 = getelementptr inbounds i8, i8* %18, i64 %107
  %109 = bitcast i8* %108 to i32 (...)***
  store i32 (...)** %104, i32 (...)*** %109, align 8, !tbaa !13
  %110 = load i32 (...)**, i32 (...)*** bitcast (i8** getelementptr inbounds ([10 x i8*], [10 x i8*]* @_ZTTSt13basic_fstreamIcSt11char_traitsIcEE, i64 0, i64 9) to i32 (...)***), align 8
  %111 = getelementptr inbounds %"class.std::basic_fstream", %"class.std::basic_fstream"* %6, i64 0, i32 0, i32 1, i32 0
  store i32 (...)** %110, i32 (...)*** %111, align 8, !tbaa !13
  call void @_ZNSt13basic_filebufIcSt11char_traitsIcEED2Ev(%"class.std::basic_filebuf"* nonnull dereferenceable(240) %86) #14
  %112 = load i32 (...)**, i32 (...)*** bitcast (i8** getelementptr inbounds ([10 x i8*], [10 x i8*]* @_ZTTSt13basic_fstreamIcSt11char_traitsIcEE, i64 0, i64 2) to i32 (...)***), align 8
  store i32 (...)** %112, i32 (...)*** %103, align 8, !tbaa !13
  %113 = load i32 (...)**, i32 (...)*** bitcast (i8** getelementptr inbounds ([10 x i8*], [10 x i8*]* @_ZTTSt13basic_fstreamIcSt11char_traitsIcEE, i64 0, i64 3) to i32 (...)***), align 8
  %114 = getelementptr i32 (...)*, i32 (...)** %112, i64 -3
  %115 = bitcast i32 (...)** %114 to i64*
  %116 = load i64, i64* %115, align 8
  %117 = getelementptr inbounds i8, i8* %18, i64 %116
  %118 = bitcast i8* %117 to i32 (...)***
  store i32 (...)** %113, i32 (...)*** %118, align 8, !tbaa !13
  %119 = getelementptr inbounds %"class.std::basic_fstream", %"class.std::basic_fstream"* %6, i64 0, i32 0, i32 0, i32 1
  store i64 0, i64* %119, align 8, !tbaa !27
  %120 = getelementptr inbounds %"class.std::basic_fstream", %"class.std::basic_fstream"* %6, i64 0, i32 2, i32 0
  call void @_ZNSt8ios_baseD2Ev(%"class.std::ios_base"* nonnull dereferenceable(216) %120) #14
  call void @llvm.lifetime.end.p0i8(i64 528, i8* nonnull %18) #14
  %121 = load i8*, i8** %13, align 8, !tbaa !9
  %122 = icmp eq i8* %121, %10
  br i1 %122, label %124, label %123

123:                                              ; preds = %101
  call void @_ZdlPv(i8* %121) #14
  br label %124

124:                                              ; preds = %101, %123
  call void @llvm.lifetime.end.p0i8(i64 32, i8* nonnull %7) #14
  ret void

125:                                              ; preds = %38, %36
  %126 = phi { i8*, i32 } [ %39, %38 ], [ %37, %36 ]
  call void @llvm.lifetime.end.p0i8(i64 528, i8* nonnull %18) #14
  %127 = load i8*, i8** %13, align 8, !tbaa !9
  %128 = icmp eq i8* %127, %10
  br i1 %128, label %130, label %129

129:                                              ; preds = %125
  call void @_ZdlPv(i8* %127) #14
  br label %130

130:                                              ; preds = %129, %125
  call void @llvm.lifetime.end.p0i8(i64 32, i8* nonnull %7) #14
  resume { i8*, i32 } %126
}

; Function Attrs: argmemonly nofree nosync nounwind willreturn
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #4

declare dso_local i32 @__gxx_personality_v0(...)

; Function Attrs: argmemonly nofree nosync nounwind willreturn
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #4

; Function Attrs: uwtable
declare dso_local void @_ZNSt13basic_fstreamIcSt11char_traitsIcEEC1ERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode(%"class.std::basic_fstream"* nonnull dereferenceable(264), %"class.std::__cxx11::basic_string"* nonnull align 8 dereferenceable(32), i32) unnamed_addr #5 align 2

; Function Attrs: inlinehint uwtable mustprogress
declare dso_local nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(%"class.std::basic_ostream"* nonnull align 8 dereferenceable(8)) local_unnamed_addr #6

; Function Attrs: noreturn nounwind
declare dso_local void @exit(i32) local_unnamed_addr #7

declare dso_local nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZNSolsEi(%"class.std::basic_ostream"* nonnull dereferenceable(8), i32) local_unnamed_addr #0

; Function Attrs: nounwind uwtable
declare dso_local void @_ZNSt13basic_fstreamIcSt11char_traitsIcEED1Ev(%"class.std::basic_fstream"* nonnull dereferenceable(264)) unnamed_addr #8 align 2

; Function Attrs: noinline uwtable mustprogress
define dso_local void @_Z6GemmExiiiiii(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5) local_unnamed_addr #9 {
  %7 = alloca i8*, align 8
  %8 = alloca i8*, align 8
  %9 = alloca i8*, align 8
  %10 = alloca i8*, align 8
  %11 = alloca i8*, align 8
  %12 = alloca i8*, align 8
  %13 = bitcast i8** %7 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %13) #14
  store i8* null, i8** %7, align 8, !tbaa !29
  %14 = bitcast i8** %8 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %14) #14
  store i8* null, i8** %8, align 8, !tbaa !29
  %15 = bitcast i8** %9 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %15) #14
  store i8* null, i8** %9, align 8, !tbaa !29
  %16 = bitcast i8** %10 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %16) #14
  store i8* null, i8** %10, align 8, !tbaa !29
  %17 = bitcast i8** %11 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %17) #14
  store i8* null, i8** %11, align 8, !tbaa !29
  %18 = bitcast i8** %12 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %18) #14
  store i8* null, i8** %12, align 8, !tbaa !29
  %19 = sext i32 %0 to i64
  %20 = shl nsw i64 %19, 2
  %21 = call i32 @hipMalloc(i8** nonnull %7, i64 %20)
  %22 = sext i32 %1 to i64
  %23 = shl nsw i64 %22, 2
  %24 = call i32 @hipMalloc(i8** nonnull %8, i64 %23)
  %25 = sext i32 %2 to i64
  %26 = shl nsw i64 %25, 2
  %27 = call i32 @hipMalloc(i8** nonnull %9, i64 %26)
  %28 = sext i32 %3 to i64
  %29 = shl nsw i64 %28, 2
  %30 = call i32 @hipMalloc(i8** nonnull %10, i64 %29)
  %31 = sext i32 %4 to i64
  %32 = shl nsw i64 %31, 2
  %33 = call i32 @hipMalloc(i8** nonnull %11, i64 %32)
  %34 = sext i32 %5 to i64
  %35 = shl nsw i64 %34, 2
  %36 = call i32 @hipMalloc(i8** nonnull %12, i64 %35)
  %37 = load i8*, i8** %7, align 8, !tbaa !29
  %38 = call i32 @hipFree(i8* %37)
  %39 = load i8*, i8** %8, align 8, !tbaa !29
  %40 = call i32 @hipFree(i8* %39)
  %41 = load i8*, i8** %9, align 8, !tbaa !29
  %42 = call i32 @hipFree(i8* %41)
  %43 = load i8*, i8** %10, align 8, !tbaa !29
  %44 = call i32 @hipFree(i8* %43)
  %45 = load i8*, i8** %11, align 8, !tbaa !29
  %46 = call i32 @hipFree(i8* %45)
  %47 = load i8*, i8** %12, align 8, !tbaa !29
  %48 = call i32 @hipFree(i8* %47)
  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %18) #14
  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %17) #14
  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %16) #14
  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %15) #14
  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %14) #14
  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %13) #14
  ret void
}

declare dso_local i32 @hipMalloc(i8**, i64) local_unnamed_addr #0

declare dso_local i32 @hipFree(i8*) local_unnamed_addr #0

; Function Attrs: norecurse uwtable mustprogress
define dso_local i32 @main() local_unnamed_addr #10 {
  %1 = alloca i8*, align 8
  tail call void @_Z6GemmExiiiiii(i32 1, i32 2, i32 3, i32 4, i32 5, i32 6)
  %2 = bitcast i8** %1 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %2) #14
  store i8* null, i8** %1, align 8, !tbaa !29
  %3 = call i32 @hipMalloc(i8** nonnull %1, i64 4)
  %4 = load i8*, i8** %1, align 8, !tbaa !29
  %5 = call i32 @hipFree(i8* %4)
  call void @_Z21output_dimension_fileiii(i32 1, i32 2, i32 3)
  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %2) #14
  ret i32 0
}

; Function Attrs: nobuiltin nounwind
declare dso_local void @_ZdlPv(i8*) local_unnamed_addr #11

declare dso_local i8* @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE9_M_createERmm(%"class.std::__cxx11::basic_string"* nonnull dereferenceable(32), i64* nonnull align 8 dereferenceable(8), i64) local_unnamed_addr #0

; Function Attrs: argmemonly nofree nosync nounwind willreturn
declare void @llvm.memcpy.p0i8.p0i8.i64(i8* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #4

; Function Attrs: nounwind uwtable
declare dso_local void @_ZNSt13basic_filebufIcSt11char_traitsIcEED2Ev(%"class.std::basic_filebuf"* nonnull dereferenceable(240)) unnamed_addr #8 align 2

declare dso_local void @_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate(%"class.std::basic_ios"* nonnull dereferenceable(264), i32) local_unnamed_addr #0

declare dso_local %"class.std::basic_filebuf"* @_ZNSt13basic_filebufIcSt11char_traitsIcEE5closeEv(%"class.std::basic_filebuf"* nonnull dereferenceable(240)) local_unnamed_addr #0

; Function Attrs: nounwind
declare dso_local void @_ZNSt8ios_baseD2Ev(%"class.std::ios_base"* nonnull dereferenceable(216)) unnamed_addr #1

declare dso_local nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(%"class.std::basic_ostream"* nonnull align 8 dereferenceable(8), i8*, i64) local_unnamed_addr #0

declare dso_local nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZNSo3putEc(%"class.std::basic_ostream"* nonnull dereferenceable(8), i8 signext) local_unnamed_addr #0

declare dso_local nonnull align 8 dereferenceable(8) %"class.std::basic_ostream"* @_ZNSo5flushEv(%"class.std::basic_ostream"* nonnull dereferenceable(8)) local_unnamed_addr #0

; Function Attrs: noreturn
declare dso_local void @_ZSt16__throw_bad_castv() local_unnamed_addr #12

declare dso_local void @_ZNKSt5ctypeIcE13_M_widen_initEv(%"class.std::ctype"* nonnull dereferenceable(570)) local_unnamed_addr #0

; Function Attrs: uwtable
define internal amdgpu_kernel void @_GLOBAL__sub_I_mark_test.hip() #13 section ".text.startup" {
  tail call void @_ZNSt8ios_base4InitC1Ev(%"class.std::ios_base::Init"* nonnull dereferenceable(1) @_ZStL8__ioinit)
  %1 = tail call i32 @__cxa_atexit(void (i8*)* bitcast (void (%"class.std::ios_base::Init"*)* @_ZNSt8ios_base4InitD1Ev to void (i8*)*), i8* getelementptr inbounds (%"class.std::ios_base::Init", %"class.std::ios_base::Init"* @_ZStL8__ioinit, i64 0, i32 0), i8* nonnull @__dso_handle) #14
  ret void
}

attributes #0 = { "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #1 = { nounwind "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #2 = { nofree nounwind }
attributes #3 = { noinline uwtable "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #4 = { argmemonly nofree nosync nounwind willreturn }
attributes #5 = { uwtable "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #6 = { inlinehint uwtable mustprogress "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #7 = { noreturn nounwind "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #8 = { nounwind uwtable "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #9 = { noinline uwtable mustprogress "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #10 = { norecurse uwtable mustprogress "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #11 = { nobuiltin nounwind "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #12 = { noreturn "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #13 = { uwtable "device-init" "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #14 = { nounwind }
attributes #15 = { noreturn nounwind }
attributes #16 = { noreturn }

!llvm.module.flags = !{!0}
!llvm.ident = !{!1}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{!"clang version 13.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-4.3.0 21295 f2943f684437d2c1143a56e418d29fc6b3314072)"}
!2 = !{!3, !4, i64 0}
!3 = !{!"_ZTSNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_Alloc_hiderE", !4, i64 0}
!4 = !{!"any pointer", !5, i64 0}
!5 = !{!"omnipotent char", !6, i64 0}
!6 = !{!"Simple C++ TBAA"}
!7 = !{!8, !8, i64 0}
!8 = !{!"long", !5, i64 0}
!9 = !{!10, !4, i64 0}
!10 = !{!"_ZTSNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE", !3, i64 0, !8, i64 8, !5, i64 16}
!11 = !{!5, !5, i64 0}
!12 = !{!10, !8, i64 8}
!13 = !{!14, !14, i64 0}
!14 = !{!"vtable pointer", !6, i64 0}
!15 = !{!16, !18, i64 32}
!16 = !{!"_ZTSSt8ios_base", !8, i64 8, !8, i64 16, !17, i64 24, !18, i64 28, !18, i64 32, !4, i64 40, !19, i64 48, !5, i64 64, !20, i64 192, !4, i64 200, !21, i64 208}
!17 = !{!"_ZTSSt13_Ios_Fmtflags", !5, i64 0}
!18 = !{!"_ZTSSt12_Ios_Iostate", !5, i64 0}
!19 = !{!"_ZTSNSt8ios_base6_WordsE", !4, i64 0, !8, i64 8}
!20 = !{!"int", !5, i64 0}
!21 = !{!"_ZTSSt6locale", !4, i64 0}
!22 = !{!23, !4, i64 240}
!23 = !{!"_ZTSSt9basic_iosIcSt11char_traitsIcEE", !4, i64 216, !5, i64 224, !24, i64 225, !4, i64 232, !4, i64 240, !4, i64 248, !4, i64 256}
!24 = !{!"bool", !5, i64 0}
!25 = !{!26, !5, i64 56}
!26 = !{!"_ZTSSt5ctypeIcE", !4, i64 16, !24, i64 24, !4, i64 32, !4, i64 40, !4, i64 48, !5, i64 56, !5, i64 57, !5, i64 313, !5, i64 569}
!27 = !{!28, !8, i64 8}
!28 = !{!"_ZTSSi", !8, i64 8}
!29 = !{!4, !4, i64 0}
#3

I found that if I use

 FunctionType * output_func_type = output_dimension_func_ptr->getFunctionType();

to get the void function return type, it’s ok to go. But what’s the difference between this and

FunctionType::get(FunctionType::getVoidTy(M.getContext()),false)
#4

FunctionType::get(FunctionType::getVoidTy(M.getContext()),false) gives you void(), that is a function returning void and taking no arguments. But the function in question seems to take three i32s.

You’d need something like

Type Args[] = { Type::getInt32Ty(Ctx), Type::getInt32Ty(Ctx), Type::getInt32Ty(Ctx) };
FunctionType::get(Type::getVoidTy(Ctx), Args, false);

to synthesize the right type from scratch.

#6

So does it mean that I can use the returned value of FunctionType::get(Type::getVoidTy(Ctx), Args, false); to create/declare a function? Otherwise I don’t see the benefit of calling such a function.

#7

Creating a function is probably the main use for it, yes. Possibly also casting a function pointer if you end up storing it as i8* in some structure for whatever reason.

But normally you’ll have the type around via other means, like the Function you seem to be using (at least for now, the “opaque pointers” project will reduce that, though not in your particular case I think).

#8

Thank you so much. That really helps me a lot.