diff options
Diffstat (limited to 'loop_cuda.cc')
| -rw-r--r-- | loop_cuda.cc | 2123 |
1 files changed, 0 insertions, 2123 deletions
diff --git a/loop_cuda.cc b/loop_cuda.cc deleted file mode 100644 index a23990d..0000000 --- a/loop_cuda.cc +++ /dev/null @@ -1,2123 +0,0 @@ -/***************************************************************************** - Copyright (C) 2009 University of Utah - All Rights Reserved. - - Purpose: - Cudaize methods - - Notes: - - History: - 1/7/10 Created by Gabe Rudy by migrating code from loop.cc - 31/1/11 Modified by Protonu Basu -*****************************************************************************/ - -#include <code_gen/code_gen.h> -#include <code_gen/CG_stringBuilder.h> -#include <code_gen/output_repr.h> -#include <code_gen/CG_outputRepr.h> -#include "loop_cuda.hh" -#include "loop.hh" -#include <math.h> -#include <useful.h> -#include "omegatools.hh" -#include "ir_cudasuif.hh" -#include "ir_suif.hh" -#include "ir_suif_utils.hh" -#include "chill_error.hh" -#include <vector> - -using namespace omega; -char *k_cuda_texture_memory; //protonu--added to track texture memory type -char *k_cuda_constant_memory; //protonu--added to track constant memory type -//extern char *omega::k_cuda_texture_memory; //protonu--added to track texture memory type -extern char *omega::k_ocg_comment; - - -static int cudaDebug; -class CudaStaticInit{ public: CudaStaticInit(){ cudaDebug=0; //Change this to 1 for debug -}}; -static CudaStaticInit junkInitInstance__; - - - -std::string& upcase(std::string& s) -{ - for(int i=0; i<s.size(); i++) - s[i] = toupper(s[i]); - return s; -} - -void printVs(const std::vector<std::string>& curOrder){ - if(!cudaDebug) return; - for(int i=0; i<curOrder.size(); i++){ - if(i>0) - printf(","); - printf("%s", curOrder[i].c_str()); - } - printf("\n"); -} - -void printVS(const std::vector<std::string>& curOrder){ - //if(!cudaDebug) return; - for(int i=0; i<curOrder.size(); i++){ - if(i>0) - printf(","); - printf("%s", curOrder[i].c_str()); - } - printf("\n"); -} - -LoopCuda::~LoopCuda() { - const int m = stmt.size(); - for (int i = 0; i < m; i++) - stmt[i].code->clear(); -} - -bool LoopCuda::symbolExists(std::string s){ - if(symtab->lookup_sym(s.c_str(), SYM_VAR, false)) - return true; - if(globals->lookup_sym(s.c_str(), SYM_VAR, false)) - return true; - for(int i=0; i<idxNames.size(); i++) - for(int j=0; j<idxNames[i].size(); j++) - if(strcmp(idxNames[i][j].c_str(), s.c_str()) == 0) - return true; - return false; -} - -void LoopCuda::addSync(int stmt_num, std::string idxName) -{ - //we store these and code-gen inserts sync to omega comments where stmt - //in loop that has idxName being generated - syncs.push_back(make_pair(stmt_num,idxName)); -} - -void LoopCuda::renameIndex(int stmt_num, std::string idx, std::string newName) -{ - int level = findCurLevel(stmt_num, idx); - if(idxNames.size() <= stmt_num || idxNames[stmt_num].size() < level) - throw std::runtime_error("Invalid statment number of index"); - idxNames[stmt_num][level-1] = newName.c_str(); -} - - - -enum Type{ Int }; - -struct VarDefs{ - std::string name; - std::string secondName; - operand size_expr; //array size as an expression (can be a product of other variables etc) - type_node * type; - var_sym* in_data; //Variable of array to copy data in from (before kernel call) - var_sym* out_data; //Variable of array to copy data out to (after kernel call) - int size_2d; //-1 if linearized, the constant size N, of a NxN 2D array otherwise - bool tex_mapped; //protonu-- true if this variable will be texture mapped, so no need to pass it as a argument - bool cons_mapped; //protonu-- true if this variable will be constant mem mapped, so no need to pass it as a argument - std::string original_name; //this is such a hack, to store the original name, to store a table to textures used - int var_ref_size ; -}; - -tree_node_list* wrapInIfFromMinBound(tree_node_list* then_part, tree_for* loop, base_symtab* symtab, var_sym* bound_sym) -{ - tree_node_list* ub = loop->ub_list(); - tree_node_list_iter upli(ub); - while(!upli.is_empty()){ - tree_node *node = upli.step(); - if(node->kind() == TREE_INSTR && ((tree_instr*)node)->instr()->format() == inf_rrr) - { - in_rrr* ins = (in_rrr*)((tree_instr*)node)->instr(); - //expect the structure: cpy( _ = min(grab_me, _)) - if(ins->opcode() == io_cpy && ins->src1_op().is_instr()){ - ins = (in_rrr*)ins->src1_op().instr(); - if(ins->opcode() == io_min){ - tree_node_list* tnl = new tree_node_list; - tnl->append(if_node(symtab, fold_sle(operand(bound_sym), ins->src1_op().instr()->clone()), then_part)); - return tnl; - } - } - } - } - return then_part; //Failed to go to proper loop level -} - -/** - * This would be better if it was done by a CHiLL xformation instead of at codegen - * - * state: - * for(...) - * for(...) - * cur_body - * stmt1 - * - * stm1 is in-between two loops that are going to be reduced. The - * solution is to put stmt1 at the end of cur_body but conditionally run - * in on the last step of the for loop. - * - * A CHiLL command that would work better: - * - * for(...) - * stmt0 - * for(for i=0; i<n; i++) - * cur_body - * stmt1 - * => - * for(...) - * for(for i=0; i<n; i++) - * if(i==0) stmt0 - * cur_body - * if(i==n-1) stmt1 - */ - -std::vector<tree_for*> findCommentedFors(const char* index, tree_node_list* tnl){ - std::vector<tree_for *> result; - - tree_node_list_iter iter(tnl); - bool next_loop_ok = false; - while (!iter.is_empty()) { - tree_node *tn = iter.step(); - if (tn->kind() == TREE_INSTR && ((tree_instr*)tn)->instr()->opcode() == io_mrk) - { - instruction* inst = ((tree_instr*)tn)->instr(); - std::string comment; - if ((inst->peek_annote(k_ocg_comment) != NULL)) - { - immed_list *data = (immed_list *)(inst->peek_annote(k_ocg_comment)); - immed_list_iter data_iter(data); - if(!data_iter.is_empty()){ - immed first_immed = data_iter.step(); - if(first_immed.kind() == im_string) - comment = first_immed.string(); - } - } - if(comment.find("~cuda~") != std::string::npos - && comment.find("preferredIdx: ") != std::string::npos){ - std::string idx = comment.substr(comment.find("preferredIdx: ")+14,std::string::npos); - if(idx.find(" ") != std::string::npos) - idx = idx.substr(0,idx.find(" ")); - if(strcmp(idx.c_str(),index) == 0) - next_loop_ok = true; - } - } - if (tn->kind() == TREE_FOR){ - if(next_loop_ok){ - //printf("found loop %s\n", static_cast<tree_for *>(tn)->index()->name()); - result.push_back(static_cast<tree_for *>(tn)); - } - else{ - //printf("looking down for loop %s\n", static_cast<tree_for *>(tn)->index()->name()); - std::vector<tree_for*> t = findCommentedFors(index, static_cast<tree_for *>(tn)->body()); - std::copy(t.begin(), t.end(), back_inserter(result)); - } - next_loop_ok = false; - } - if (tn->kind() == TREE_IF) { - //printf("looking down if\n"); - tree_if *tni = static_cast<tree_if *>(tn); - std::vector<tree_for*> t = findCommentedFors(index, tni->then_part()); - std::copy(t.begin(), t.end(), back_inserter(result)); - } - } - - return result; -} - -tree_node_list* forReduce(tree_for* loop, var_sym* reduceIndex, proc_symtab* proc_syms) -{ - //We did the replacements all at once with recursiveFindPreferedIdxs - //replacements r; - //r.oldsyms.append(loop->index()); - //r.newsyms.append(reduceIndex); - //tree_for* new_loop = (tree_for*)loop->clone_helper(&r, true); - tree_for* new_loop = loop; - - //return body one loops in - tree_node_list* tnl = loop_body_at_level(new_loop, 1); - //wrap in conditional if necessary - tnl = wrapInIfFromMinBound(tnl, new_loop, proc_syms, reduceIndex); - return tnl; -} - -void recursiveFindRefs(tree_node_list* code, proc_symtab* proc_syms, replacements* r) -{ - if(code->parent() && code->scope()->is_block()) - ((block_symtab*)code->scope())->find_exposed_refs(proc_syms, r); - tree_node_list_iter tnli(code); - while (!tnli.is_empty()) { - tree_node *node = tnli.step(); - //printf("node kind: %d\n", node->kind()); - if(node->is_instr()) - { - tree_instr* t_instr = (tree_instr*)node; - t_instr->find_exposed_refs(proc_syms, r); - } - if(node->is_block()){ - recursiveFindRefs(static_cast<tree_block *>(node)->body(), proc_syms, r); - } - else if(node->is_for()){ - tree_for* tn_for = static_cast<tree_for *>(node); - //Find refs in statemetns and body - tn_for->find_exposed_refs(proc_syms, r); - //recursiveFindRefs(tn_for->body(), proc_syms, r); - } - } -} - -tree_node_list* recursiveFindReplacePreferedIdxs(tree_node_list* code, proc_symtab* proc_syms, - proc_sym* cudaSync, func_type* unkown_func, - std::map<std::string, var_sym*>& loop_idxs) -{ - tree_node_list* tnl = new tree_node_list; - tree_node_list_iter tnli(code); - var_sym* idxSym=0; - bool sync = false; - std::vector<tree_node*> r1; - std::vector<tree_node_list*> r2; - while (!tnli.is_empty()) { - tree_node *node = tnli.step(); - //printf("node kind: %d\n", node->kind()); - if(node->is_instr()) - { - if(((tree_instr*)node)->instr()->format() == inf_rrr){ - in_rrr* inst = (in_rrr*)((tree_instr*)node)->instr(); - if(inst->opcode() == io_mrk){ - std::string comment; - if ((inst->peek_annote(k_ocg_comment) != NULL)) - { - immed_list *data = (immed_list *)(inst->peek_annote(k_ocg_comment)); - immed_list_iter data_iter(data); - if(!data_iter.is_empty()){ - immed first_immed = data_iter.step(); - if(first_immed.kind() == im_string) - comment = first_immed.string(); - } - } - if(comment.find("~cuda~") != std::string::npos - && comment.find("preferredIdx: ") != std::string::npos){ - std::string idx = comment.substr(comment.find("preferredIdx: ")+14,std::string::npos); - if(idx.find(" ") != std::string::npos) - idx = idx.substr(0,idx.find(" ")); - //printf("sym_tab preferred index: %s\n", idx.c_str()); - if(loop_idxs.find(idx) != loop_idxs.end()) - idxSym = loop_idxs.find(idx)->second; - //Get the proc variable sybol for this preferred index - if(idxSym == 0){ - idxSym = (var_sym*)proc_syms->lookup_sym(idx.c_str(), SYM_VAR, false); - //printf("idx not found: lookup %p\n", idxSym); - if(!idxSym){ - idxSym = new var_sym(type_s32, (char*)idx.c_str()); - proc_syms->add_sym(idxSym); - //printf("idx created and inserted\n"); - } - //Now insert into our map for future - loop_idxs.insert(make_pair(idx, idxSym)); - } - //See if we have a sync as well - if(comment.find("sync") != std::string::npos){ - //printf("Inserting sync after current block\n"); - sync = true; - } - } - } - } - tnl->append(node); - } - else if(node->is_block()){ - tree_block* b = static_cast<tree_block *>(node); - b->set_body(recursiveFindReplacePreferedIdxs(b->body(), proc_syms, cudaSync, unkown_func, loop_idxs)); - tnl->append(b); - } - else if(node->is_for()){ - tree_for* tn_for = static_cast<tree_for *>(node); - if(idxSym){ - //Replace the current tn_for's index variable with idxSym - //printf("replacing sym %s -> %s\n", tn_for->index()->name(), idxSym->name()); - replacements r; - r.oldsyms.append(tn_for->index()); - r.newsyms.append(idxSym); - tree_for* new_loop = (tree_for*)tn_for->clone_helper(&r, true); - idxSym = 0; //Reset for more loops in this tnl - new_loop->set_body(recursiveFindReplacePreferedIdxs(new_loop->body(), proc_syms, cudaSync, unkown_func, loop_idxs)); - tnl->append(new_loop); - - if(sync){ - in_cal *the_call = - new in_cal(type_s32, operand(), operand(new in_ldc(unkown_func->ptr_to(), operand(), immed(cudaSync))), 0); - tnl->append(new tree_instr(the_call)); - //tnl->print(); - sync = true; - } - }else{ - tn_for->set_body(recursiveFindReplacePreferedIdxs(tn_for->body(), proc_syms, cudaSync, unkown_func, loop_idxs)); - tnl->append(tn_for); - } - }else if (node->kind() == TREE_IF) { - tree_if *tni = static_cast<tree_if *>(node); - tni->set_then_part(recursiveFindReplacePreferedIdxs(tni->then_part(), proc_syms, cudaSync, unkown_func, loop_idxs)); - tnl->append(tni); - } - } - //Do this after the loop to not screw up the pointer interator - /* - for(int i=0; i<r1.size(); i++){ - swap_node_for_node_list(r1[i],r2[i]); - }*/ - return tnl; -} - -// loop_vars -> array references -// loop_idxs -> <idx_name,idx_sym> map for when we encounter a loop with a different preferredIndex -// dim_vars -> out param, fills with <old,new> var_sym pair for 2D array dimentions (messy stuff) -tree_node_list* swapVarReferences(tree_node_list* code, replacements* r, CG_suifBuilder *ocg, - std::map<std::string, var_sym*>& loop_vars, - proc_symtab *proc_syms, - std::vector< std::pair<var_sym*,var_sym*> >& dim_vars) -{ - //Iterate over every expression, looking up each variable and type - //reference used and possibly replacing it or adding it to our symbol - //table - // - //We use the built-in cloning helper methods to seriously help us with this! - - //Need to do a recursive mark - recursiveFindRefs(code, proc_syms, r); - - - //We can't rely on type_node->clone() to do the heavy lifting when the - //old type is a two dimentional array with variable upper bounds as - //that requires creating and saveing variable references to the upper - //bounds. So we do one pass over the oldtypes doing this type of - //conversion, putting results in the fixed_types map for a second pass - //to pick up. - std::map<type_node*,type_node*> fixed_types; //array_types needing their upper bound installed - type_node_list_iter tlip(&r->oldtypes); - while(!tlip.is_empty()) - { - type_node* old_tn = tlip.step(); - type_node* new_tn = 0; - type_node* base_type = old_tn; - std::vector< std::pair<var_sym*, type_node*> > variable_upper_bouneds; - if(old_tn->is_ptr()){ - while (base_type->is_array() || base_type->is_ptr()) { - if (base_type->is_array()){ - array_bound ub = ((array_type*)base_type)->upper_bound(); - if(ub.is_variable()){ - var_sym* old_ub = (var_sym*)ub.variable(); - var_sym *new_ub = proc_syms->new_unique_var(type_s32); - dim_vars.push_back(std::pair<var_sym* , var_sym*>(old_ub, new_ub)); - variable_upper_bouneds.push_back( std::pair<var_sym*, type_node*>(new_ub, base_type) ); - } - base_type = static_cast<array_type *>(base_type)->elem_type(); - } - else if (base_type->is_ptr()) - base_type = static_cast<ptr_type *>(base_type)->ref_type(); - } - } - for (int i = variable_upper_bouneds.size()-1; i >= 0; i--) { - var_sym *var_ub = variable_upper_bouneds[i].first; - type_node* old_tn = variable_upper_bouneds[i].second; - if(new_tn == 0) - new_tn = new array_type(base_type, array_bound(1), array_bound(var_ub)); - else - new_tn = new array_type(new_tn, array_bound(1), array_bound(var_ub)); - proc_syms->add_type(new_tn); - fixed_types.insert(std::pair<type_node*,type_node*>(old_tn, new_tn)); - } - if(new_tn){ - if(old_tn->is_ptr()){ - new_tn = new ptr_type(new_tn); - proc_syms->add_type(new_tn); - } - fixed_types.insert(std::pair<type_node*,type_node*>(old_tn, new_tn)); - } - } - - //Quickly look for modifiers on our our array types (__shared__ float [][]) - type_node_list_iter tliq(&r->oldtypes); - while(!tliq.is_empty()) - { - type_node* old_tn = tliq.step(); - if(old_tn->is_modifier()){ - type_node* base_type = static_cast<modifier_type *>(old_tn)->base(); - if(fixed_types.find(base_type) != fixed_types.end()){ - type_node* fixed_base = (*fixed_types.find(base_type)).second; - //printf("Fix modifier with fixed base\n"); - //This should work to copy over the annotations, but apparently doesn't work so well - type_node* new_tn = new modifier_type(static_cast<modifier_type*>(old_tn)->op(), fixed_base); - old_tn->copy_annotes(new_tn); - fixed_types.insert(std::pair<type_node*,type_node*>(old_tn, new_tn)); - } - } - } - - //Run through the types and create entries in r->newtypes but don't install - type_node_list_iter tli(&r->oldtypes); - while(!tli.is_empty()) - { - type_node* old_tn = tli.step(); - type_node* new_tn = 0; - - //If we recorded this as fixed by our special case, use that type - //instead of cloning. - if(fixed_types.find(old_tn) != fixed_types.end()){ - new_tn = (*fixed_types.find(old_tn)).second; - //printf("Reusing fixed typ %u: ", new_tn->type_id()); - }else{ - new_tn = old_tn->clone(); - //printf("Cloning type %u: ", old_tn->type_id()); - } - new_tn = proc_syms->install_type(new_tn); - - //Ok, there is a weird case where an array type that has var_sym as - //their upper bounds can't be covered fully in this loop or the - //var_sym loop, so we need special code. - /* - if(old_tn->op() == TYPE_PTR && ((ptr_type*)old_tn)->ref_type()->op() == TYPE_ARRAY){ - array_type* outer_array = (array_type*)((ptr_type*)old_tn)->ref_type(); - array_bound ub = outer_array->upper_bound(); - if(ub.is_variable()){ - var_sym* old_ub = (var_sym*)ub.variable(); - var_sym* new_ub = (var_sym*)((array_type*)((ptr_type*)new_tn)->ref_type())->upper_bound().variable(); - //r->oldsyms.append(old_ub); - fix_ub.insert(std::pair<var_sym*,array_type*>(old_ub, (array_type*)((ptr_type*)new_tn)->ref_type())); - dim_vars.push_back(std::pair<var_sym* , var_sym*>(old_ub, new_ub)); - printf("array var_sym: %p\n", new_ub); - } - if(outer_array->elem_type()->op() == TYPE_ARRAY) - { - array_type* inner_array = (array_type*)outer_array->elem_type(); - array_bound ub = inner_array->upper_bound(); - if(ub.is_variable()){ - var_sym* old_ub = (var_sym*)ub.variable(); - var_sym* new_ub = (var_sym*)((array_type*)((array_type*)((ptr_type*)new_tn)->ref_type())->elem_type())->upper_bound().variable(); - dim_vars.push_back(std::pair<var_sym* , var_sym*>(old_ub, new_ub)); - printf("array var_sym: %p\n", new_ub); - //r->oldsyms.append(old_ub); - fix_ub.insert(std::pair<var_sym*,array_type*>(old_ub, (array_type*)((array_type*)((ptr_type*)new_tn)->ref_type())->elem_type())); - } - } - } - */ - r->newtypes.append(new_tn); - } - - //printf("proc_syms symbol run through\n"); - //proc_syms->print(); - - //Run through the syms creating new copies - sym_node_list_iter snli(&r->oldsyms); - while(!snli.is_empty()) - { - sym_node *old_sn = snli.step(); - - if(loop_vars.count(std::string(old_sn->name())) > 0) - { - r->newsyms.append(loop_vars[std::string(old_sn->name())]); - //printf("def exists: %s\n", old_sn->name()); - }else{ - sym_node *new_sn = old_sn->copy(); - if(new_sn->is_var()){ - var_sym* var = (var_sym*)new_sn; - type_node* new_type = var->type()->clone_helper(r); - - //TODO: Have a tagged list of variables to make shared - //Make local 2D arrays __shared__ - if(new_type->op() == TYPE_ARRAY && ((array_type*)new_type)->elem_type()->op() == TYPE_ARRAY){ - //protonu--changes suggested by Malik - //printf("Adding __shared__ annotation to : %s\n", new_sn->name()); - //new_type = ocg->ModifyType(new_type, "__shared__"); - //proc_syms->add_type(new_type); - } - var->set_type(new_type); - } - proc_syms->add_sym(new_sn); - r->newsyms.append(new_sn); - //printf("def new: %s\n", new_sn->name()); - } - } - - //printf("proc_syms var runthrough\n"); - //proc_syms->print(); - return code->clone_helper(r); -} - -bool LoopCuda::validIndexes(int stmt, const std::vector<std::string>& idxs){ - for(int i=0; i<idxs.size(); i++){ - bool found = false; - for(int j=0; j<idxNames[stmt].size(); j++){ - if(strcmp(idxNames[stmt][j].c_str(), idxs[i].c_str()) == 0){ - found=true; - } - } - if(!found){ - return false; - } - } - return true; -} - - -bool LoopCuda::cudaize_v2(std::string kernel_name, std::map<std::string, int> array_dims, - std::vector<std::string> blockIdxs, std::vector<std::string> threadIdxs) -{ - int stmt_num = 0; - if(cudaDebug){ - printf("cudaize_v2(%s, {", kernel_name.c_str()); - //for( - printf("}, blocks={"); printVs(blockIdxs); printf("}, thread={"); printVs(threadIdxs); printf("})\n"); - } - - this->array_dims = array_dims; - if(!validIndexes(stmt_num, blockIdxs)){ - throw std::runtime_error("One of the indexes in the block list was not " - "found in the current set of indexes."); - } - if(!validIndexes(stmt_num, threadIdxs)){ - throw std::runtime_error("One of the indexes in the thread list was not " - "found in the current set of indexes."); - } - if(blockIdxs.size() ==0) - throw std::runtime_error("Cudaize: Need at least one block dimention"); - int block_level=0; - //Now, we will determine the actual size (if possible, otherwise - //complain) for the block dimentions and thread dimentions based on our - //indexes and the relations for our stmt; - for(int i=0; i<blockIdxs.size(); i++){ - int level = findCurLevel(stmt_num, blockIdxs[i]); - int ub,lb; - extractCudaUB(stmt_num,level,ub,lb); - if(lb!= 0){ - //attempt to "normalize" the loop with an in-place tile and then re-check our bounds - if(cudaDebug) printf("Cudaize: doing tile at level %d to try and normalize lower bounds\n", level); - tile(stmt_num,level,1,level,CountedTile); - idxNames[stmt_num].insert(idxNames[stmt_num].begin()+(level),"");//TODO: possibly handle this for all sibling stmts - extractCudaUB(stmt_num,level,ub,lb); - } - if(lb != 0){ - char buf[1024]; - sprintf(buf, "Cudaize: Loop at level %d does not have 0 as it's lower bound", level); - throw std::runtime_error(buf); - } - if(ub < 0){ - char buf[1024]; - sprintf(buf, "Cudaize: Loop at level %d does not have a hard upper bound", level); - throw std::runtime_error(buf); - } - if(cudaDebug) printf("block idx %s level %d lb: %d ub %d\n", blockIdxs[i].c_str(), level, lb, ub); - if(i == 0){ - block_level = level; - cu_bx = ub+1; - idxNames[stmt_num][level-1] = "bx"; - } - else if(i == 1){ - cu_by = ub+1; - idxNames[stmt_num][level-1] = "by"; - } - } - if(!cu_by) - block_level=0; - int thread_level1 = 0; - int thread_level2 = 0; - for(int i=0; i<threadIdxs.size(); i++){ - int level = findCurLevel(stmt_num, threadIdxs[i]); - int ub,lb; - extractCudaUB(stmt_num,level,ub,lb); - if(lb!= 0){ - //attempt to "normalize" the loop with an in-place tile and then re-check our bounds - if(cudaDebug) printf("Cudaize: doing tile at level %d to try and normalize lower bounds\n", level); - tile(stmt_num,level,1,level,CountedTile); - idxNames[stmt_num].insert(idxNames[stmt_num].begin()+(level),""); - extractCudaUB(stmt_num,level,ub,lb); - } - if(lb != 0){ - char buf[1024]; - sprintf(buf, "Cudaize: Loop at level %d does not have 0 as it's lower bound", level); - throw std::runtime_error(buf); - } - if(ub < 0){ - char buf[1024]; - sprintf(buf, "Cudaize: Loop at level %d does not have a hard upper bound", level); - throw std::runtime_error(buf); - } - - if(cudaDebug) printf("thread idx %s level %d lb: %d ub %d\n", threadIdxs[i].c_str(), level, lb, ub); - if(i == 0){ - thread_level1 = level; - cu_tx = ub+1; - idxNames[stmt_num][level-1] = "tx"; - } - else if(i == 1){ - thread_level2 = level; - cu_ty = ub+1; - idxNames[stmt_num][level-1] = "ty"; - } - else if(i == 2){ - cu_tz = ub+1; - idxNames[stmt_num][level-1] = "tz"; - } - } - if(!cu_ty) - thread_level1 = 0; - if(!cu_tz) - thread_level2 = 0; - - //Make changes to nonsplitlevels - const int m = stmt.size(); - for (int i = 0; i < m; i++) { - if(block_level){ - //stmt[i].nonSplitLevels.append((block_level)*2); - stmt_nonSplitLevels[i].append((block_level)*2); - } - if(thread_level1){ - //stmt[i].nonSplitLevels.append((thread_level1)*2); - stmt_nonSplitLevels[i].append((thread_level1)*2); - } - if(thread_level2){ - //stmt[i].nonSplitLevels.append((thread_level1)*2); - stmt_nonSplitLevels[i].append((thread_level1)*2); - } - } - - if(cudaDebug) { - printf("Codegen: current names: "); - printVS(idxNames[stmt_num]); - } - //Set codegen flag - code_gen_flags |= GenCudaizeV2; - - //Save array dimention sizes - this->array_dims = array_dims; - cu_kernel_name = kernel_name.c_str(); - -} - -tree_node_list* LoopCuda::cudaize_codegen_v2() -{ - //printf("cudaize codegen V2\n"); - CG_suifBuilder *ocg = dynamic_cast<CG_suifBuilder*>(ir->builder()); - if(!ocg) return false; - - //protonu--adding an annote to track texture memory type - ANNOTE(k_cuda_texture_memory, "cuda texture memory", TRUE); - ANNOTE(k_cuda_constant_memory, "cuda constant memory", TRUE); - int tex_mem_on = 0; - int cons_mem_on = 0; - - - - CG_outputRepr* repr; - std::vector<VarDefs> arrayVars; - std::vector<VarDefs> localScopedVars; - - std::vector<IR_ArrayRef *> ro_refs; - std::vector<IR_ArrayRef *> wo_refs; - std::set<std::string> uniqueRefs; - std::set<std::string> uniqueWoRefs; - //protonu--let's try a much simpler approach of a map instead - //we also keep a map for constant memories - std::map<std::string , var_sym *>tex_ref_map; - std::map<std::string , var_sym *>cons_ref_map; - - for(int j=0; j<stmt.size(); j++) - { - std::vector<IR_ArrayRef *> refs = ir->FindArrayRef(stmt[j].code); - for (int i = 0; i < refs.size(); i++) - { - //printf("ref %s wo %d\n", static_cast<const char*>(refs[i]->name()), refs[i]->is_write()); - var_sym* var = symtab->lookup_var((char*)refs[i]->name().c_str(),false); - //If the array is not a parameter, then it's a local array and we - //want to recreate it as a stack variable in the kernel as opposed to - //passing it in. - if(!var->is_param()) - continue; - if (uniqueRefs.find(refs[i]->name()) == uniqueRefs.end()) - { - uniqueRefs.insert(refs[i]->name()); - if(refs[i]->is_write()){ - uniqueWoRefs.insert(refs[i]->name()); - wo_refs.push_back(refs[i]); - } - else - ro_refs.push_back(refs[i]); - } - if (refs[i]->is_write() && uniqueWoRefs.find(refs[i]->name()) == uniqueWoRefs.end()){ - uniqueWoRefs.insert(refs[i]->name()); - wo_refs.push_back(refs[i]); - //printf("adding %s to wo\n", static_cast<const char*>(refs[i]->name())); - } - } - } - - // printf("reading from array "); - // for(int i=0; i<ro_refs.size(); i++) - // printf("'%s' ", ro_refs[i]->name().c_str()); - // printf("and writting to array "); - // for(int i=0; i<wo_refs.size(); i++) - // printf("'%s' ", wo_refs[i]->name().c_str()); - // printf("\n"); - - const char* gridName = "dimGrid"; - const char* blockName = "dimBlock"; - - //TODO: Could allow for array_dims_vars to be a mapping from array - //references to to variable names that define their length. - var_sym* dim1 = 0; - var_sym* dim2 = 0; - - for(int i=0; i<wo_refs.size(); i++) - { - //TODO: Currently assume all arrays are floats of one or two dimentions - var_sym* outArray = 0; - std::string name = wo_refs[i]->name(); - outArray = symtab->lookup_var((char*)name.c_str(),false); - - VarDefs v; - v.size_2d = -1; - char buf[32]; - snprintf(buf, 32, "devO%dPtr", i+1); - v.name = buf; - if(outArray->type()->is_ptr()) - if(((ptr_type *)(outArray->type()))->ref_type()->is_array()) - v.type = ((array_type *)(((ptr_type *)(outArray->type()))->ref_type()))->elem_type(); - else - v.type = ((ptr_type *)(outArray->type()))->ref_type(); - else - v.type = type_f32; - v.tex_mapped = false; - v.cons_mapped = false; - v.original_name = wo_refs[i]->name(); - //Size of the array = dim1 * dim2 * num bytes of our array type - - //If our input array is 2D (non-linearized), we want the actual - //dimentions of the array - CG_outputRepr* size; - //Lookup in array_dims - std::map<std::string, int>::iterator it = array_dims.find(name.c_str()); - if(outArray->type()->is_ptr() && outArray->type()->ref_type(0)->is_array()) - { - array_type* t = (array_type*)outArray->type()->ref_type(0); - v.size_2d = t->upper_bound().constant()+1; - printf("Detected 2D array sized of %d for %s\n", v.size_2d, (char*)wo_refs[i]->name().c_str()); - size = ocg->CreateInt(v.size_2d * v.size_2d); - }else if(it != array_dims.end()){ - int ref_size = it->second; - v.var_ref_size = ref_size; - size = ocg->CreateInt(ref_size); - } - else{ - if(dim1){ - size = ocg->CreateTimes(new CG_suifRepr(operand(dim1)), - new CG_suifRepr(operand(dim2))); - }else{ - char buf[1024]; - sprintf(buf, "CudaizeCodeGen: Array reference %s does not have a " - "detectable size or specififed dimentions", name.c_str()); - throw std::runtime_error(buf); - } - } - v.size_expr = operand(static_cast<CG_suifRepr*>(ocg->CreateTimes( - size, - ocg->CreateInt(v.type->size()/8)))->GetExpression()); - v.in_data = 0; - v.out_data = outArray; - //Check for in ro_refs and remove it at this point - std::vector<IR_ArrayRef *>::iterator it_; - for(it_ = ro_refs.begin(); it_ != ro_refs.end(); it_++) - { - if((*it_)->name() == wo_refs[i]->name()){ - break; - } - } - if(it_ != ro_refs.end()) - { - v.in_data = outArray; - ro_refs.erase(it_); - } - - arrayVars.push_back(v); - - } - - //protonu-- assuming that all texture mapped memories were originally read only mems - //there should be safety checks for that, will implement those later - - int cs_ref_size = 0; - - for(int i=0; i<ro_refs.size(); i++) - { - var_sym* inArray = 0; - std::string name = ro_refs[i]->name(); - inArray = symtab->lookup_var((char*)name.c_str(),false); - VarDefs v; - v.size_2d = -1; - char buf[32]; - snprintf(buf, 32, "devI%dPtr", i+1); - v.name = buf; - if(inArray->type()->is_ptr()) - if(((ptr_type *)(inArray->type()))->ref_type()->is_array()) - v.type = ((array_type *)(((ptr_type *)(inArray->type()))->ref_type()))->elem_type(); - else - v.type = ((ptr_type *)(inArray->type()))->ref_type(); - else - v.type = type_f32; - v.tex_mapped = false; - v.cons_mapped = false; - v.original_name = ro_refs[i]->name(); - if ( texture != NULL) - v.tex_mapped = (texture->is_array_tex_mapped(name.c_str()))? true:false; //protonu-track tex mapped vars - if (v.tex_mapped){ - printf("this variable %s is mapped to texture memory", name.c_str()); - } - if ( constant_mem != NULL) - v.cons_mapped = (constant_mem->is_array_cons_mapped(name.c_str()))? true:false; //protonu-track tex mapped vars - if (v.cons_mapped){ - printf("this variable %s is mapped to constant memory", name.c_str()); - } - - //Size of the array = dim1 * dim2 * num bytes of our array type - - //If our input array is 2D (non-linearized), we want the actual - //dimentions of the array (as it might be less than cu_n - CG_outputRepr* size; - //Lookup in array_dims - std::map<std::string, int>::iterator it = array_dims.find(name.c_str()); - int ref_size = 0; - if(inArray->type()->is_ptr() && inArray->type()->ref_type(0)->is_array()) - { - array_type* t = (array_type*)inArray->type()->ref_type(0); - v.size_2d = t->upper_bound().constant()+1; - printf("Detected 2D array sized of %d for %s\n", v.size_2d, (char*)ro_refs[i]->name().c_str()); - size = ocg->CreateInt(v.size_2d * v.size_2d); - }else if(it != array_dims.end()){ - ref_size = it->second; - v.var_ref_size = ref_size; - size = ocg->CreateInt(ref_size); - }else{ - if(dim1){ - size = ocg->CreateTimes(new CG_suifRepr(operand(dim1)), - new CG_suifRepr(operand(dim2))); - }else{ - char buf[1024]; - sprintf(buf, "CudaizeCodeGen: Array reference %s does not have a " - "detectable size or specififed dimentions", name.c_str()); - throw std::runtime_error(buf); - } - } - - - - v.size_expr = operand(static_cast<CG_suifRepr*>(ocg->CreateTimes( - size, - ocg->CreateInt(v.type->size()/8)))->GetExpression()); - - v.in_data = inArray; - v.out_data = 0; - arrayVars.push_back(v); - } - - - if(arrayVars.size() < 2) - { - fprintf(stderr, "cudaize error: Did not find two arrays being accessed\n"); - return false; - } - - //protonu--debugging tool--the printf statement - //tex_mem_on signals use of tex mem - for(int i=0; i<arrayVars.size(); i++) - { - //printf("var name %s, tex_mem used %s\n", arrayVars[i].name.c_str(), (arrayVars[i].tex_mapped)?"true":"false"); - if (arrayVars[i].tex_mapped ) tex_mem_on ++; - if (arrayVars[i].cons_mapped ) cons_mem_on ++; - } - - //Add CUDA function extern prototypes and function types - func_type* unkown_func = new func_type(type_s32); //function on unkown args that returns a i32 - unkown_func = (func_type*)symtab->install_type(unkown_func); - func_type* void_func = new func_type(type_void); //function on unkown args that returns a void - void_func = (func_type*)globals->install_type(void_func); - func_type* float_func = new func_type(type_f32); //function on unkown args that returns a float - float_func = (func_type*)globals->install_type(float_func); - - type_node* result = ocg->ModifyType(type_void, "__global__"); - result = globals->install_type(result); - func_type* kernel_type = new func_type(result); //function returns a '__global__ void' - - int numArgs = arrayVars.size() + (dim1 ? 2 : 0) + localScopedVars.size(); - //protonu--need to account for texture memory here, reduce the #args - if( tex_mem_on ) numArgs -= tex_mem_on; - if( cons_mem_on ) numArgs -= cons_mem_on; - kernel_type->set_num_args(numArgs); - int argCount = 0; - for(int i=0; i<arrayVars.size(); i++) - { - type_node* fptr; - if(arrayVars[i].in_data) - fptr = arrayVars[i].in_data->type()->clone(); - else - fptr = arrayVars[i].out_data->type()->clone(); - //protonu--skip this for texture mems - if( arrayVars[i].tex_mapped != true && arrayVars[i].cons_mapped !=true ) - kernel_type->set_arg_type(argCount++, fptr); - } - if(dim1){ - kernel_type->set_arg_type(argCount++, type_s32); //width x height dimentions - kernel_type->set_arg_type(argCount++, type_s32); - } - kernel_type = (func_type*)globals->install_type(kernel_type); - - proc_sym* cudaMalloc = globals->new_proc(unkown_func, src_c, "cudaMalloc"); - proc_sym* cudaMemcpy = globals->new_proc(unkown_func, src_c, "cudaMemcpy"); - proc_sym* cudaFree = globals->new_proc(unkown_func, src_c, "cudaFree"); - proc_sym* cudaSync = globals->new_proc(void_func, src_c, "__syncthreads"); - proc_sym* cudaBind = globals->new_proc(unkown_func, src_c, "cudaBindTexture"); - proc_sym* cudaMemcpySym = globals->new_proc(unkown_func, src_c, "cudaMemcpyToSymbol"); - - - //protonu-removing Gabe's function, introducing mine, this is pretty cosmetic - //proc_sym* cudaFetch = globals->new_proc(float_func, src_c, "tex1Dfetch"); - proc_sym* tex1D = globals->new_proc(float_func, src_c, "tex1Dfetch"); - - var_sym *cudaMemcpyHostToDevice = new var_sym(type_s32, "cudaMemcpyHostToDevice"); - var_sym *cudaMemcpyDeviceToHost = new var_sym(type_s32, "cudaMemcpyDeviceToHost"); - cudaMemcpyDeviceToHost->set_param(); - cudaMemcpyHostToDevice->set_param(); - globals->add_sym(cudaMemcpyHostToDevice); - globals->add_sym(cudaMemcpyDeviceToHost); - - //protonu--adding the bool tex_mem to the structure struct_type - //to bypass the re-naming of struct texture, this is a hack fix - struct_type* texType = new struct_type(TYPE_GROUP, 0, "texture<float, 1, cudaReadModeElementType>", 0, true); - immed_list *iml_tex = new immed_list; - iml_tex->append(immed("texture memory")); - texType->append_annote(k_cuda_texture_memory, iml_tex); - //protonu--end my changes - texType = (struct_type*)globals->install_type(texType); - //protonu--should register the locals later on - //when we do the bind operation - //var_sym* texRef = new var_sym(texType, "texRef"); - //globals->add_sym(texRef); - - //Add our mallocs (and input array memcpys) - for(int i=0; i<arrayVars.size(); i++) - { - //protonu--check if the variable is not a tex-mapped variable. If it is tex mapped - // allow a malloc and memcpy operation, and a bind, but only if it is tex mapped, but dont call - // the kernel with it as an argument. - - //Make a pointer of type a[i].type - //type_node* fptr = new ptr_type(arrayVars[i].type->clone()); - //protonu--temporary change - type_node* fptr = new ptr_type(arrayVars[i].type); - fptr = symtab->install_type(fptr); - var_sym *dvs = new var_sym(fptr, const_cast<char*>( - arrayVars[i].name.c_str())); - dvs->set_addr_taken(); - symtab->add_sym(dvs); - - //cudaMalloc args - //protonu--no cudaMalloc required for constant memory - tree_node_list* tnl = new tree_node_list; - if(arrayVars[i].cons_mapped != true ) - { - in_cal *the_call = - new in_cal(type_s32, operand(), operand(new in_ldc(unkown_func->ptr_to(), operand(), immed(cudaMalloc))), 2); - the_call->set_argument(0, operand(new in_ldc(type_void->ptr_to()->ptr_to(), operand(), immed(dvs)))); - the_call->set_argument(1, arrayVars[i].size_expr); - - tnl->append(new tree_instr(the_call)); - setup_code = ocg->StmtListAppend(setup_code, - new CG_suifRepr(tnl)); - } - if(arrayVars[i].in_data) - { - //cudaMemcpy args - //protonu-- no cudaMemcpy required for constant memory - if ( arrayVars[i].cons_mapped != true ) - { - in_cal *the_call = - new in_cal(type_s32, operand(), operand(new in_ldc(unkown_func->ptr_to(), operand(), immed(cudaMemcpy))), 4); - the_call->set_argument(0, operand(dvs)); - the_call->set_argument(1, operand(arrayVars[i].in_data)); - the_call->set_argument(2, arrayVars[i].size_expr.clone()); - the_call->set_argument(3, operand(cudaMemcpyHostToDevice)); - - tnl = new tree_node_list; - tnl->append(new tree_instr(the_call)); - setup_code = ocg->StmtListAppend(setup_code, - new CG_suifRepr(tnl)); - } - - //protonu--check if the arrayvar is tex mapped - if(arrayVars[i].tex_mapped == true) - { - //Need a texture reference variable - char buf[32]; - snprintf(buf, 32, "tex%dRef", i+1); - arrayVars[i].secondName = buf; - - var_sym* texRef = new var_sym(texType, buf); - //printf("\n putting in %s\n", arrayVars[i].original_name.c_str()); - tex_ref_map[arrayVars[i].original_name] = texRef; - globals->add_sym(texRef); - //protonu--added the above two lines - - in_cal *the_call = - new in_cal(type_s32, operand(), operand(new in_ldc(unkown_func->ptr_to(), operand(), immed(cudaBind))), 4); - in_ldc *ins = new in_ldc(type_s32, operand(), immed(0)); - the_call->set_argument(0, operand(ins)); - the_call->set_argument(1, operand(texRef));//protonu--change to add the new sym - the_call->set_argument(2, operand(dvs)); - the_call->set_argument(3, arrayVars[i].size_expr.clone()); - - tnl = new tree_node_list; - tnl->append(new tree_instr(the_call)); - setup_code = ocg->StmtListAppend(setup_code, - new CG_suifRepr(tnl)); - } - - //protonu--if arrayvar is mapped to constant memory - if(arrayVars[i].cons_mapped == true) - { - char buf[32]; - snprintf(buf, 32, "cs%dRef", i+1); - //arrayVars[i].secondName = buf; - array_bound low (0); - array_bound high (arrayVars[i].var_ref_size -1); - array_type *arr = new array_type(arrayVars[i].type,low, high); - type_node* cons_arr = ocg->ModifyType(arr, "__device__ __constant__"); - cons_arr = globals->install_type(cons_arr); - var_sym* consRef = new var_sym(cons_arr, buf); - cons_ref_map[arrayVars[i].original_name] = consRef; - globals->add_sym(consRef); - - - - in_cal *the_call = - new in_cal(type_s32, operand(), operand(new in_ldc(unkown_func->ptr_to(), operand(), immed(cudaMemcpySym))), 3); - the_call->set_argument(0, operand(new in_ldc(type_void->ptr_to(), operand(), immed(consRef)))); - the_call->set_argument(1, operand(arrayVars[i].in_data)); - the_call->set_argument(2, arrayVars[i].size_expr.clone()); - - tnl = new tree_node_list; - tnl->append(new tree_instr(the_call)); - setup_code = ocg->StmtListAppend(setup_code, - new CG_suifRepr(tnl)); - - } - } - } - - //Build dimGrid dim3 variables based on loop dimentions and ti/tj - char blockD1[120]; - char blockD2[120]; - if(dim1){ - snprintf(blockD1, 120, "%s/%d", dim1->name(), cu_tx); - snprintf(blockD2, 120, "%s/%d", dim2->name(), cu_ty); - }else{ - snprintf(blockD1, 120, "%d", cu_bx); - snprintf(blockD2, 120, "%d", cu_by); - //snprintf(blockD1, 120, "%d/%d", cu_nx, cu_tx); - //snprintf(blockD2, 120, "%d/%d", cu_ny, cu_ty); - } - repr = ocg->CreateDim3(immed((char*)gridName), - immed(blockD1), - immed(blockD2)); - setup_code = ocg->StmtListAppend(setup_code, repr); - - repr = ocg->CreateDim3(immed((char*)blockName), immed(cu_tx),immed(cu_ty)); - - if(cu_tz > 1) - repr = ocg->CreateDim3(immed((char*)blockName), immed(cu_tx), immed(cu_ty), immed(cu_tz)); - else - repr = ocg->CreateDim3(immed((char*)blockName), immed(cu_tx), immed(cu_ty)); - setup_code = ocg->StmtListAppend(setup_code, repr); - - //call kernel function with name loop_name - //like: transpose_k<<<dimGrid,dimBlock>>>(devOPtr, devIPtr , width, height); - char dims[120]; - snprintf(dims,120,"<<<%s,%s>>>",gridName, blockName); - immed_list *iml = new immed_list; - iml->append(immed((char*)cu_kernel_name.c_str())); - iml->append(immed(dims)); - //printf("%s %s\n", static_cast<const char*>(cu_kernel_name), dims); - for(int i=0; i<arrayVars.size(); i++) - //Throw in a type cast if our kernel takes 2D array notation - //like (float(*) [1024]) - { - //protonu--throwing in another hack to stop the caller from passing tex mapped - //vars to the kernel. - if(arrayVars[i].tex_mapped == true || arrayVars[i].cons_mapped == true ) - continue; - if(arrayVars[i].size_2d >= 0) - { - snprintf(dims,120,"(float(*) [%d])%s", arrayVars[i].size_2d, - const_cast<char*>(arrayVars[i].name.c_str())); - //printf("%d %s\n", i, dims); - iml->append(immed(dims)); - }else{ - //printf("%d %s\n", i, static_cast<const char*>(arrayVars[i].name)); - iml->append(immed(const_cast<char*>( - arrayVars[i].name.c_str()))); - } - } - if(dim1){ - iml->append(immed(dim1)); - iml->append(immed(dim2)); - } - repr = ocg->CreateKernel(iml);//kernel call - setup_code = ocg->StmtListAppend(setup_code, repr); - - //cuda free variables - for(int i=0; i<arrayVars.size(); i++) - { - if(arrayVars[i].out_data) - { - //cudaMemcpy args - in_cal *the_call = - new in_cal(type_s32, operand(), operand(new in_ldc(unkown_func->ptr_to(), operand(), immed(cudaMemcpy))), 4); - the_call->set_argument(0, operand(arrayVars[i].out_data)); - the_call->set_argument(1, operand(symtab->lookup_var(const_cast<char*>( - arrayVars[i].name.c_str())))); - the_call->set_argument(2, arrayVars[i].size_expr.clone()); - the_call->set_argument(3, operand(cudaMemcpyDeviceToHost)); - - tree_node_list* tnl = new tree_node_list; - tnl->append(new tree_instr(the_call)); - teardown_code = ocg->StmtListAppend(teardown_code, - new CG_suifRepr(tnl)); - } - - in_cal *the_call = - new in_cal(type_s32, operand(), operand(new in_ldc(unkown_func->ptr_to(), operand(), immed(cudaFree))), 1); - the_call->set_argument(0, operand(symtab->lookup_var(const_cast<char*>( - arrayVars[i].name.c_str())))); - - tree_node_list* tnl = new tree_node_list; - tnl->append(new tree_instr(the_call)); - teardown_code = ocg->StmtListAppend(teardown_code, - new CG_suifRepr(tnl)); - } - - // --------------- - // BUILD THE KERNEL - // --------------- - - //Extract out kernel body - tree_node_list* code = getCode(); - //Get rid of wrapper if that original() added - if(code->head()->contents->kind() == TREE_IF) - { - tree_if* ifn = (tree_if*)code->head()->contents; - code = ifn->then_part(); - } - - //Create kernel function body - proc_sym *new_psym = globals->new_proc(kernel_type, src_c, (char*)cu_kernel_name.c_str()); - proc_symtab *new_proc_syms = new proc_symtab(new_psym->name()); - globals->add_child(new_proc_syms); - - //Add Params - std::map<std::string, var_sym*> loop_vars; - //In-Out arrays - type_node* fptr; - for(int i=0; i<arrayVars.size(); i++) - { - if(arrayVars[i].in_data) - //fptr = arrayVars[i].in_data->type()->clone(); - fptr = arrayVars[i].in_data->type(); - else - //fptr = arrayVars[i].out_data->type()->clone(); - fptr = arrayVars[i].out_data->type(); - fptr = new_proc_syms->install_type(fptr); - std::string name = arrayVars[i].in_data ? arrayVars[i].in_data->name() : arrayVars[i].out_data->name(); - var_sym* sym = new var_sym(fptr, arrayVars[i].in_data ? arrayVars[i].in_data->name() : arrayVars[i].out_data->name()); - //protonu--adding a check to ensure that texture memories are not passed in as arguments - if(arrayVars[i].tex_mapped != true && arrayVars[i].cons_mapped !=true ) - { - sym->set_param(); - new_proc_syms->params()->append(sym); - new_proc_syms->add_sym(sym);//protonu--added to suppress the addition of the redundant var in the kernel - } - if (arrayVars[i].cons_mapped == true) - { - sym->set_param(); - new_proc_syms->add_sym(sym); - } - //printf("inserting name: %s\n", static_cast<const char*>(name)); - loop_vars.insert(std::pair<std::string, var_sym*>(std::string(name), sym)); - } - - if(dim1) - { - //Array dimentions - var_sym* kdim1 = new var_sym(dim1->type(), dim1->name()); - kdim1->set_param(); - new_proc_syms->add_sym(kdim1); - loop_vars.insert(std::pair<std::string, var_sym*>(std::string(dim1->name()), kdim1)); - var_sym* kdim2 = new var_sym(dim2->type(), dim2->name()); - kdim2->set_param(); - new_proc_syms->add_sym(kdim2); - loop_vars.insert(std::pair<std::string, var_sym*>(std::string(dim2->name()), kdim2)); - new_proc_syms->params()->append(kdim1); - new_proc_syms->params()->append(kdim2); - } - //Put block and thread implicit variables into scope - std::vector<var_sym *> index_syms; - /* Currently we don't use the block dimentions - var_sym* blockDim_x = new var_sym(type_s32, "blockDim.x"); - blockDim_x->set_param(); - new_proc_syms->add_sym(blockDim_x); - var_sym* blockDim_y = new var_sym(type_s32, "blockDim.y"); - blockDim_y->set_param(); - new_proc_syms->add_sym(blockDim_y); - */ - if(cu_bx > 1){ - var_sym* blockIdx_x = new var_sym(type_s32, "blockIdx.x"); - blockIdx_x->set_param(); - new_proc_syms->add_sym(blockIdx_x); - index_syms.push_back(blockIdx_x); - } - if(cu_by > 1){ - var_sym* blockIdx_y = new var_sym(type_s32, "blockIdx.y"); - blockIdx_y->set_param(); - new_proc_syms->add_sym(blockIdx_y); - index_syms.push_back(blockIdx_y); - } - if(cu_tx > 1){ - var_sym* threadIdx_x = new var_sym(type_s32, "threadIdx.x"); - threadIdx_x->set_param(); - new_proc_syms->add_sym(threadIdx_x); - index_syms.push_back(threadIdx_x); - } - if(cu_ty > 1){ - var_sym* threadIdx_y = new var_sym(type_s32, "threadIdx.y"); - threadIdx_y->set_param(); - new_proc_syms->add_sym(threadIdx_y); - index_syms.push_back(threadIdx_y); - } - - if(cu_tz > 1){ - var_sym* threadIdx_z = new var_sym(type_s32, "threadIdx.z"); - threadIdx_z->set_param(); - new_proc_syms->add_sym(threadIdx_z); - index_syms.push_back(threadIdx_z); - } - - //Figure out which loop variables will be our thread and block dimention variables - std::vector<var_sym *> loop_syms; - //Get our indexes - std::vector<const char*> indexes;// = get_loop_indexes(code,cu_num_reduce); - int threadsPos=0; - if(cu_bx > 1) - indexes.push_back("bx"); - if(cu_by > 1) - indexes.push_back("by"); - if(cu_tx > 1){ - threadsPos = indexes.size(); - indexes.push_back("tx"); - } - if(cu_ty > 1) - indexes.push_back("ty"); - if(cu_tz > 1) - indexes.push_back("tz"); - for(int i=0; i<indexes.size(); i++) - { - //printf("indexes[%d] = %s\n", i, (char*)indexes[i]); - loop_syms.push_back(new var_sym(type_s32, (char*)indexes[i])); - new_proc_syms->add_sym(loop_syms[i]); - //loop_vars.insert(std::pair<std::string, var_sym*>(std::string(indexes[i]), loop_syms[i])); - } - - //Generate this code - //int bx = blockIdx.x - //int by = blockIdx.y - //int tx = threadIdx.x - //int ty = threadIdx.y - CG_outputRepr *body=NULL; - for(int i=0; i<indexes.size(); i++){ - CG_outputRepr *lhs = new CG_suifRepr(operand(loop_syms[i])); - //body = ocg->StmtListAppend(body, ocg->CreateStmtList( - // ocg->CreateAssignment(0, lhs, new CG_suifRepr(operand(index_syms[i]))))); - body = ocg->StmtListAppend(body, ocg->StmtListAppend( - ocg->CreateAssignment(0, lhs, new CG_suifRepr(operand(index_syms[i]))), NULL)); - } - - //Get our inital code prepped for loop reduction. First we need to swap - //out internal SUIF variable references to point to the new local - //function symbol table. - std::map<std::string, var_sym*> loop_idxs; //map from idx names to their new syms - std::vector< std::pair<var_sym*, var_sym*> > dim_vars; //pair is of <old,new> var_sym (for 2D array size initializations) - replacements r; - tree_node_list* swapped = swapVarReferences(code, &r, ocg, loop_vars, new_proc_syms, dim_vars); - //printf("\n code before recursiveFindReplacePreferedIdxs :\n"); - //swapped->print(); - swapped = recursiveFindReplacePreferedIdxs(swapped, new_proc_syms, cudaSync, void_func, loop_idxs);//in-place swapping - //printf("\n code after recursiveFindReplacePreferedIdxs :\n"); - //swapped->print(); - - for(int i=0; i<indexes.size(); i++){ - std::vector<tree_for*> tfs = findCommentedFors(indexes[i], swapped); - for(int k=0; k<tfs.size(); k++){ - //printf("replacing %p tfs for index %s\n", tfs[k], indexes[i]); - tree_node_list* newBlock = forReduce(tfs[k], loop_idxs[indexes[i]], new_proc_syms); - //newBlock->print(); - swap_node_for_node_list(tfs[k], newBlock); - //printf("AFTER SWAP\n"); newBlock->print(); - } - } - //printf("AFTER REDUCE\n"); swapped->print(); - - if(static_cast<const IR_cudasuifCode *>(ir)->init_code()){ - tree_node_list* orig_init_code = static_cast<CG_suifRepr *>(static_cast<const IR_cudasuifCode *>(ir)->init_code())->GetCode(); - for(int i=0; i<dim_vars.size(); i++){ - //We have a map of var_sym from the original function body and we know - //that these var_syms have initialization statements which define the - //array size. We need to mimic these initialization statements. - - //First find the assignment and pull out the constant initialization - //value - int value = -1; - tree_node_list_iter tnli(orig_init_code); - while (!tnli.is_empty()) { - tree_node *node = tnli.step(); - if(node->kind() == TREE_INSTR && ((tree_instr*)node)->instr()->format() == inf_rrr) - { - in_rrr* inst = (in_rrr*)((tree_instr*)node)->instr(); - //expect the structure: cpy( _ = min(grab_me, _)) - if(inst->opcode() == io_cpy && inst->dst_op().is_symbol()){ - //printf("looking at instruction: "); - //inst->print(); - var_sym* dest = inst->dst_op().symbol(); - if(dest == dim_vars[i].first) - { - if(inst->src1_op().is_instr() && inst->src1_op().instr()->format() == inf_ldc){ - value = ((in_ldc*)inst->src1_op().instr())->value().integer(); - } - } - } - } - } - if(value < 0){ - fprintf(stderr, "ERROR: Could not find initializing statement for variable used in upper_bound of array type"); - } - CG_outputRepr *lhs = new CG_suifRepr(operand(dim_vars[i].second)); - //body = ocg->StmtListAppend(body, ocg->CreateStmtList(ocg->CreateAssignment(0, lhs, ocg->CreateInt(value)))); - body = ocg->StmtListAppend(body, ocg->StmtListAppend(ocg->CreateAssignment(0, lhs, ocg->CreateInt(value)), NULL)); - } - } - - - body = ocg->StmtListAppend(body, new CG_suifRepr(swapped)); - - //protonu--lets try creating our function definiton here - var_sym *tsym = NULL; - - - std::vector<IR_ArrayRef *> refs = ir->FindArrayRef(body); - for(int i=0; i<refs.size(); i++) - { - //check if the array is tex mapped - if(texture != NULL && texture->is_array_tex_mapped(refs[i]->name().c_str())) - { - //protonu--our new tex lookup function - in_cal *tex_lookup = - new in_cal(type_f32, operand(), operand(new in_ldc(float_func->ptr_to(), operand(), immed(tex1D))), 2); - - //printf("name of the array to be mapped is %s\n", refs[i]->name().c_str()); - tsym = tex_ref_map[refs[i]->name()]; - tex_lookup->set_argument(0, operand(tsym)); - - - int array_dims = ((IR_suifArrayRef *)refs[i])->ia_->dims(); - - if (array_dims == 1){ - tex_lookup->set_argument(1, ((IR_suifArrayRef *)refs[i])->ia_->index(0).clone()); - }else if (array_dims > 2) { - printf(" \n we don't handle more than 2D arrays mapped to textures yet\n"); - }else if (array_dims == 2) { - - IR_ArraySymbol *sym = refs[i]->symbol(); - CG_outputRepr *sz = sym->size(1); - delete sym; // free the wrapper object only - // find the builder ocg - CG_outputRepr *expr = ocg->CreateTimes(sz->clone(),refs[i]->index(0)); - delete sz; // free the wrapper object only - expr = ocg->CreatePlus(expr, refs[i]->index(1)); - // expr holds the 1D access expression and take it out - tex_lookup->set_argument(1, ((CG_suifRepr *)expr)->GetExpression()); - } - - //using chun's function to replace the array look up with the function call - ((IR_suifCode *)ir)->ReplaceExpression(refs[i] , new CG_suifRepr(operand(tex_lookup))); - } - - } - - - tsym = NULL; - //protonu--now let's try what we did above for constant memory - for(int i=0; i<refs.size(); i++) - { - //check if the array is tex mapped - if(constant_mem != NULL && constant_mem->is_array_cons_mapped(refs[i]->name().c_str())) - { - - //printf("name of the array to be cons mapped is %s\n", refs[i]->name().c_str()); - tsym = cons_ref_map[refs[i]->name()]; - //we should create a IR_SuifArray here - IR_ArraySymbol *ar_sym = new IR_suifArraySymbol(ir,tsym); - std::vector<CG_outputRepr *> ar_index; - ar_index.push_back(((IR_suifArrayRef *)refs[i])->index(0)); - IR_ArrayRef *ar_ref = ((IR_suifCode *)ir)->CreateArrayRef(ar_sym, ar_index); - //using chun's function to replace the array look up with the function call - ((IR_suifCode *)ir)->ReplaceExpression(refs[i] , new CG_suifRepr(operand(((IR_suifArrayRef *)ar_ref)->ia_))); - - } - } - - - tree_proc *new_body = new tree_proc(static_cast<CG_suifRepr*>(body)->GetCode(), new_proc_syms); - //globals->add_child(new_proc_syms); - new_psym->set_block(new_body); - new_procs.push_back(new_psym); - - return swapped; -} - -//Order taking out dummy variables -std::vector<std::string> cleanOrder(std::vector<std::string> idxNames){ - std::vector<std::string> results; - for(int j=0; j<idxNames.size(); j++){ - if(idxNames[j].length() != 0) - results.push_back(idxNames[j]); - } - return results; -} - -//First non-dummy level in ascending order -int LoopCuda::nonDummyLevel(int stmt, int level){ - //level comes in 1-basd and should leave 1-based - for(int j=level-1; j<idxNames[stmt].size(); j++){ - if(idxNames[stmt][j].length() != 0){ - //printf("found non dummy level of %d with idx: %s when searching for %d\n", j+1, (const char*) idxNames[stmt][j], level); - return j+1; - } - } - char buf[128]; sprintf(buf, "%d", level); - throw std::runtime_error(std::string("Unable to find a non-dummy level starting from ") + std::string(buf)); -} - -int LoopCuda::findCurLevel(int stmt, std::string idx){ - for(int j=0; j<idxNames[stmt].size(); j++){ - if(strcmp(idxNames[stmt][j].c_str(),idx.c_str()) == 0) - return j+1; - } - throw std::runtime_error(std::string("Unable to find index ") + idx + std::string(" in current list of indexes")); -} - -void LoopCuda::permute_cuda(int stmt, const std::vector<std::string>& curOrder) -{ - //printf("curOrder: "); - //printVs(curOrder); - //printf("idxNames: "); - //printVS(idxNames[stmt]); - std::vector<std::string> cIdxNames = cleanOrder(idxNames[stmt]); - bool same=true; - std::vector<int> pi; - for(int i=0; i<curOrder.size(); i++){ - bool found = false; - for(int j=0; j<cIdxNames.size(); j++){ - if(strcmp(cIdxNames[j].c_str(), curOrder[i].c_str()) == 0){ - pi.push_back(j+1); - found=true; - if(j!=i) - same=false; - } - } - if(!found){ - throw std::runtime_error("One of the indexes in the permute order where not " - "found in the current set of indexes."); - } - } - for(int i=curOrder.size(); i<cIdxNames.size(); i++){ - pi.push_back(i); - } - if(same) - return; - permute(stmt, pi); - //Set old indexe names as new - for(int i=0; i<curOrder.size(); i++){ - idxNames[stmt][i] = curOrder[i].c_str(); //what about sibling stmts? - } -} - - -bool LoopCuda::permute(int stmt_num, const std::vector<int> &pi) -{ -// check for sanity of parameters - if (stmt_num >= stmt.size() || stmt_num < 0) - throw std::invalid_argument("invalid statement " + to_string(stmt_num)); - const int n = stmt[stmt_num].xform.n_out(); - if (pi.size() > (n-1)/2) - throw std::invalid_argument("iteration space dimensionality does not match permute dimensionality"); - int first_level = 0; - int last_level = 0; - for (int i = 0; i < pi.size(); i++) { - if (pi[i] > (n-1)/2 || pi[i] <= 0) - throw std::invalid_argument("invalid loop level " + to_string(pi[i]) + " in permuation"); - - if (pi[i] != i+1) { - if (first_level == 0) - first_level = i+1; - last_level = i+1; - } - } - if (first_level == 0) - return true; - - std::vector<int> lex = getLexicalOrder(stmt_num); - std::set<int> active = getStatements(lex, 2*first_level-2); - Loop::permute(active, pi); -} - - -void LoopCuda::tile_cuda(int stmt, int level, int outer_level) -{ - tile_cuda(stmt,level,1,outer_level,"","",CountedTile); -} -void LoopCuda::tile_cuda(int level, int tile_size, int outer_level, std::string idxName, - std::string ctrlName, TilingMethodType method){ - tile_cuda(0, level, tile_size, outer_level, idxName, ctrlName, method); -} - -void LoopCuda::tile_cuda(int stmt, int level, int tile_size, int outer_level, std::string idxName, - std::string ctrlName, TilingMethodType method){ - //Do regular tile but then update the index and control loop variable - //names as well as the idxName to reflect the current state of things. - //printf("tile(%d,%d,%d,%d)\n", stmt, level, tile_size, outer_level); - //printf("idxNames before: "); - //printVS(idxNames[stmt]); - - tile(stmt, level, tile_size, outer_level, method); - - if(idxName.size()) - idxNames[stmt][level-1] = idxName.c_str(); - if(tile_size == 1){ - //potentially rearrange loops - if(outer_level < level){ - std::string tmp = idxNames[stmt][level-1]; - for(int i=level-1; i>outer_level-1; i--){ - if(i-1 >= 0) - idxNames[stmt][i] = idxNames[stmt][i-1]; - } - idxNames[stmt][outer_level-1] = tmp; - } - //TODO: even with a tile size of one, you need a insert (of a dummy loop) - idxNames[stmt].insert(idxNames[stmt].begin()+(level),""); - }else{ - if(!ctrlName.size()) - throw std::runtime_error("No ctrl loop name for tile"); - //insert - idxNames[stmt].insert(idxNames[stmt].begin()+(outer_level-1),ctrlName.c_str()); - } - - //printf("idxNames after: "); - //printVS(idxNames[stmt]); -} - - -bool LoopCuda::datacopy_privatized_cuda(int stmt_num, int level, const std::string &array_name, const std::vector<int> &privatized_levels, bool allow_extra_read , int fastest_changing_dimension , int padding_stride , int padding_alignment , bool cuda_shared) -{ - int old_stmts =stmt.size(); - //datacopy_privatized(stmt_num, level, array_name, privatized_levels, allow_extra_read, fastest_changing_dimension, padding_stride, padding_alignment, cuda_shared); - if(cuda_shared) - datacopy_privatized(stmt_num, level, array_name, privatized_levels, allow_extra_read, fastest_changing_dimension, padding_stride, padding_alignment, 1); - else - datacopy_privatized(stmt_num, level, array_name, privatized_levels, allow_extra_read, fastest_changing_dimension, padding_stride, padding_alignment, 0); - - - //Adjust idxNames to reflect updated state - std::vector<std::string> cIdxNames = cleanOrder(idxNames[stmt_num]); - int new_stmts = stmt.size(); - for(int i=old_stmts; i<new_stmts; i++){ - //printf("fixing up statement %d\n", i); - std::vector<std::string> idxs; - - - //protonu-making sure the vector of nonSplitLevels grows along with - //the statement structure - stmt_nonSplitLevels.push_back(omega::Tuple<int>()); - - //Indexes up to level will be the same - for(int j=0; j<level-1; j++) - idxs.push_back(cIdxNames[j]); - - //Expect privatized_levels to match - for(int j=0; j<privatized_levels.size(); j++) - idxs.push_back(cIdxNames[privatized_levels[j]-1]);//level is one-based - - //all further levels should match order they are in originally - if(privatized_levels.size()){ - int last_privatized = privatized_levels.back(); - int top_level = last_privatized + (stmt[i].IS.n_set()-idxs.size()); - //printf("last privatized_levels: %d top_level: %d\n", last_privatized, top_level); - for(int j=last_privatized; j<top_level; j++){ - idxs.push_back(cIdxNames[j]); - //printf("pushing back: %s\n", (const char*)cIdxNames[j]); - } - } - idxNames.push_back(idxs); - } -} - -bool LoopCuda::datacopy_cuda(int stmt_num, int level, const std::string &array_name, std::vector<std::string> new_idxs, bool allow_extra_read, int fastest_changing_dimension, int padding_stride, int padding_alignment, bool cuda_shared) -{ - - int old_stmts =stmt.size(); - //datacopy(stmt_num,level,array_name,allow_extra_read,fastest_changing_dimension,padding_stride,padding_alignment,cuda_shared); - if(cuda_shared) - datacopy(stmt_num,level,array_name,allow_extra_read,fastest_changing_dimension,padding_stride,padding_alignment, 1); - else - datacopy(stmt_num,level,array_name,allow_extra_read,fastest_changing_dimension,padding_stride,padding_alignment, 0); - //Adjust idxNames to reflect updated state - std::vector<std::string> cIdxNames = cleanOrder(idxNames[stmt_num]); - int new_stmts = stmt.size(); - for(int i=old_stmts; i<new_stmts; i++){ - //printf("fixing up statement %d\n", i); - std::vector<std::string> idxs; - - //protonu-making sure the vector of nonSplitLevels grows along with - //the statement structure - stmt_nonSplitLevels.push_back(omega::Tuple<int>()); - - //protonu--lets dump out the code from each statement here - //printf("\n dumping statement :%d", i); - //stmt[i].code->Dump(); - - //Indexes up to level will be the same - for(int j=0; j<level-1; j++) - idxs.push_back(cIdxNames[j]); - - //all further levels should get names from new_idxs - int top_level = stmt[i].IS.n_set(); - //printf("top_level: %d level: %d\n", top_level, level); - if(new_idxs.size() < top_level-level+1) - throw std::runtime_error("Need more new index names for new datacopy loop levels"); - - for(int j=level-1; j<top_level; j++){ - idxs.push_back(new_idxs[j-level+1].c_str()); - //printf("pushing back: %s\n", new_idxs[j-level+1].c_str()); - } - idxNames.push_back(idxs); - } -} - -bool LoopCuda::unroll_cuda(int stmt_num, int level, int unroll_amount) -{ - int old_stmts =stmt.size(); - //bool b= unroll(stmt_num, , unroll_amount); - - - int dim = 2*level-1; - std::vector<int> lex = getLexicalOrder(stmt_num); - std::set<int> same_loop = getStatements(lex, dim-1); - - level = nonDummyLevel(stmt_num,level); - //printf("unrolling %d at level %d\n", stmt_num,level); - - //protonu--using the new version of unroll, which returns - //a set of ints instead of a bool. To keep Gabe's logic - //I'll check the size of the set, if it's 0 return true - //bool b= unroll(stmt_num, level, unroll_amount); - std::set<int> b_set= unroll(stmt_num, level, unroll_amount); - bool b = false; - if (b_set.size() == 0) b = true; - //end--protonu - - //Adjust idxNames to reflect updated state - std::vector<std::string> cIdxNames = cleanOrder(idxNames[stmt_num]); - std::vector<std::string> origSource = idxNames[stmt_num];; - //Drop index names at level - if(unroll_amount == 0){ - //For all statements that were in this unroll together, drop index name for unrolled level - idxNames[stmt_num][level-1] = ""; - for (std::set<int>::iterator i = same_loop.begin(); i != same_loop.end(); i++) { - //printf("in same loop as %d is %d\n", stmt_num, (*i)); - //idxNames[(*i)][level-1] = ""; - idxNames[(*i)] = idxNames[stmt_num]; - } - } - - lex = getLexicalOrder(stmt_num); - same_loop = getStatements(lex, dim-1); - - bool same_as_source = false; - int new_stmts = stmt.size(); - for(int i=old_stmts; i<new_stmts; i++){ - //Check whether we had a sync for the statement we are unrolling, if - //so, propogate that to newly created statements so that if they are - //in a different loop structure, they will also get a syncthreads - int size = syncs.size(); - for(int j=0; j<size; j++){ - if(syncs[j].first == stmt_num) - syncs.push_back(make_pair(i,syncs[j].second)); - } - - //protonu-making sure the vector of nonSplitLevels grows along with - //the statement structure - stmt_nonSplitLevels.push_back(omega::Tuple<int>()); - - - //We expect that new statements have a constant for the variable in - //stmt[i].IS at level (as seen with print_with_subs), otherwise there - //will be a for loop at level and idxNames should match stmt's - //idxNames pre-unrolled - Relation IS = stmt[i].IS; - //Ok, if you know how the hell to get anything out of a Relation, you - //should probably be able to do this more elegantly. But for now, I'm - //hacking it. - std::string s = IS.print_with_subs_to_string(); - //s looks looks like - //{[_t49,8,_t51,_t52,128]: 0 <= _t52 <= 3 && 0 <= _t51 <= 15 && 0 <= _t49 && 64_t49+16_t52+_t51 <= 128} - //where level == 5, you see a integer in the input set - - //If that's not an integer and this is the first new statement, then - //we think codegen will have a loop at that level. It's not perfect, - //not sure if it can be determined without round-tripping to codegen. - int sIdx = 0; - int eIdx = 0; - for(int j=0; j<level-1; j++){ - sIdx = s.find(",",sIdx+1); - if(sIdx < 0) break; - } - if(sIdx > 0){ - eIdx = s.find("]"); - int tmp = s.find(",",sIdx+1); - if(tmp > 0 && tmp < eIdx) - eIdx = tmp; //", before ]" - if(eIdx > 0){ - sIdx++; - std::string var = s.substr(sIdx,eIdx-sIdx); - //printf("%s\n", s.c_str()); - //printf("set var for stmt %d at level %d is %s\n", i, level, var.c_str()); - if(atoi(var.c_str()) == 0 && i ==old_stmts){ - //TODO:Maybe do see if this new statement would be in the same - //group as the original and if it would, don't say - //same_as_source - if(same_loop.find(i) == same_loop.end()){ - printf("stmt %d level %d, newly created unroll statement should have same level indexes as source\n", i, level); - same_as_source = true; - } - } - } - } - - - //printf("fixing up statement %d n_set %d with %d levels\n", i, stmt[i].IS.n_set(), level-1); - if(same_as_source) - idxNames.push_back(origSource); - else - idxNames.push_back(idxNames[stmt_num]); - } - - return b; -} - -void LoopCuda::copy_to_texture(const char *array_name) -{ - //protonu--placeholder for now - //set the bool for using cuda memory as true - //in a vector of strings, put the names of arrays to tex mapped - if ( !texture ) - texture = new texture_memory_mapping(true, array_name); - else - texture->add(array_name); - - -} - - -void LoopCuda::copy_to_constant(const char *array_name) -{ - //protonu--placeholder for now - //set the bool for using cuda memory as true - //in a vector of strings, put the names of arrays to tex mapped - if ( !constant_mem ) - constant_mem = new constant_memory_mapping(true, array_name); - else - constant_mem->add(array_name); -} - -//protonu--moving this from Loop -tree_node_list* LoopCuda::codegen() -{ - if(code_gen_flags & GenCudaizeV2) - return cudaize_codegen_v2(); - //Do other flagged codegen methods, return plain vanilla generated code - return getCode(); -} - -//These three are in Omega code_gen.cc and are used as a massive hack to -//get out some info from MMGenerateCode. Yea for nasty side-effects. -namespace omega{ - extern int checkLoopLevel; - extern int stmtForLoopCheck; - extern int upperBoundForLevel; - extern int lowerBoundForLevel; -} - - -void LoopCuda::extractCudaUB(int stmt_num, int level, int &outUpperBound, int &outLowerBound){ - // check for sanity of parameters - const int m = stmt.size(); - if (stmt_num >= m || stmt_num < 0) - throw std::invalid_argument("invalid statement " + to_string(stmt_num)); - const int n = stmt[stmt_num].xform.n_out(); - if (level > (n-1)/2 || level <= 0) - throw std::invalid_argument("invalid loop level " + to_string(level)); - - int dim = 2*level-1; - - std::vector<int> lex = getLexicalOrder(stmt_num); - std::set<int> same_loop = getStatements(lex, dim-1); - - // extract the intersection of the iteration space to be considered - Relation hull; - { - hull = Relation::True(n); - for (std::set<int>::iterator i = same_loop.begin(); i != same_loop.end(); i++) { - hull = Intersection(hull, project_onto_levels(getNewIS(*i), dim+1, true)); - hull.simplify(2, 4); - } - - for (int i = 2; i <= dim+1; i+=2) { - //std::string name = std::string("_t") + to_string(t_counter++); - std::string name = std::string("_t") + to_string(tmp_loop_var_name_counter++); - hull.name_set_var(i, name); - } - hull.setup_names(); - } - - // extract the exact loop bound of the dimension to be unrolled - if (is_single_iteration(hull, dim)){ - throw std::runtime_error("No loop availabe at level to extract upper bound."); - } - Relation bound = get_loop_bound(hull, dim); - if (!bound.has_single_conjunct() || !bound.is_satisfiable() || bound.is_tautology()) - throw loop_error("loop error: unable to extract loop bound for cudaize"); - - // extract the loop stride - EQ_Handle stride_eq; - int stride = 1; - { - bool simple_stride = true; - int strides = countStrides(bound.query_DNF()->single_conjunct(), bound.set_var(dim+1), stride_eq, simple_stride); - if (strides > 1) - throw loop_error("loop error: too many strides"); - else if (strides == 1) { - int sign = stride_eq.get_coef(bound.set_var(dim+1)); -// assert(sign == 1 || sign == -1); - Constr_Vars_Iter it(stride_eq, true); - stride = abs((*it).coef/sign); - } - } - if(stride != 1){ - char buf[1024]; - sprintf(buf, "Cudaize: Loop at level %d has non-one stride of %d", level, stride); - throw std::runtime_error(buf); - } - - //Use code generation system to build tell us our bound information. We - //need a hard upper bound a 0 lower bound. - - checkLoopLevel = level*2; - stmtForLoopCheck = stmt_num; - upperBoundForLevel = -1; - lowerBoundForLevel = -1; - printCode(1,false); - checkLoopLevel = 0; - - outUpperBound = upperBoundForLevel; - outLowerBound = lowerBoundForLevel; - return; -} - - -void LoopCuda::printCode(int effort, bool actuallyPrint) const { - const int m = stmt.size(); - if (m == 0) - return; - const int n = stmt[0].xform.n_out(); - - - - Tuple<Relation> IS(m); - Tuple<Relation> xform(m); - Tuple<IntTuple > nonSplitLevels(m); - for (int i = 0; i < m; i++) { - IS[i+1] = stmt[i].IS; - xform[i+1] = stmt[i].xform; - nonSplitLevels[i+1] = stmt_nonSplitLevels[i]; - //nonSplitLevels[i+1] = stmt[i].nonSplitLevels; - } - - Tuple< Tuple<std::string> > idxTupleNames; - if(useIdxNames){ - for(int i=0; i<idxNames.size(); i++){ - Tuple<std::string> idxs; - for(int j=0; j<idxNames[i].size(); j++) - idxs.append(idxNames[i][j]); - idxTupleNames.append( idxs ); - } - } - - Relation known = Extend_Set(copy(this->known), n - this->known.n_set()); - CG_stringBuilder *ocg = new CG_stringBuilder(); - Tuple<CG_outputRepr *> nameInfo; - for (int i = 1; i <= m; i++) - nameInfo.append(new CG_stringRepr("s" + to_string(i))); - CG_outputRepr* repr = MMGenerateCode(ocg, xform, IS, nameInfo, known, nonSplitLevels, syncs, idxTupleNames, effort); - if(actuallyPrint) - std::cout << GetString(repr); -/* - for (int i = 1; i <= m; i++) - delete nameInfo[i]; -*/ - - delete ocg; -} - - - -void LoopCuda::printRuntimeInfo() const { - for(int i=0; i<stmt.size(); i++){ - Relation IS = stmt[i].IS; - Relation xform = stmt[i].xform; - printf("stmt[%d]\n", i); - printf("IS\n"); - IS.print_with_subs(); - - printf("xform[%d]\n", i); - xform.print_with_subs(); - - //printf("code\n"); - //static_cast<CG_suifRepr *>(stmt[i].code)->GetCode()->print_expr(); - } -} - -void LoopCuda::printIndexes() const { - for(int i=0; i<stmt.size(); i++){ - printf("stmt %d nset %d ", i, stmt[i].IS.n_set()); - - for(int j=0; j<idxNames[i].size(); j++){ - if(j>0) - printf(","); - printf("%s", idxNames[i][j].c_str()); - } - printf("\n"); - } -} - -tree_node_list* LoopCuda::getCode(int effort) const { - const int m = stmt.size(); - if (m == 0) - return new tree_node_list; - const int n = stmt[0].xform.n_out(); - - - - Tuple<CG_outputRepr *> ni(m); - Tuple<Relation> IS(m); - Tuple<Relation> xform(m); - Tuple< IntTuple > nonSplitLevels(m); - for (int i = 0; i < m; i++) { - ni[i+1] = stmt[i].code; - IS[i+1] = stmt[i].IS; - xform[i+1] = stmt[i].xform; - nonSplitLevels[i+1] = stmt_nonSplitLevels[i]; - //nonSplitLevels[i+1] = stmt[i].nonSplitLevels; - } - - - Relation known = Extend_Set(copy(this->known), n - this->known.n_set()); -#ifdef DEBUG -// std::cout << GetString(MMGenerateCode(new CG_stringBuilder(), xform, IS, known, effort)); -#endif - Tuple< Tuple<std::string> > idxTupleNames; - if(useIdxNames){ - for(int i=0; i<idxNames.size(); i++){ - Tuple<std::string> idxs; - for(int j=0; j<idxNames[i].size(); j++) - idxs.append(idxNames[i][j]); - idxTupleNames.append( idxs ); - } - } - - CG_outputBuilder *ocg = ir->builder(); - CG_outputRepr *repr = MMGenerateCode(ocg, xform, IS, ni, known, nonSplitLevels, syncs, idxTupleNames, effort); - - //CG_outputRepr *overflow_initialization = ocg->CreateStmtList(); - //protonu--using the new function CG_suifBuilder::StmtListAppend - CG_outputRepr *overflow_initialization = ocg->StmtListAppend(NULL, NULL); - for (std::map<int, std::vector<Free_Var_Decl *> >::const_iterator i = overflow.begin(); i != overflow.end(); i++) - for (std::vector<Free_Var_Decl *>::const_iterator j = i->second.begin(); j != i->second.end(); j++) - //overflow_initialization = ocg->StmtListAppend(overflow_initialization, ocg->CreateStmtList(ocg->CreateAssignment(0, ocg->CreateIdent((*j)->base_name()), ocg->CreateInt(0)))); - overflow_initialization = ocg->StmtListAppend(overflow_initialization, ocg->StmtListAppend(ocg->CreateAssignment(0, ocg->CreateIdent((*j)->base_name()), ocg->CreateInt(0)), NULL)); - - repr = ocg->StmtListAppend(overflow_initialization, repr); - tree_node_list *tnl = static_cast<CG_suifRepr *>(repr)->GetCode(); - - delete repr; - /* - for (int i = 1; i <= m; i++) - delete ni[i]; - */ - - return tnl; -} - - -//protonu--adding constructors for the new derived class -LoopCuda::LoopCuda():Loop(), code_gen_flags(GenInit){} - -LoopCuda::LoopCuda(IR_Control *irc, int loop_num) - :Loop(irc) -{ - setup_code = NULL; - teardown_code = NULL; - code_gen_flags = 0; - cu_bx = cu_by = cu_tx = cu_ty = cu_tz = 1; - cu_num_reduce = 0; - cu_mode = GlobalMem; - texture = NULL; - constant_mem = NULL; - - int m=stmt.size(); - //printf("\n the size of stmt(initially) is: %d\n", stmt.size()); - for(int i=0; i<m; i++) - stmt_nonSplitLevels.push_back(omega::Tuple<int>()); - - - //protonu--setting up - //proc_symtab *symtab - //global_symtab *globals - - globals = ((IR_cudasuifCode *)ir)->gsym_ ; - std::vector<tree_for *> tf = ((IR_cudasuifCode *)ir)->get_loops(); - - symtab = tf[loop_num]->proc()->block()->proc_syms(); - - std::vector<tree_for *> deepest = find_deepest_loops(tf[loop_num]); - - for (int i = 0; i < deepest.size(); i++){ - index.push_back(deepest[i]->index()->name()); //reflects original code index names - } - - for(int i=0; i< stmt.size(); i++) - idxNames.push_back(index); //refects prefered index names (used as handles in cudaize v2) - useIdxNames=false; - -} - |