From bc89f8c7ebc77db6a28f5600d17c120fe693e988 Mon Sep 17 00:00:00 2001 From: J Date: Wed, 8 Aug 2018 16:59:28 -0700 Subject: brev implementation --- src/cuda-sim/instructions.cc | 29 ++++++++++++++++++++++++++++- 1 file changed, 28 insertions(+), 1 deletion(-) (limited to 'src') diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index fe36902..109cfa7 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1459,7 +1459,34 @@ void breakaddr_impl( const ptx_instruction *pI, ptx_thread_info *thread ) assert(pI->has_pred() == false); // pdom analysis cannot handle if this instruction is predicated } -void brev_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void brev_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + ptx_reg_t src1_data, data; + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + + unsigned msb; + switch(i_type){ + case B32_TYPE: + msb = 31; + for (unsigned i=0; i<=msb; i++) { + if((src1_data.u32 & (1 << i))) + data.u32 |= 1 << (msb - i); + } + break; + case B64_TYPE: + msb = 63; + for (unsigned i=0; i<=msb; i++) { + if((src1_data.u64 & (1 << i))) + data.u64 |= 1 << (msb - i); + } + break; + default: assert(0); + } + thread->set_operand_value(dst,data, i_type, thread, pI); +} void brkpt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) -- cgit v1.3 From 1a3de9dba3ddf4adab4256fadf5e815b410862ae Mon Sep 17 00:00:00 2001 From: Jonathan Date: Fri, 10 Aug 2018 16:27:26 -0700 Subject: fixes ptxjitplus and cuLaunchKernel --- .../WatchYourStep/ptxjitplus/ptxjitplus.cpp | 12 ------ src/cuda-sim/cuda-sim.cc | 44 +--------------------- src/cuda-sim/ptx_parser.cc | 4 +- 3 files changed, 5 insertions(+), 55 deletions(-) (limited to 'src') diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp index 1065870..554e831 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp @@ -61,7 +61,6 @@ const char *sSDKname = "PTX Just In Time (JIT) Compilation Plus"; char *wys_exec_path; char *wys_exec_name; char *wys_launch_num; -bool gpgpusim = false; dim3 gridDim, blockDim; std::string kernelName; @@ -138,10 +137,6 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl void initializeData(std::vector& v_params) { - char *gpgpusim_env = getenv("GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN"); - if (gpgpusim_env!=NULL&&gpgpusim_env[0] == '1'){ - gpgpusim=true; - } wys_exec_path = getenv("WYS_EXEC_PATH"); assert(wys_exec_path!=NULL); wys_exec_name = getenv("WYS_EXEC_NAME"); @@ -272,23 +267,16 @@ int main(int argc, char **argv) unsigned char **d_data = (unsigned char **) malloc(sizeof(unsigned char **)); checkCudaErrors(cudaMalloc((void**)d_data, p->size)); checkCudaErrors(cudaMemcpy((void*)*d_data,(void*)p->data,p->size,cudaMemcpyHostToDevice)); - if (gpgpusim){ - checkCudaErrors(cuParamSetv(hKernel, p->offset, d_data, sizeof(*d_data))); - } paramKernels[index] = (void*)d_data; m_device_data[index]=*d_data; m_cleanup[index]=d_data; paramOffset = p->offset + 8; }else{ - if (gpgpusim){ - checkCudaErrors(cuParamSetv(hKernel, p->offset, p->data, p->size)); - } paramKernels[index] = (void*)p->data; paramOffset = p->offset + p->size; } index ++; } - checkCudaErrors(cuParamSetSize(hKernel, paramOffset)); // Launch the kernel (Driver API_) CUDAAPI cuLaunchKernel(hKernel, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index e0cd3f9..5b80616 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1232,6 +1232,7 @@ void function_info::ptx_jit_config(std::map mallocPt std::vector< std::pair > param_data; std::vector offsets; std::vector paramIsPointer; + char buff[1024]; char * gpgpusim_path = getenv("GPGPUSIM_ROOT"); assert(gpgpusim_path!=NULL); @@ -1240,46 +1241,6 @@ void function_info::ptx_jit_config(std::map mallocPt 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); - } - if(param.find("*")!=std::string::npos){ - paramIsPointer.push_back(true); - }else{ - paramIsPointer.push_back(false); - } - } - tok = strtok(NULL, ","); - } - - for( std::map::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(); @@ -1288,9 +1249,8 @@ void function_info::ptx_jit_config(std::map mallocPt param_t param_value = p.get_value(); offsets.push_back((unsigned)p.get_offset()); - if (paramIsPointer[i->first]){ + if(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){ //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()){ diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index e6d6325..3f3485c 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -571,7 +571,9 @@ void add_function_arg() if( g_func_info ) { PTX_PARSE_DPRINTF("add_function_arg \"%s\"", g_last_symbol->name().c_str() ); g_func_info->add_arg(g_last_symbol); - g_func_info->add_config_param( g_size, g_alignment_spec ); + unsigned alignment = (g_alignment_spec==-1) ? g_size : g_alignment_spec; + assert(alignment<=8); + g_func_info->add_config_param( g_size, alignment); } } -- cgit v1.3 From cbe8e3ddbca3895a499882c31b080a58a7545d25 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Thu, 16 Aug 2018 14:10:55 -0700 Subject: fix alignment bug in parser, modified ptxjitconfig, minor fixes --- src/cuda-sim/cuda-sim.cc | 45 ++++++++++++++++++++++++++++++++++++++++++-- src/cuda-sim/instructions.cc | 8 ++++---- src/cuda-sim/ptx_ir.h | 1 + src/cuda-sim/ptx_parser.cc | 2 +- 4 files changed, 49 insertions(+), 7 deletions(-) (limited to 'src') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 5b80616..b85ba95 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1232,7 +1232,6 @@ void function_info::ptx_jit_config(std::map mallocPt std::vector< std::pair > param_data; std::vector offsets; std::vector paramIsPointer; - char buff[1024]; char * gpgpusim_path = getenv("GPGPUSIM_ROOT"); assert(gpgpusim_path!=NULL); @@ -1241,6 +1240,47 @@ void function_info::ptx_jit_config(std::map mallocPt 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::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(); @@ -1249,8 +1289,9 @@ void function_info::ptx_jit_config(std::map mallocPt param_t param_value = p.get_value(); offsets.push_back((unsigned)p.get_offset()); - if(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){ + 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()){ diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 109cfa7..31fc434 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1471,17 +1471,17 @@ void brev_impl( const ptx_instruction *pI, ptx_thread_info *thread ) switch(i_type){ case B32_TYPE: msb = 31; - for (unsigned i=0; i<=msb; i++) { + for (unsigned i=0; i<=msb; i++) { if((src1_data.u32 & (1 << i))) data.u32 |= 1 << (msb - i); - } + } break; case B64_TYPE: msb = 63; - for (unsigned i=0; i<=msb; i++) { + for (unsigned i=0; i<=msb; i++) { if((src1_data.u64 & (1 << i))) data.u64 |= 1 << (msb - i); - } + } break; default: assert(0); } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index e726ab9..768b3e7 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1297,6 +1297,7 @@ public: unsigned offset = 0; if (m_param_configs.size()>0){ unsigned offset_nom = m_param_configs.back().first + m_param_configs.back().second; + //ensure offset matches alignment requirements offset = offset_nom%alignment ? (offset_nom/alignment + 1) * alignment : offset_nom; } m_param_configs.push_back(std::pair(size, offset)); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 3f3485c..0ba9a5b 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -572,7 +572,7 @@ void add_function_arg() PTX_PARSE_DPRINTF("add_function_arg \"%s\"", g_last_symbol->name().c_str() ); g_func_info->add_arg(g_last_symbol); unsigned alignment = (g_alignment_spec==-1) ? g_size : g_alignment_spec; - assert(alignment<=8); + assert(alignment==1||alignment==2||alignment==4||alignment==8||alignment==16);//known valid alignment values g_func_info->add_config_param( g_size, alignment); } -- cgit v1.3