summaryrefslogtreecommitdiff
path: root/src/cuda-sim/cuda-sim.cc
diff options
context:
space:
mode:
Diffstat (limited to 'src/cuda-sim/cuda-sim.cc')
-rw-r--r--src/cuda-sim/cuda-sim.cc596
1 files changed, 521 insertions, 75 deletions
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 9246613..4c0fc58 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -33,7 +33,7 @@
#include "ptx.tab.h"
#include "ptx_sim.h"
#include <stdio.h>
-
+#include <sstream>
#include "opcodes.h"
#include "../statwrapper.h"
#include <set>
@@ -58,12 +58,14 @@ int g_debug_execution = 0;
int g_debug_thread_uid = 0;
addr_t g_debug_pc = 0xBEEF1518;
// Output debug information to file options
+int cp_count;
+int cp_cta_resume;
unsigned g_ptx_sim_num_insn = 0;
unsigned gpgpu_param_num_shaders = 0;
-char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu;
-char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp,*opcode_initiation_sfu;
+char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor;
+char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp,*opcode_initiation_sfu,*opcode_initiation_tensor;
char *cdp_latency_str;
unsigned cdp_latency[5];
@@ -84,6 +86,10 @@ void ptx_opcocde_latency_options (option_parser_t opp) {
"Opcode latencies for SFU instructions"
"Default 8",
"8");
+ option_parser_register(opp, "-ptx_opcode_latency_tesnor", OPT_CSTR, &opcode_latency_tensor,
+ "Opcode latencies for Tensor instructions"
+ "Default 64",
+ "64");
option_parser_register(opp, "-ptx_opcode_initiation_int", OPT_CSTR, &opcode_initiation_int,
"Opcode initiation intervals for integers <ADD,MAX,MUL,MAD,DIV>"
"Default 1,1,4,4,32",
@@ -100,6 +106,10 @@ void ptx_opcocde_latency_options (option_parser_t opp) {
"Opcode initiation intervals for sfu instructions"
"Default 8",
"8");
+ option_parser_register(opp, "-ptx_opcode_initiation_tensor", OPT_CSTR, &opcode_initiation_tensor,
+ "Opcode initiation intervals for tensor instructions"
+ "Default 64",
+ "64");
option_parser_register(opp, "-cdp_latency", OPT_CSTR, &cdp_latency_str,
"CDP API latency <cudaStreamCreateWithFlags, \
cudaGetParameterBufferV2_init_perWarp, cudaGetParameterBufferV2_perKernel, \
@@ -113,22 +123,35 @@ static address_type get_converge_point(address_type pc);
void gpgpu_t::gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref, int dim, int readmode, int ext)
{
std::string texname(name);
- m_NameToTextureRef[texname] = texref;
+ if (m_NameToTextureRef.find(texname)==m_NameToTextureRef.end()){
+ m_NameToTextureRef[texname] = std::set<const struct textureReference*>();
+ }else{
+ const struct textureReference* tr = *m_NameToTextureRef[texname].begin();
+ assert(tr!=NULL);
+ //asserts that all texrefs in set have same fields
+ assert(tr->normalized==texref->normalized&&
+ tr->filterMode==texref->filterMode&&
+ tr->addressMode[0]==texref->addressMode[0]&&
+ tr->addressMode[1]==texref->addressMode[1]&&
+ tr->addressMode[2]==texref->addressMode[2]&&
+ tr->channelDesc.x==texref->channelDesc.x&&
+ tr->channelDesc.y==texref->channelDesc.y&&
+ tr->channelDesc.z==texref->channelDesc.z&&
+ tr->channelDesc.w==texref->channelDesc.w&&
+ tr->channelDesc.f==texref->channelDesc.f
+ );
+ }
+ m_NameToTextureRef[texname].insert(texref);
+ m_TextureRefToName[texref] = texname;
const textureReferenceAttr *texAttr = new textureReferenceAttr(texref, dim, (enum cudaTextureReadMode)readmode, ext);
- m_TextureRefToAttribute[texref] = texAttr;
+ m_NameToAttribute[texname] = texAttr;
}
const char* gpgpu_t::gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref)
{
- std::map<std::string, const struct textureReference*>::iterator itr = m_NameToTextureRef.begin();
- while (itr != m_NameToTextureRef.end()) {
- if ((*itr).second == texref) {
- const char *p = ((*itr).first).c_str();
- return p;
- }
- itr++;
- }
- return NULL;
+ std::map<const struct textureReference*, std::string>::const_iterator t=m_TextureRefToName.find(texref);
+ assert( t != m_TextureRefToName.end() );
+ return t->second.c_str();
}
unsigned int intLOGB2( unsigned int v ) {
@@ -148,7 +171,14 @@ unsigned int intLOGB2( unsigned int v ) {
void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array)
{
- m_TextureRefToCudaArray[texref] = array;
+ std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref);
+
+ std::map<std::string,const struct cudaArray*>::const_iterator t=m_NameToCudaArray.find(texname);
+ //check that there's nothing there first
+ if(t != m_NameToCudaArray.end()){
+ printf("GPGPU-Sim PTX: Warning: binding to texref associated with %s, which was previously bound.\nImplicitly unbinding texref associated to %s first\n", texname.c_str(), texname.c_str());
+ }
+ m_NameToCudaArray[texname] = array;
unsigned int texel_size_bits = array->desc.w + array->desc.x + array->desc.y + array->desc.z;
unsigned int texel_size = texel_size_bits/8;
unsigned int Tx, Ty;
@@ -188,7 +218,15 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te
texInfo->Ty_numbits = intLOGB2(Ty);
texInfo->texel_size = texel_size;
texInfo->texel_size_numbits = intLOGB2(texel_size);
- m_TextureRefToTexureInfo[texref] = texInfo;
+ m_NameToTextureInfo[texname] = texInfo;
+}
+
+void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref)
+{
+ //assumes bind-use-unbind-bind-use-unbind pattern
+ std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref);
+ m_NameToCudaArray.erase(texname);
+ m_NameToTextureInfo.erase(texname);
}
unsigned g_assemble_code_next_pc=0;
@@ -220,7 +258,9 @@ void function_info::ptx_assemble()
m_start_PC = PC;
addr_t n=0; // offset in m_instr_mem
- s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + MAX_INST_SIZE*m_instructions.size());
+ //Why s_g_pc_to_insn.size() is needed to reserve additional memory for insts? reserve is cumulative.
+ //s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + MAX_INST_SIZE*m_instructions.size());
+ s_g_pc_to_insn.reserve(MAX_INST_SIZE*m_instructions.size());
for ( i=m_instructions.begin(); i != m_instructions.end(); i++ ) {
ptx_instruction *pI = *i;
if ( pI->is_label() ) {
@@ -258,11 +298,13 @@ void function_info::ptx_assemble()
target.set_type(label_t);
}
}
-
+ m_n = n;
printf(" done.\n");
fflush(stdout);
- printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", m_name.c_str() );
+ //disable pdom analysis here and do it at runtime
+#if 0
+ printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", m_name.c_str() );
create_basic_blocks();
connect_basic_blocks();
bool modified = false;
@@ -296,6 +338,7 @@ void function_info::ptx_assemble()
fflush(stdout);
m_assembled = true;
+#endif
}
addr_t shared_to_generic( unsigned smid, addr_t addr )
@@ -536,7 +579,7 @@ void ptx_instruction::set_mul_div_or_other_archop(){
sp_op=FP_EXP_OP;
break;
default:
- if(op==ALU_OP)
+ if((op==ALU_OP)||(op==TENSOR_CORE_OP))
sp_op=FP__OP;
break;
@@ -558,7 +601,7 @@ void ptx_instruction::set_mul_div_or_other_archop(){
sp_op=INT_DIV_OP;
break;
default:
- if(op==ALU_OP)
+ if((op==ALU_OP))
sp_op=INT__OP;
break;
}
@@ -597,6 +640,9 @@ void ptx_instruction::set_bar_type()
abort();
}
}
+ else if(m_opcode==SST_OP) {
+ bar_type = SYNC;
+ }
}
@@ -606,10 +652,12 @@ void ptx_instruction::set_opcode_and_latency()
unsigned fp_latency[5];
unsigned dp_latency[5];
unsigned sfu_latency;
+ unsigned tensor_latency;
unsigned int_init[5];
unsigned fp_init[5];
unsigned dp_init[5];
unsigned sfu_init;
+ unsigned tensor_init;
/*
* [0] ADD,SUB
* [1] MAX,Min
@@ -628,6 +676,8 @@ void ptx_instruction::set_opcode_and_latency()
&dp_latency[3],&dp_latency[4]);
sscanf(opcode_latency_sfu, "%u",
&sfu_latency);
+ sscanf(opcode_latency_tensor, "%u",
+ &tensor_latency);
sscanf(opcode_initiation_int, "%u,%u,%u,%u,%u",
&int_init[0],&int_init[1],&int_init[2],
&int_init[3],&int_init[4]);
@@ -639,6 +689,8 @@ void ptx_instruction::set_opcode_and_latency()
&dp_init[3],&dp_init[4]);
sscanf(opcode_initiation_sfu, "%u",
&sfu_init);
+ sscanf(opcode_initiation_tensor, "%u",
+ &tensor_init);
sscanf(cdp_latency_str, "%u,%u,%u,%u,%u",
&cdp_latency[0],&cdp_latency[1],&cdp_latency[2],
&cdp_latency[3],&cdp_latency[4]);
@@ -662,13 +714,16 @@ void ptx_instruction::set_opcode_and_latency()
if ( has_memory_write() ) op = STORE_OP;
break;
case LD_OP: op = LOAD_OP; break;
+ case MMA_LD_OP: op = TENSOR_CORE_LOAD_OP; break;
case LDU_OP: op = LOAD_OP; break;
case ST_OP: op = STORE_OP; break;
+ case MMA_ST_OP: op = TENSOR_CORE_STORE_OP; break;
case BRA_OP: op = BRANCH_OP; break;
case BREAKADDR_OP: op = BRANCH_OP; break;
case TEX_OP: op = LOAD_OP; mem_op=TEX; break;
case ATOM_OP: op = LOAD_OP; break;
case BAR_OP: op = BARRIER_OP; break;
+ case SST_OP: op = BARRIER_OP; break;
case MEMBAR_OP: op = MEMORY_BARRIER_OP; break;
case CALL_OP:
{
@@ -695,6 +750,7 @@ void ptx_instruction::set_opcode_and_latency()
case F32_TYPE:
latency = fp_latency[0];
initiation_interval = fp_init[0];
+ op = SP_OP;
break;
case F64_TYPE:
case FF64_TYPE:
@@ -708,6 +764,7 @@ void ptx_instruction::set_opcode_and_latency()
default: //Use int settings for default
latency = int_latency[0];
initiation_interval = int_init[0];
+ op = INTP_OP;
break;
}
break;
@@ -717,6 +774,7 @@ void ptx_instruction::set_opcode_and_latency()
case F32_TYPE:
latency = fp_latency[1];
initiation_interval = fp_init[1];
+ op = SP_OP;
break;
case F64_TYPE:
case FF64_TYPE:
@@ -730,6 +788,7 @@ void ptx_instruction::set_opcode_and_latency()
default: //Use int settings for default
latency = int_latency[1];
initiation_interval = int_init[1];
+ op = INTP_OP;
break;
}
break;
@@ -739,6 +798,7 @@ void ptx_instruction::set_opcode_and_latency()
case F32_TYPE:
latency = fp_latency[2];
initiation_interval = fp_init[2];
+ op = SP_OP;
break;
case F64_TYPE:
case FF64_TYPE:
@@ -752,7 +812,7 @@ void ptx_instruction::set_opcode_and_latency()
default: //Use int settings for default
latency = int_latency[2];
initiation_interval = int_init[2];
- op = SFU_OP;
+ op = INTP_OP;
break;
}
break;
@@ -762,6 +822,7 @@ void ptx_instruction::set_opcode_and_latency()
case F32_TYPE:
latency = fp_latency[3];
initiation_interval = fp_init[3];
+ op = SP_OP;
break;
case F64_TYPE:
case FF64_TYPE:
@@ -775,7 +836,7 @@ void ptx_instruction::set_opcode_and_latency()
default: //Use int settings for default
latency = int_latency[3];
initiation_interval = int_init[3];
- op = SFU_OP;
+ op = INTP_OP;
break;
}
break;
@@ -802,13 +863,17 @@ void ptx_instruction::set_opcode_and_latency()
}
break;
case SQRT_OP: case SIN_OP: case COS_OP: case EX2_OP: case LG2_OP: case RSQRT_OP: case RCP_OP:
- //Using double to approximate those
latency = sfu_latency;
initiation_interval = sfu_init;
op = SFU_OP;
break;
+ case MMA_OP:
+ latency = tensor_latency;
+ initiation_interval = tensor_init;
+ op=TENSOR_CORE_OP;
+ break;
case SHFL_OP:
- latency = 32;
+ latency = 4;
initiation_interval = 4;
break;
default:
@@ -863,10 +928,14 @@ void ptx_instruction::pre_decode()
{
pc = m_PC;
isize = m_inst_size;
- for( unsigned i=0; i<4; i++) {
+ for(unsigned i=0; i<MAX_OUTPUT_VALUES; i++) {
out[i] = 0;
+ }
+ for(unsigned i=0; i<MAX_INPUT_VALUES; i++) {
in[i] = 0;
}
+ incount=0;
+ outcount=0;
is_vectorin = 0;
is_vectorout = 0;
std::fill_n(arch_reg.src, MAX_REG_OPERANDS, -1);
@@ -898,6 +967,7 @@ void ptx_instruction::pre_decode()
switch( m_cache_option ) {
case CA_OPTION: cache_op = CACHE_ALL; break;
+ case NC_OPTION: cache_op = CACHE_L1; break;
case CG_OPTION: cache_op = CACHE_GLOBAL; break;
case CS_OPTION: cache_op = CACHE_STREAMING; break;
case LU_OPTION: cache_op = CACHE_LAST_USE; break;
@@ -905,9 +975,11 @@ void ptx_instruction::pre_decode()
case WB_OPTION: cache_op = CACHE_WRITE_BACK; break;
case WT_OPTION: cache_op = CACHE_WRITE_THROUGH; break;
default:
- if( m_opcode == LD_OP || m_opcode == LDU_OP )
+ //if( m_opcode == LD_OP || m_opcode == LDU_OP )
+ if( m_opcode == MMA_LD_OP || m_opcode == LD_OP || m_opcode == LDU_OP )
cache_op = CACHE_ALL;
- else if( m_opcode == ST_OP )
+ //else if( m_opcode == ST_OP )
+ else if( m_opcode == MMA_ST_OP || m_opcode == ST_OP )
cache_op = CACHE_WRITE_BACK;
else if( m_opcode == ATOM_OP )
cache_op = CACHE_GLOBAL;
@@ -933,6 +1005,10 @@ void ptx_instruction::pre_decode()
if( num_elem >= 2 ) out[1] = o.reg2_num();
if( num_elem >= 3 ) out[2] = o.reg3_num();
if( num_elem >= 4 ) out[3] = o.reg4_num();
+ if( num_elem >= 5 ) out[4] = o.reg5_num();
+ if( num_elem >= 6 ) out[5] = o.reg6_num();
+ if( num_elem >= 7 ) out[6] = o.reg7_num();
+ if( num_elem >= 8 ) out[7] = o.reg8_num();
for (int i = 0; i < num_elem; i++)
arch_reg.dst[i] = o.arch_reg_num(i);
}
@@ -951,16 +1027,29 @@ void ptx_instruction::pre_decode()
//assert(m == 0); //only support 1 vector operand (for textures) right now
is_vectorout = 1;
unsigned num_elem = o.get_vect_nelem();
- if( num_elem >= 1 ) in[0] = o.reg1_num();
- if( num_elem >= 2 ) in[1] = o.reg2_num();
- if( num_elem >= 3 ) in[2] = o.reg3_num();
- if( num_elem >= 4 ) in[3] = o.reg4_num();
+ if( num_elem >= 1 ) in[m+0] = o.reg1_num();
+ if( num_elem >= 2 ) in[m+1] = o.reg2_num();
+ if( num_elem >= 3 ) in[m+2] = o.reg3_num();
+ if( num_elem >= 4 ) in[m+3] = o.reg4_num();
+ if( num_elem >= 5 ) in[m+4] = o.reg5_num();
+ if( num_elem >= 6 ) in[m+5] = o.reg6_num();
+ if( num_elem >= 7 ) in[m+6] = o.reg7_num();
+ if( num_elem >= 8 ) in[m+7] = o.reg8_num();
for (int i = 0; i < num_elem; i++)
- arch_reg.src[i] = o.arch_reg_num(i);
- m+=4;
+ arch_reg.src[m+i] = o.arch_reg_num(i);
+ m+=num_elem;
}
}
}
+
+ //Setting number of input and output operands which is required for scoreboard check
+ for(int i=0;i<MAX_OUTPUT_VALUES;i++)
+ if(out[i]>0)
+ outcount++;
+
+ for(int i=0;i<MAX_INPUT_VALUES;i++)
+ if(in[i]>0)
+ incount++;
// Get predicate
if(has_pred()) {
@@ -1158,14 +1247,19 @@ void function_info::finalize( memory_space *param_mem )
}
// copy the parameter over word-by-word so that parameter that crosses a memory page can be copied over
//Jin: copy parameter using aligned rules
+ const type_info *paramtype = param->type();
+ int align_amount = paramtype->get_key().get_alignment_spec();
+ align_amount = (align_amount == -1) ? size : align_amount;
+ param_address = (param_address + align_amount - 1) / align_amount * align_amount; //aligned
+
const size_t word_size = 4;
- //param_address = (param_address + size - 1) / size * size; //aligned with size TODO: align not correct
+ //param_address = (param_address + size - 1) / size * size; //aligned with size
for (size_t idx = 0; idx < size; idx += word_size) {
const char *pdata = reinterpret_cast<const char*>(param_value.pdata) + idx; // cast to char * for ptr arithmetic
param_mem->write(param_address + idx, word_size, pdata,NULL,NULL);
}
unsigned offset = p.get_offset();
- //assert(offset == param_address);
+ assert(offset == param_address);
param->set_address(param_address);
param_address += size;
}
@@ -1212,6 +1306,165 @@ void function_info::list_param( FILE *fout ) const
fflush(fout);
}
+void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim)
+{
+ static unsigned long long counter = 0;
+ std::vector< std::pair<size_t, unsigned char*> > param_data;
+ std::vector<unsigned> offsets;
+ std::vector<bool> paramIsPointer;
+
+ char * gpgpusim_path = getenv("GPGPUSIM_ROOT");
+ assert(gpgpusim_path!=NULL);
+ char * wys_exec_path = getenv("WYS_EXEC_PATH");
+ assert(wys_exec_path!=NULL);
+ std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data";
+ std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter));
+
+ //initialize paramList
+ char buff[1024];
+ std::string filename_c(filename+"_c");
+ snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str());
+ system(buff);
+ FILE *fp = fopen(filename_c.c_str(), "r");
+ fgets(buff, 1024, fp);
+ fclose(fp);
+ std::string fn(buff);
+ size_t pos1, pos2;
+ pos1 = fn.find_last_of("(");
+ pos2 = fn.find(")", pos1);
+ assert(pos2>pos1&&pos1>0);
+ strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str());
+ char *tok;
+ tok = strtok(buff, ",");
+ std::string tmp;
+ while(tok!=NULL){
+ std::string param(tok);
+ if(param.find("<")!=std::string::npos){
+ assert(param.find(">")==std::string::npos);
+ assert(param.find("*")==std::string::npos);
+ tmp = param;
+ } else {
+ if (tmp.length()>0){
+ tmp = "";
+ assert(param.find(">")!=std::string::npos);
+ assert(param.find("<")==std::string::npos);
+ assert(param.find("*")==std::string::npos);
+ }
+ printf("%s\n", param.c_str());
+ if(param.find("*")!=std::string::npos){
+ paramIsPointer.push_back(true);
+ }else{
+ paramIsPointer.push_back(false);
+ }
+ }
+ tok = strtok(NULL, ",");
+ }
+
+
+ for( std::map<unsigned,param_info>::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
+ param_info &p = i->second;
+ std::string name = p.get_name();
+ symbol *param = m_symtab->lookup(name.c_str());
+ addr_t param_addr = param->get_address();
+ param_t param_value = p.get_value();
+ offsets.push_back((unsigned)p.get_offset());
+
+ if (paramIsPointer[i->first] && (*(unsigned long long*)param_value.pdata != 0)){
+ //is pointer
+ assert(param_value.size==sizeof(void*)&&"MisID'd this param as pointer");
+ size_t array_size = 0;
+ unsigned long long param_pointer = *(unsigned long long*)param_value.pdata;
+ if(mallocPtr_Size.find(param_pointer)!=mallocPtr_Size.end()){
+ array_size = mallocPtr_Size[param_pointer];
+ }else{
+ for( std::map<unsigned long long, size_t>::iterator j=mallocPtr_Size.begin(); j!=mallocPtr_Size.end(); j++ ) {
+ if(param_pointer>j->first&&param_pointer<j->first + j->second){
+ array_size = j->first + j->second - param_pointer;
+ break;
+ }
+ }
+ assert(array_size>0&&"pointer was not previously malloc'd");
+ }
+
+ unsigned char* val = (unsigned char*) malloc(param_value.size);
+ param_mem->read(param_addr,param_value.size,(void*)val);
+ unsigned char* array_val = (unsigned char*) malloc(array_size);
+ gpu->get_global_memory()->read(*(unsigned*)((void*)val),array_size,(void*)array_val);
+ param_data.push_back(std::pair<size_t, unsigned char*>(array_size,array_val));
+ paramIsPointer.push_back(true);
+ }else{
+ unsigned char* val = (unsigned char*) malloc(param_value.size);
+ param_mem->read(param_addr,param_value.size,(void*)val);
+ param_data.push_back(std::pair<size_t, unsigned char*>(param_value.size,val));
+ paramIsPointer.push_back(false);
+ }
+ }
+
+ FILE *fout = fopen (filename.c_str(), "w");
+ printf("Writing data to %s ...\n", filename.c_str());
+ fprintf(fout, "%s\n", get_name().c_str());
+ fprintf(fout, "%u,%u,%u %u,%u,%u\n", gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z);
+ size_t index = 0;
+ for( std::vector< std::pair<size_t,unsigned char*> >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) {
+ if (paramIsPointer[index]){
+ fprintf(fout, "*");
+ }
+ fprintf(fout, "%lu :", i->first);
+ for (size_t j = 0; j<i->first; j++){
+ fprintf(fout, " %u", i->second[j]);
+ }
+ fprintf(fout, " : %u", offsets[index]);
+ free (i->second);
+ fprintf(fout, "\n");
+ index++;
+ }
+ fflush(fout);
+ fclose(fout);
+
+ //ptx config
+ std::string ptx_config_fn(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/ptx.config" + std::to_string(counter));
+ snprintf(buff, 1024, "grep -rn \".entry %s\" %s/*.ptx | cut -d \":\" -f 1-2 > %s", get_name().c_str(), wys_exec_path, ptx_config_fn.c_str());
+ if (system(buff)!=0){
+ printf("WARNING: Failed to execute grep to find ptx source \n");
+ printf("Problematic call: %s", buff);
+ abort();
+ }
+ FILE *fin = fopen(ptx_config_fn.c_str(), "r");
+ char ptx_source[256];
+ unsigned line_number;
+ int numscanned = fscanf(fin, "%[^:]:%u", ptx_source, &line_number);
+ assert(numscanned == 2);
+ fclose(fin);
+ snprintf(buff, 1024, "grep -rn \".version\" %s | cut -d \":\" -f 1 | xargs -I \"{}\" awk \"NR>={}&&NR<={}+2\" %s > %s", ptx_source, ptx_source, ptx_config_fn.c_str());
+ if (system(buff)!=0){
+ printf("WARNING: Failed to execute grep to find ptx header \n");
+ printf("Problematic call: %s", buff);
+ abort();
+ }
+ fin = fopen(ptx_source, "r");
+ assert(fin!=NULL);
+ printf("Writing data to %s ...\n", ptx_config_fn.c_str());
+ fout = fopen(ptx_config_fn.c_str(), "a");
+ assert(fout!=NULL);
+ for (unsigned i = 0; i<line_number; i++){
+ fgets(buff, 1024, fin);
+ assert(!feof(fin));
+ }
+ fprintf(fout, "\n\n");
+ do{
+ fprintf(fout, "%s", buff);
+ fgets(buff, 1024, fin);
+ if(feof(fin)){
+ break;
+ }
+ } while(strstr(buff, "entry")==NULL);
+
+ fclose(fin);
+ fflush(fout);
+ fclose(fout);
+ counter++;
+}
+
template<int activate_level>
bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc)
{
@@ -1254,13 +1507,27 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th
std::string texname = src1.name();
gpgpu_t *gpu = thread->get_gpu();
- const struct textureReference* texref = gpu->get_texref(texname);
- const struct textureInfo* texInfo = gpu->get_texinfo(texref);
+ /*
+ For programs with many streams, textures can be bound and unbound
+ asynchronously. This means we need to use the kernel's "snapshot" of
+ the state of the texture mappings when it was launched (so that we
+ don't try to access the incorrect texture mapping if it's been updated,
+ or that we don't access a mapping that has been unbound).
+ */
+ kernel_info_t& k = thread->get_kernel();
+ const struct textureInfo* texInfo = k.get_texinfo(texname);
unsigned data_size = texInfo->texel_size;
return data_size;
}
+int tensorcore_op(int inst_opcode){
+
+ if((inst_opcode==MMA_OP)||(inst_opcode==MMA_LD_OP)||(inst_opcode==MMA_ST_OP))
+ return 1;
+ else
+ return 0;
+}
void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
{
@@ -1269,6 +1536,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
addr_t pc = next_instr();
assert( pc == inst.pc ); // make sure timing model and functional model are in sync
const ptx_instruction *pI = m_func_info->get_instruction(pc);
+
set_npc( pc + pI->inst_size() );
@@ -1301,6 +1569,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
skip = !pred_lookup(pI->get_pred_mod(), pred_value.pred & 0x000F);
}
}
+ int inst_opcode=pI->get_opcode();
if( skip ) {
inst.set_not_active(lane_id);
@@ -1312,13 +1581,25 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
*((warp_inst_t*)pJ) = inst; // copy active mask information
pI = pJ;
}
- switch ( pI->get_opcode() ) {
-#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break;
-#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break;
-#include "opcodes.def"
-#undef OP_DEF
-#undef OP_W_DEF
- default: printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() ); break;
+
+ if(((inst_opcode==MMA_OP||inst_opcode==MMA_LD_OP||inst_opcode==MMA_ST_OP))){
+ if(inst.active_count()!=MAX_WARP_SIZE)
+ {
+ printf("Tensor Core operation are warp synchronous operation. All the threads needs to be active.");
+ assert(0);
+ }
+ }
+
+ //Tensorcore is warp synchronous operation. So these instructions needs to be executed only once. To make the simulation faster removing the redundant tensorcore operation
+ if(!tensorcore_op(inst_opcode)||(tensorcore_op(inst_opcode))&&(lane_id==0)){
+ switch ( inst_opcode ) {
+ #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break;
+ #define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break;
+ #include "opcodes.def"
+ #undef OP_DEF
+ #undef OP_W_DEF
+ default: printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() ); break;
+ }
}
delete pJ;
pI = pI_saved;
@@ -1361,13 +1642,16 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
_memory_op_t insn_memory_op = no_memory_op;
unsigned insn_data_size = 0;
if ( (pI->has_memory_read() || pI->has_memory_write()) ) {
- insn_memaddr = last_eaddr();
- insn_space = last_space();
- unsigned to_type = pI->get_type();
- insn_data_size = datatype2size(to_type);
- insn_memory_op = pI->has_memory_read() ? memory_load : memory_store;
+ if(!((inst_opcode==MMA_LD_OP||inst_opcode==MMA_ST_OP)))
+ {
+ insn_memaddr = last_eaddr();
+ insn_space = last_space();
+ unsigned to_type = pI->get_type();
+ insn_data_size = datatype2size(to_type);
+ insn_memory_op = pI->has_memory_read() ? memory_load : memory_store;
+ }
}
-
+
if ( pI->get_opcode() == BAR_OP && pI->barrier_op() == RED_OPTION) {
inst.add_callback( lane_id, last_callback().function, last_callback().instruction, this,false /*not atomic*/);
}
@@ -1439,12 +1723,15 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
// "Return values"
if(!skip) {
- inst.space = insn_space;
- inst.set_addr(lane_id, insn_memaddr);
- inst.data_size = insn_data_size; // simpleAtomicIntrinsics
- assert( inst.memory_op == insn_memory_op );
- }
-
+ if(!((inst_opcode==MMA_LD_OP||inst_opcode==MMA_ST_OP)))
+ {
+ inst.space = insn_space;
+ inst.set_addr(lane_id, insn_memaddr);
+ inst.data_size = insn_data_size; // simpleAtomicIntrinsics
+ assert( inst.memory_op == insn_memory_op );
+ }
+ }
+
} catch ( int x ) {
printf("GPGPU-Sim PTX: ERROR (%d) executing intruction (%s:%u)\n", x, pI->source_file(), pI->source_line() );
printf("GPGPU-Sim PTX: '%s'\n", pI->get_source() );
@@ -1483,6 +1770,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
std::list<ptx_thread_info *> &active_threads = kernel.active_threads();
static std::map<unsigned,memory_space*> shared_memory_lookup;
+ static std::map<unsigned,memory_space*> sstarr_memory_lookup;
static std::map<unsigned,ptx_cta_info*> ptx_cta_lookup;
static std::map<unsigned,ptx_warp_info*> ptx_warp_lookup;
static std::map<unsigned,std::map<unsigned,memory_space*> > local_memory_lookup;
@@ -1527,6 +1815,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
//initializing new CTA
ptx_cta_info *cta_info = NULL;
memory_space *shared_mem = NULL;
+ memory_space *sstarr_mem = NULL;
unsigned cta_size = kernel.threads_per_cta();
unsigned max_cta_per_sm = num_threads/cta_size; // e.g., 256 / 48 = 5
@@ -1544,6 +1833,9 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
snprintf(buf,512,"shared_%u", sid);
shared_mem = new memory_space_impl<16*1024>(buf,4);
shared_memory_lookup[sm_idx] = shared_mem;
+ snprintf(buf,512,"sstarr_%u", sid);
+ sstarr_mem = new memory_space_impl<16*1024>(buf,4);
+ sstarr_memory_lookup[sm_idx] = sstarr_mem;
cta_info = new ptx_cta_info(sm_idx);
ptx_cta_lookup[sm_idx] = cta_info;
} else {
@@ -1552,6 +1844,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
sm_idx, sid, max_cta_per_sm );
}
shared_mem = shared_memory_lookup[sm_idx];
+ sstarr_mem = sstarr_memory_lookup[sm_idx];
cta_info = ptx_cta_lookup[sm_idx];
cta_info->check_cta_thread_status_and_reset();
}
@@ -1564,7 +1857,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
kernel.increment_thread_id();
new_tid += tid;
ptx_thread_info *thd = new ptx_thread_info(kernel);
-
ptx_warp_info *warp_info = NULL;
if ( ptx_warp_lookup.find(hw_warp_id) == ptx_warp_lookup.end() ) {
warp_info = new ptx_warp_info();
@@ -1593,9 +1885,11 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
thd->cpy_tid_to_reg(tid3d);
thd->set_valid();
thd->m_shared_mem = shared_mem;
+ thd->m_sstarr_mem = sstarr_mem;
function_info *finfo = thd->func_info();
symbol_table *st = finfo->get_symtab();
thd->func_info()->param_to_shared(thd->m_shared_mem,st);
+ thd->func_info()->param_to_shared(thd->m_sstarr_mem,st);
thd->m_cta_info = cta_info;
cta_info->add_thread(thd);
thd->m_local_mem = local_mem;
@@ -1632,7 +1926,7 @@ kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
struct dim3 blockDim,
gpgpu_t *gpu )
{
- kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+ kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping());
unsigned argcount=args.size();
unsigned argn=1;
for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
@@ -1813,6 +2107,38 @@ ptx_cta_info *g_func_cta_info = NULL;
#define MAX(a,b) (((a)>(b))?(a):(b))
+unsigned max_cta (const struct gpgpu_ptx_sim_info *kernel_info, unsigned threads_per_cta, unsigned int warp_size, unsigned int n_thread_per_shader, unsigned int gpgpu_shmem_size, unsigned int gpgpu_shader_registers, unsigned int max_cta_per_core)
+{
+
+ unsigned int padded_cta_size = threads_per_cta;
+ if (padded_cta_size%warp_size)
+ padded_cta_size = ((padded_cta_size/warp_size)+1)*(warp_size);
+ unsigned int result_thread = n_thread_per_shader / padded_cta_size;
+
+ unsigned int result_shmem = (unsigned)-1;
+ if (kernel_info->smem > 0)
+ result_shmem = gpgpu_shmem_size / kernel_info->smem;
+ unsigned int result_regs = (unsigned)-1;
+ if (kernel_info->regs > 0)
+ result_regs = gpgpu_shader_registers / (padded_cta_size * ((kernel_info->regs+3)&~3));
+ printf("padded cta size is %d and %d and %d",padded_cta_size, kernel_info->regs, ((kernel_info->regs+3)&~3) );
+ //Limit by CTA
+ unsigned int result_cta = max_cta_per_core;
+
+ unsigned result = result_thread;
+ result = gs_min2(result, result_shmem);
+ result = gs_min2(result, result_regs);
+ result = gs_min2(result, result_cta);
+
+ printf ("GPGPU-Sim uArch: CTA/core = %u, limited by:", result);
+ if (result == result_thread) printf (" threads");
+ if (result == result_shmem) printf (" shmem");
+ if (result == result_regs) printf (" regs");
+ if (result == result_cta) printf (" cta_limit");
+ printf ("\n");
+
+ return result;
+}
/*!
This function simulates the CUDA code functionally, it takes a kernel_info_t parameter
which holds the data for the CUDA kernel to be executed
@@ -1823,20 +2149,71 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL )
//using a shader core object for book keeping, it is not needed but as most function built for performance simulation need it we use it here
extern gpgpu_sim *g_the_gpu;
+ //before we execute, we should do PDOM analysis for functional simulation scenario.
+ function_info *kernel_func_info = kernel.entry();
+ const struct gpgpu_ptx_sim_info *kernel_info = ptx_sim_kernel_info(kernel_func_info);
+ checkpoint *g_checkpoint;
+ g_checkpoint = new checkpoint();
+
+ if (kernel_func_info->is_pdom_set()) {
+ printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kernel.name().c_str() );
+ } else {
+ printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kernel.name().c_str() );
+ kernel_func_info->do_pdom();
+ kernel_func_info->set_pdom();
+ }
+
+ unsigned max_cta_tot = max_cta(kernel_info,kernel.threads_per_cta(), g_the_gpu->getShaderCoreConfig()->warp_size, g_the_gpu->getShaderCoreConfig()->n_thread_per_shader, g_the_gpu->getShaderCoreConfig()->gpgpu_shmem_size, g_the_gpu->getShaderCoreConfig()->gpgpu_shader_registers, g_the_gpu->getShaderCoreConfig()->max_cta_per_core);
+ printf("Max CTA : %d\n",max_cta_tot);
+
+
+
+
+
+ int inst_count=50;
+ int cp_op= g_the_gpu->checkpoint_option;
+ int cp_CTA = g_the_gpu->checkpoint_CTA;
+ int cp_kernel= g_the_gpu->checkpoint_kernel;
+ cp_count= g_the_gpu->checkpoint_insn_Y;
+ cp_cta_resume= g_the_gpu->checkpoint_CTA_t;
+ int cta_launched =0;
//we excute the kernel one CTA (Block) at the time, as synchronization functions work block wise
while(!kernel.no_more_ctas_to_run()){
- functionalCoreSim cta(
- &kernel,
- g_the_gpu,
- g_the_gpu->getShaderCoreConfig()->warp_size
- );
- cta.execute();
+ unsigned temp=kernel.get_next_cta_id_single();
+
-#if (CUDART_VERSION >= 5000)
- launch_all_device_kernels();
-#endif
+ if(cp_op==0 || (cp_op==1 && cta_launched<cp_cta_resume && kernel.get_uid()==cp_kernel) || kernel.get_uid()< cp_kernel) // just fro testing
+ {
+ functionalCoreSim cta(
+ &kernel,
+ g_the_gpu,
+ g_the_gpu->getShaderCoreConfig()->warp_size
+ );
+ cta.execute(cp_count,temp);
+
+ #if (CUDART_VERSION >= 5000)
+ launch_all_device_kernels();
+ #endif
+ }
+ else
+ {
+ kernel.increment_cta_id();
+ }
+ cta_launched++;
}
+
+
+
+ if(cp_op==1)
+ {
+ char f1name[2048];
+ snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", kernel.get_uid() );
+ g_checkpoint->store_global_mem(g_the_gpu->get_global_memory(), f1name , "%08x");
+ }
+
+
+
//registering this kernel as done
@@ -1874,9 +2251,10 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL )
fflush(stdout);
}
-void functionalCoreSim::initializeCTA()
+void functionalCoreSim::initializeCTA(unsigned ctaid_cp)
{
int ctaLiveThreads=0;
+ symbol_table * symtab= m_kernel->entry()->get_symtab();
for(int i=0; i< m_warp_count; i++){
m_warpAtBarrier[i]=false;
@@ -1889,9 +2267,13 @@ void functionalCoreSim::initializeCTA()
for(unsigned i=0; i<m_kernel->threads_per_cta();i++) {
ptx_sim_init_thread(*m_kernel,&m_thread[i],0,i,m_kernel->threads_per_cta()-i,m_kernel->threads_per_cta(),this,0,i/m_warp_size,(gpgpu_t*)m_gpu, true);
assert(m_thread[i]!=NULL && !m_thread[i]->is_done());
+ char fname[2048];
+ snprintf(fname,2048,"checkpoint_files/thread_%d_0_reg.txt",i );
+ if(cp_cta_resume==1)
+ m_thread[i]->resume_reg_thread(fname,symtab);
ctaLiveThreads++;
}
-
+
for(int k=0;k<m_warp_count;k++)
createWarp(k);
}
@@ -1908,27 +2290,91 @@ void functionalCoreSim::createWarp(unsigned warpId)
assert(m_thread[warpId*m_warp_size]!=NULL);
m_simt_stack[warpId]->launch(m_thread[warpId*m_warp_size]->get_pc(),initialMask);
+ char fname[2048];
+ snprintf(fname,2048,"checkpoint_files/warp_%d_0_simt.txt",warpId );
+
+ if(cp_cta_resume==1)
+ {
+ unsigned pc,rpc;
+ m_simt_stack[warpId]->resume(fname);
+ m_simt_stack[warpId]->get_pdom_stack_top_info(&pc,&rpc);
+ for(int i=warpId*m_warp_size; i<warpId*m_warp_size+m_warp_size;i++){
+ m_thread[i]->set_npc(pc);
+ m_thread[i]->update_pc();
+ }
+
+ }
m_liveThreadCount[warpId]= liveThreadsCount;
}
-void functionalCoreSim::execute()
+void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp)
{
- initializeCTA();
+ cp_count= m_gpu->checkpoint_insn_Y;
+ cp_cta_resume= m_gpu->checkpoint_CTA_t;
+ initializeCTA(ctaid_cp);
- //start executing the CTA
+ int count=0;
while(true){
bool someOneLive= false;
bool allAtBarrier = true;
for(unsigned i=0;i<m_warp_count;i++){
executeWarp(i,allAtBarrier,someOneLive);
+ count++;
}
+
+ if(inst_count>0 && count>inst_count && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cp<m_gpu->checkpoint_CTA_t) && m_gpu->checkpoint_option==1)
+ {
+ someOneLive=false;
+ break;
+ }
if(!someOneLive) break;
if(allAtBarrier){
for(unsigned i=0;i<m_warp_count;i++)
m_warpAtBarrier[i]=false;
}
}
- }
+
+ checkpoint *g_checkpoint;
+ g_checkpoint = new checkpoint();
+
+ symbol * sym;
+ ptx_reg_t regval;
+ regval.u64= 123;
+ symbol_table * symtab= m_kernel->entry()->get_symtab();
+
+
+ unsigned ctaid =m_kernel->get_next_cta_id_single();
+ if(m_gpu->checkpoint_option==1 && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cp<m_gpu->checkpoint_CTA_t))
+ {
+ char fname[2048];
+ snprintf(fname,2048,"checkpoint_files/shared_mem_%d.txt",ctaid-1 );
+ g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , "%08x");
+ for(int i=0; i<32*m_warp_count;i++)
+ {
+ char fname[2048];
+ snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i,ctaid-1 );
+ m_thread[i]->print_reg_thread(fname);
+ char f1name[2048];
+ snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i,ctaid-1 );
+ g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , "%08x");
+ m_thread[i]->set_done();
+ m_thread[i]->exitCore();
+ m_thread[i]->registerExit();
+ }
+
+ for(int i=0;i<m_warp_count;i++)
+ {
+
+ char fname[2048];
+ snprintf(fname,2048,"checkpoint_files/warp_%d_%d_simt.txt",i,ctaid-1 );
+ FILE * fp = fopen(fname,"w");
+ assert(fp!=NULL);
+ m_simt_stack[i]->print_checkpoint(fp);
+ fclose(fp);
+ }
+ }
+
+}
void functionalCoreSim::executeWarp(unsigned i, bool &allAtBarrier, bool & someOneLive)
{
@@ -2091,7 +2537,7 @@ struct rec_pts {
int s_num_recon;
};
-struct std::map<function_info*,rec_pts> g_rpts;
+class std::map<function_info*,rec_pts> g_rpts;
struct rec_pts find_reconvergence_points( function_info *finfo )
{
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 00:27:11 +0000