From 46aad91327a265c2fea2cfe629cc38eadb629200 Mon Sep 17 00:00:00 2001 From: speverel Date: Thu, 2 Jun 2016 11:28:15 -0700 Subject: Added handling of .cc option for arithmetic instructions. NOTE: Only made changes to parse instructions. Carry functionality NOT fully implemented; .cc instructions function like their unmodified ueqivelents. Also modified GTX750Ti config to model L1 data cache as simply not being used for global loads (instead of not existing at all). Changed ptxinfo parsing to avoid crashing when info includes texture information. --- src/cuda-sim/cuda-sim.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 715be98..15417d1 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -719,7 +719,7 @@ void ptx_instruction::set_opcode_and_latency() break; } break; - case MAD_OP: case MADP_OP: + case MAD_OP: case MADC_OP: //MAD latency switch(get_type()){ case F32_TYPE: -- cgit v1.3 From 547ce656018a6439e658be3c283721a961b0217f Mon Sep 17 00:00:00 2001 From: sspenst Date: Mon, 13 Jun 2016 01:08:36 -0700 Subject: If ptxas notices any duplicate errors, they now automatically get resolved and the program continues with the duplicate function/variable declarations removed. --- src/cuda-sim/cuda-sim.cc | 19 +++++++ src/cuda-sim/ptx_loader.cc | 126 +++++++++++++++++++++++++++++++++++++++++++-- src/cuda-sim/ptxinfo.l | 4 ++ src/cuda-sim/ptxinfo.y | 16 +++++- 4 files changed, 160 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 15417d1..fba3a59 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1864,10 +1864,14 @@ unsigned translate_pc_to_ptxlineno(unsigned pc) // ptxinfo parser +extern std::map get_duplicate(); + int g_ptxinfo_error_detected; static char *g_ptxinfo_kname = NULL; static struct gpgpu_ptx_sim_info g_ptxinfo; +static std::map g_duplicate; +static const char *g_last_dup_type; const char *get_ptxinfo_kname() { @@ -1897,6 +1901,21 @@ struct gpgpu_ptx_sim_info get_ptxinfo() return g_ptxinfo; } +std::map get_duplicate() +{ + return g_duplicate; +} + +void ptxinfo_linenum( unsigned linenum ) +{ + g_duplicate[linenum] = g_last_dup_type; +} + +void ptxinfo_dup_type( const char *dup_type ) +{ + g_last_dup_type = dup_type; +} + void ptxinfo_function(const char *fname ) { clear_ptxinfo(); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index f7bf70e..9bb5008 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -46,6 +46,8 @@ bool g_override_embedded_ptx = false; extern int ptx_parse(); extern int ptx__scan_string(const char*); +extern std::map get_duplicate(); + const char *g_ptxinfo_filename; extern int ptxinfo_parse(); extern int ptxinfo_debug; @@ -182,6 +184,109 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source return symtab; } +void fix_duplicate_errors(char fname2[1024]) { + char tempfile[1024] = "_temp_ptx"; + char commandline[1024]; + + // change the name of the ptx file to _temp_ptx + snprintf(commandline,1024,"mv %s %s",fname2,tempfile); + printf("Running: %s\n", commandline); + int result = system(commandline); + if (result != 0) { + printf("GPGPU-Sim PTX: ERROR ** while changing filename from %s to %s", fname2, tempfile); + exit(1); + } + + // store all of the ptx into a char array + FILE *ptxsource = fopen(tempfile,"r"); + fseek(ptxsource, 0, SEEK_END); + long filesize = ftell(ptxsource); + rewind(ptxsource); + char *ptxdata = (char*)malloc((filesize+1)*sizeof(char)); + fread(ptxdata, filesize, 1, ptxsource); + fclose(ptxsource); + + FILE *ptxdest = fopen(fname2,"w"); + std::map duplicate = get_duplicate(); + unsigned offset; + unsigned oldlinenum = 1; + unsigned linenum; + char *startptr = ptxdata; + char *funcptr; + char *tempptr = ptxdata - 1; + char *lineptr = ptxdata - 1; + + // recreate the ptx file without duplications + for ( std::map::iterator iter = duplicate.begin(); + iter != duplicate.end(); + iter++){ + // find the line of the next error + linenum = iter->first; + for (int i = oldlinenum; i < linenum; i++) { + lineptr = strchr(lineptr + 1, '\n'); + } + + // find the end of the current section to be copied over + // then find the start of the next section that will be copied + if (strcmp("function", iter->second) == 0) { + // get location of most recent .func + while (tempptr < lineptr && tempptr != NULL) { + funcptr = tempptr; + tempptr = strstr(funcptr + 1, ".func"); + } + + // get the start of the previous line + offset = 0; + while (*(funcptr - offset) != '\n') offset++; + + fwrite(startptr, sizeof(char), funcptr - offset + 1 - startptr, ptxdest); + + //find next location of startptr + if (*(lineptr + 3) == ';') { + // for function definitions + startptr = lineptr + 5; + } else if (*(lineptr + 3) == '{') { + // for functions enclosed with curly brackets + offset = 5; + unsigned bracket = 1; + while (bracket != 0) { + if (*(lineptr + offset) == '{') bracket++; + else if (*(lineptr + offset) == '}') bracket--; + offset++; + } + startptr = lineptr + offset + 1; + } else { + printf("GPGPU-Sim PTX: ERROR ** Unrecognized function format\n"); + abort(); + } + } else if (strcmp("variable", iter->second) == 0) { + fwrite(startptr, sizeof(char), (int)(lineptr + 1 - startptr), ptxdest); + + //find next location of startptr + offset = 1; + while (*(lineptr + offset) != '\n') offset++; + startptr = lineptr + offset + 1; + } else { + printf("GPGPU-Sim PTX: ERROR ** Unsupported duplicate type: %s\n", iter->second); + } + + oldlinenum = linenum; + } + // copy over the rest of the file + fwrite(startptr, sizeof(char), ptxdata + filesize - startptr, ptxdest); + + // cleanup + free(ptxdata); + fclose(ptxdest); + snprintf(commandline,1024,"rm -f %s",tempfile); + printf("Running: %s\n", commandline); + result = system(commandline); + if (result != 0) { + printf("GPGPU-Sim PTX: ERROR ** while deleting %s", tempfile); + exit(1); + } +} + void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version ) { char fname[1024]; @@ -225,14 +330,29 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline); result = system(commandline); if( result != 0 ) { - printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result); - printf(" Ensure ptxas is in your path.\n"); - exit(1); + // 65280 = duplicate errors + if (result == 65280) { + ptxinfo_in = fopen(tempfile_ptxinfo,"r"); + g_ptxinfo_filename = tempfile_ptxinfo; + ptxinfo_parse(); + + fix_duplicate_errors(fname2); + snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file /dev/null 2> %s", + extra_flags, fname2, tempfile_ptxinfo); + printf("GPGPU-Sim PTX: regenerating ptxinfo using \"%s\"\n", commandline); + result = system(commandline); + } + if (result != 0) { + printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result); + printf(" Ensure ptxas is in your path.\n"); + exit(1); + } } ptxinfo_in = fopen(tempfile_ptxinfo,"r"); g_ptxinfo_filename = tempfile_ptxinfo; ptxinfo_parse(); + if( ! g_save_embedded_ptx ) { snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo); printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline); diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l index f9b6846..33c2748 100644 --- a/src/cuda-sim/ptxinfo.l +++ b/src/cuda-sim/ptxinfo.l @@ -60,6 +60,10 @@ unsigned ptxinfo_col = 0; "line" TC; return LINE; "for" TC; return FOR; "textures" TC; return TEXTURES; +"error : Duplicate definition of" TC; return DUPLICATE; +"function" TC; ptxinfo_lval.string_value = strdup(yytext); return FUNCTION; +"variable" TC; ptxinfo_lval.string_value = strdup(yytext); return VARIABLE; +"fatal : Ptx assembly aborted due to errors" TC; return FATAL; [_A-Za-z$%][_0-9A-Za-z$]* TC; ptxinfo_lval.string_value = strdup(yytext); return IDENTIFIER; [-]{0,1}[0-9]+ TC; ptxinfo_lval.int_value = atoi(yytext); return INT_OPERAND; diff --git a/src/cuda-sim/ptxinfo.y b/src/cuda-sim/ptxinfo.y index faa33eb..37092f4 100644 --- a/src/cuda-sim/ptxinfo.y +++ b/src/cuda-sim/ptxinfo.y @@ -55,11 +55,15 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %token WARNING %token FOR %token TEXTURES +%token DUPLICATE +%token FUNCTION +%token VARIABLE +%token FATAL %{ #include #include - + static unsigned g_declared; static unsigned g_system; int ptxinfo_lex(void); @@ -71,6 +75,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. void ptxinfo_smem( unsigned declared, unsigned system ); void ptxinfo_cmem( unsigned nbytes, unsigned bank ); int ptxinfo_error(const char*); + void ptxinfo_linenum( unsigned ); + void ptxinfo_dup_type( const char* ); %} %% @@ -82,6 +88,8 @@ input: /* empty */ line: HEADER INFO COLON line_info | HEADER IDENTIFIER COMMA LINE INT_OPERAND SEMICOLON WARNING | HEADER WARNING { printf("GPGPU-Sim: ptxas %s\n", $2); } + | HEADER IDENTIFIER COMMA LINE INT_OPERAND SEMICOLON DUPLICATE duplicate { ptxinfo_linenum($5); } + | HEADER FATAL ; line_info: function_name @@ -89,7 +97,7 @@ line_info: function_name ; function_name: FUNC QUOTE IDENTIFIER QUOTE { ptxinfo_function($3); } - | FUNC QUOTE IDENTIFIER QUOTE FOR QUOTE IDENTIFIER QUOTE {ptxinfo_function($3); } + | FUNC QUOTE IDENTIFIER QUOTE FOR QUOTE IDENTIFIER QUOTE { ptxinfo_function($3); } ; function_info: info @@ -110,6 +118,10 @@ info: USED INT_OPERAND REGS { ptxinfo_regs($2); } tuple: INT_OPERAND PLUS INT_OPERAND BYTES { g_declared=$1; g_system=$3; } +duplicate: FUNCTION QUOTE IDENTIFIER QUOTE { ptxinfo_dup_type($1); } + | VARIABLE QUOTE IDENTIFIER QUOTE { ptxinfo_dup_type($1); } + ; + %% -- cgit v1.3 From 2683b8bd7ba9950e0aa174915ef9ff64e0a20421 Mon Sep 17 00:00:00 2001 From: speverel Date: Mon, 4 Jul 2016 15:07:50 -0700 Subject: Restored madp instruction. --- src/cuda-sim/cuda-sim.cc | 2 +- src/cuda-sim/instructions.cc | 5 +++++ src/cuda-sim/opcodes.def | 1 + src/cuda-sim/ptx.l | 1 + 4 files changed, 8 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index fba3a59..09e9a81 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -719,7 +719,7 @@ void ptx_instruction::set_opcode_and_latency() break; } break; - case MAD_OP: case MADC_OP: + case MAD_OP: case MADC_OP: case MADP_OP: //MAD latency switch(get_type()){ case F32_TYPE: diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 02ce01c..7b0f4fa 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -2496,6 +2496,11 @@ void mad_impl( const ptx_instruction *pI, ptx_thread_info *thread ) mad_def(pI, thread, false); } +void madp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + mad_def(pI, thread, true); +} + void madc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { mad_def(pI, thread, true); diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 874acc7..2ee6976 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -70,6 +70,7 @@ OP_DEF(LG2_OP,lg2_impl,"lg2",1,4) OP_DEF(MAD24_OP,mad24_impl,"mad24",1,2) OP_DEF(MAD_OP,mad_impl,"mad",1,2) OP_DEF(MADC_OP,madc_impl,"madc",1,2) +OP_DEF(MADP_OP,madp_impl,"madp",1,2) OP_DEF(MAX_OP,max_impl,"max",1,1) OP_DEF(MEMBAR_OP,membar_impl,"membar",1,3) OP_DEF(MIN_OP,min_impl,"min",1,1) diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index b8ce497..88ccf6a 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -87,6 +87,7 @@ lg2 TC; ptx_lval.int_value = LG2_OP; return OPCODE; mad24 TC; ptx_lval.int_value = MAD24_OP; return OPCODE; mad TC; ptx_lval.int_value = MAD_OP; return OPCODE; madc TC; ptx_lval.int_value = MADC_OP; return OPCODE; +madp TC; ptx_lval.int_value = MADP_OP; return OPCODE; max TC; ptx_lval.int_value = MAX_OP; return OPCODE; membar TC; ptx_lval.int_value = MEMBAR_OP; return OPCODE; min TC; ptx_lval.int_value = MIN_OP; return OPCODE; -- cgit v1.3 From 70e02ee5283cb96f0edcb46a15edf0ab6e1d0697 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 13 Oct 2014 18:58:37 -0400 Subject: ADD: add cudaGetParameterBufferV2 and add cudaLaunchDeviceV2 implementation. Kernel launch to stream not yet implemented --- src/cuda-sim/Makefile | 3 +- src/cuda-sim/cuda-sim.cc | 10 +++ src/cuda-sim/cuda_device_runtime.cc | 175 ++++++++++++++++++++++++++++++++++++ src/cuda-sim/cuda_device_runtime.h | 7 ++ src/cuda-sim/instructions.cc | 12 ++- src/cuda-sim/ptx_ir.h | 4 +- src/cuda-sim/ptx_parser.cc | 14 +++ 7 files changed, 219 insertions(+), 6 deletions(-) create mode 100644 src/cuda-sim/cuda_device_runtime.cc create mode 100644 src/cuda-sim/cuda_device_runtime.h (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile index 166e256..f479294 100644 --- a/src/cuda-sim/Makefile +++ b/src/cuda-sim/Makefile @@ -62,7 +62,7 @@ ifeq ($(GNUC_CPP0X),1) endif endif -OBJS := $(OUTPUT_DIR)/ptx_parser.o $(OUTPUT_DIR)/ptx_loader.o $(OUTPUT_DIR)/cuda_device_printf.o $(OUTPUT_DIR)/instructions.o $(OUTPUT_DIR)/cuda-sim.o $(OUTPUT_DIR)/ptx_ir.o $(OUTPUT_DIR)/ptx_sim.o $(OUTPUT_DIR)/memory.o $(OUTPUT_DIR)/ptx-stats.o $(OUTPUT_DIR)/decuda_pred_table/decuda_pred_table.o $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o +OBJS := $(OUTPUT_DIR)/ptx_parser.o $(OUTPUT_DIR)/ptx_loader.o $(OUTPUT_DIR)/cuda_device_printf.o $(OUTPUT_DIR)/instructions.o $(OUTPUT_DIR)/cuda-sim.o $(OUTPUT_DIR)/ptx_ir.o $(OUTPUT_DIR)/ptx_sim.o $(OUTPUT_DIR)/memory.o $(OUTPUT_DIR)/ptx-stats.o $(OUTPUT_DIR)/decuda_pred_table/decuda_pred_table.o $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o $(OUTPUT_DIR)/cuda_device_runtime.o OPT += -DCUDART_VERSION=$(CUDART_VERSION) @@ -145,5 +145,6 @@ $(OUTPUT_DIR)/ptx_sim.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/cuda-sim.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/lex.ptxinfo_.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/lex.ptx_.o: $(OUTPUT_DIR)/ptx.tab.c +$(OUTPUT_DIR)/cuda_device_runtime.o: $(OUTPUT_DIR)/ptx.tab.c include $(OUTPUT_DIR)/Makefile.makedepend diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 715be98..4933029 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1065,6 +1065,16 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg } } +unsigned function_info::get_args_aligned_size() { + unsigned int align_size = 4; // a word + unsigned int total_size = 0; + for(unsigned int i = 0; i < num_args(); i++) { + total_size += ((m_args[i]->get_size_in_bytes() + align_size - 1) / align_size) * align_size; + } + return total_size; +} + + void function_info::finalize( memory_space *param_mem ) { unsigned param_address = 0; diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc new file mode 100644 index 0000000..937eec8 --- /dev/null +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -0,0 +1,175 @@ +//Jin: cuda_device_runtime.cc +//Defines CUDA device runtime APIs for CDP support + +#include +#include + +#define __CUDA_RUNTIME_API_H__ + +#include +#include +#include "../gpgpu-sim/gpu-sim.h" +#include "cuda-sim.h" +#include "ptx_ir.h" +#include "../stream_manager.h" +#include "cuda_device_runtime.h" + +#define DEV_RUNTIME_REPORT(a) \ + if( g_debug_execution ) { \ + std::cout << __FILE__ << ", " << __LINE__ << ": " << a << "\n"; \ + std::cout.flush(); \ + } + +std::map g_cuda_device_launch_map; +struct CUstream_st * g_device_default_stream = NULL; +extern stream_manager *g_stream_manager; + +//Handling device runtime api: +//void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) +void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) +{ + DEV_RUNTIME_REPORT("Calling cudaGetParameterBufferV2"); + + unsigned n_return = target_func->has_return(); + assert(n_return); + unsigned n_args = target_func->num_args(); + assert( n_args == 4 ); + + function_info * child_kernel_entry; + struct dim3 gridDim, blockDim; + unsigned int sharedMem; + + for( unsigned arg=0; arg < n_args; arg ++ ) { + const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# + const symbol *formal_param = target_func->get_arg(arg); //cudaGetParameterBufferV2_param_# + unsigned size=formal_param->get_size_in_bytes(); + assert( formal_param->is_param_local() ); + assert( actual_param_op.is_param_local() ); + addr_t from_addr = actual_param_op.get_symbol()->get_address(); + + if(arg == 0) {//function_info* for the child kernel + unsigned long long buf; + assert(size == sizeof(function_info *)); + thread->m_local_mem->read(from_addr, size, &buf); + child_kernel_entry = (function_info *)buf; + assert(child_kernel_entry); + DEV_RUNTIME_REPORT("child kernel name " << child_kernel_entry->get_name()); + } + else if(arg == 1) { //dim3 gridDim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & gridDim); + DEV_RUNTIME_REPORT("grid (" << gridDim.x << ", " << gridDim.y << ", " << gridDim.z << ")"); + } + else if(arg == 2) { //dim3 blockDim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & blockDim); + DEV_RUNTIME_REPORT("block (" << blockDim.x << ", " << blockDim.y << ", " << blockDim.z << ")"); + } + else if(arg == 3) { //unsigned int sharedMem + assert(size == sizeof(unsigned int)); + thread->m_local_mem->read(from_addr, size, & sharedMem); + DEV_RUNTIME_REPORT("shared memory " << sharedMem); + } + } + + //get total child kernel argument size and malloc buffer in global memory + unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); + void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); + DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); + + //create child kernel_info_t and index it with parameter_buffer address + kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); + assert(g_cuda_device_launch_map.find(param_buffer) == g_cuda_device_launch_map.end()); + g_cuda_device_launch_map[param_buffer] = child_grid; + + //copy the buffer address to retval0 + const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 + const symbol *formal_return = target_func->get_return_var(); //void * + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *)); + addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); + thread->m_local_mem->write(ret_param_addr, return_size, ¶m_buffer, NULL, NULL); + +} + +//Handling device runtime api: +//cudaError_t cudaLaunchDeviceV2(void *parameterBuffer, cudaStream_t stream) +void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) { + DEV_RUNTIME_REPORT("Calling cudaLaunchDeviceV2"); + + unsigned n_return = target_func->has_return(); + assert(n_return); + unsigned n_args = target_func->num_args(); + assert( n_args == 2 ); + + kernel_info_t * child_grid = NULL; + function_info * child_kernel_entry = NULL; + void * parameter_buffer; + struct CUstream_st * child_stream; + for( unsigned arg=0; arg < n_args; arg ++ ) { + const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# + const symbol *formal_param = target_func->get_arg(arg); //cudaLaunchDeviceV2_param_# + unsigned size=formal_param->get_size_in_bytes(); + assert( formal_param->is_param_local() ); + assert( actual_param_op.is_param_local() ); + addr_t from_addr = actual_param_op.get_symbol()->get_address(); + + if(arg == 0) {//paramter buffer for child kernel (in global memory) + //get parameter_buffer from the cudaDeviceLaunchV2_param0 + assert(size == sizeof(void *)); + thread->m_local_mem->read(from_addr, size, ¶meter_buffer); + assert((size_t)parameter_buffer >= GLOBAL_HEAP_START); + DEV_RUNTIME_REPORT("Parameter buffer locating at global memory " << parameter_buffer); + + //get child grid info through parameter_buffer address + assert(g_cuda_device_launch_map.find(parameter_buffer) != g_cuda_device_launch_map.end()); + child_grid = g_cuda_device_launch_map[parameter_buffer]; + child_kernel_entry = child_grid->entry(); + DEV_RUNTIME_REPORT("find child kernel " << child_kernel_entry->get_name()); + + //copy data in parameter_buffer to child kernel param memory + unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); + DEV_RUNTIME_REPORT("child_kernel_arg_size " << child_kernel_arg_size); + memory_space *child_kernel_param_mem = child_grid->get_param_memory(); + size_t param_start_address = 0; + for(unsigned n = 0; n < child_kernel_arg_size; n++) { + unsigned char one_byte; + thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte); + child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); + } + } + else if(arg == 1) { //cudaStream for the child kernel + assert(size == sizeof(cudaStream_t)); + thread->m_local_mem->read(from_addr, size, &child_stream); + if(child_stream == 0) { //default stream on device + if(!g_device_default_stream) { + //g_device_default_stream = new struct CUstream_st(); + //g_stream_manager->add_stream(g_device_default_stream); + } + child_stream = g_device_default_stream; + } +// DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid()); + } + + } + + //launch child kernel + stream_operation op(child_grid, g_ptx_sim_mode, child_stream); +// g_stream_manager->push(op); +// g_cuda_device_launch_map.erase(parameter_buffer); + + //set retval0 + const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 + const symbol *formal_return = target_func->get_return_var(); //cudaError_t + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size + && return_size == sizeof(cudaError_t)); + cudaError_t error = cudaSuccess; + addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); + +} diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h new file mode 100644 index 0000000..1b10407 --- /dev/null +++ b/src/cuda-sim/cuda_device_runtime.h @@ -0,0 +1,7 @@ +//Jin: cuda_device_runtime.h +//Defines CUDA device runtime APIs for CDP support + +#pragma once + +void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); +void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 254427b..1c47ad3 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -41,6 +41,9 @@ #include "../gpgpu-sim/gpu-sim.h" #include "../gpgpu-sim/shader.h" +//Jin: include device runtime for CDP +#include "cuda_device_runtime.h" + #include unsigned ptx_instruction::g_num_ptx_inst_uid=0; @@ -149,8 +152,7 @@ ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_in } else if ( op.is_local() ) { result.u64 = op.get_symbol()->get_address(); } else if ( op.is_function_address() ) { - result.u64 = op.get_symbol()->get_pc()->get_start_PC(); - printf("Get pc for kernel function %u\n", op.get_symbol()->get_pc()->get_start_PC()); + result.u64 = (size_t)op.get_symbol()->get_pc(); } else { const char *name = op.name().c_str(); printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name ); @@ -1411,12 +1413,14 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpusim_cuda_vprintf(pI, thread, target_func); return; } + + //Jin: handle device runtime apis for CDP else if(fname == "cudaGetParameterBufferV2") { - printf("calling cudaGetParameterBufferV2\n"); + gpgpusim_cuda_getParameterBufferV2(pI, thread, target_func); return; } else if(fname == "cudaLaunchDeviceV2") { - printf("calling cudaLaunchDeviceV2\n"); + gpgpusim_cuda_launchDeviceV2(pI, thread, target_func); return; } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 7325e5f..a7ca27e 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -699,7 +699,7 @@ public: } bool is_immediate_address() const { - return m_immediate_address; + return m_immediate_address; } bool is_literal() const { return m_type == int_t || @@ -1209,6 +1209,8 @@ public: { return m_args.size(); } + unsigned get_args_aligned_size(); + const symbol* get_arg( unsigned n ) const { assert( n < m_args.size() ); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 39257da..baa3bcd 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -120,6 +120,20 @@ symbol_table *init_parser( const char *ptx_filename ) #define DEF(X,Y) g_ptx_token_decode[X] = Y; #include "ptx_parser_decode.def" #undef DEF + g_ptx_token_decode[undefined_space] = "undefined_space"; + g_ptx_token_decode[undefined_space] = "undefined_space=0"; + g_ptx_token_decode[reg_space] = "reg_space"; + g_ptx_token_decode[local_space] = "local_space"; + g_ptx_token_decode[shared_space] = "shared_space"; + g_ptx_token_decode[param_space_unclassified] = "param_space_unclassified"; + g_ptx_token_decode[param_space_kernel] = "param_space_kernel"; + g_ptx_token_decode[param_space_local] = "param_space_local"; + g_ptx_token_decode[const_space] = "const_space"; + g_ptx_token_decode[tex_space] = "tex_space"; + g_ptx_token_decode[surf_space] = "surf_space"; + g_ptx_token_decode[global_space] = "global_space"; + g_ptx_token_decode[generic_space] = "generic_space"; + g_ptx_token_decode[instruction_space] = "instruction_space"; return g_global_symbol_table; } -- cgit v1.3 From 582106ec595bb7745db092f2f4aa2b0fb9521b16 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sat, 18 Oct 2014 23:07:53 -0400 Subject: MOD: add child kernel stream and scheduling support --- src/abstract_hardware_model.cc | 77 +++++++++++++++++++++++++++++++++++++ src/abstract_hardware_model.h | 39 +++++++++++++++++++ src/cuda-sim/cuda-sim.cc | 9 ++++- src/cuda-sim/cuda_device_runtime.cc | 31 +++++++++------ src/cuda-sim/ptx_sim.h | 3 ++ src/gpgpu-sim/gpu-sim.cc | 4 ++ src/gpgpu-sim/gpu-sim.h | 20 ++++++++++ src/gpgpusim_entrypoint.cc | 9 +++++ src/stream_manager.cc | 39 ++++++++++--------- 9 files changed, 199 insertions(+), 32 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 84d165c..b19b9d1 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -565,6 +565,10 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_num_cores_running=0; m_uid = m_next_uid++; m_param_mem = new memory_space_impl<8192>("param",64*1024); + + //Jin: parent and child kernel management for CDP + m_parent_kernel = NULL; + } kernel_info_t::~kernel_info_t() @@ -578,6 +582,79 @@ std::string kernel_info_t::name() const return m_kernel_entry->get_name(); } +//Jin: parent and child kernel management for CDP +void kernel_info_t::set_parent(kernel_info_t * parent, + dim3 parent_ctaid, dim3 parent_tid) { + m_parent_kernel = parent; + m_parent_ctaid = parent_ctaid; + m_parent_tid = parent_tid; + parent->set_child(this); +} + +void kernel_info_t::set_child(kernel_info_t * child) { + m_child_kernels.push_back(child); +} + +bool kernel_info_t::is_finished() { + if(done() && children_all_finished()) + return true; + else + return false; +} + +bool kernel_info_t::children_all_finished() { + for(auto child = m_child_kernels.begin(); child != m_child_kernels.end(); child++) + { + if(!(*child)->is_finished()) + return false; + } + return true; +} + +void kernel_info_t::notify_parent_finished() { + if(m_parent_kernel) { + g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); + } +} + +CUstream_st * kernel_info_t::create_stream_cta(dim3 ctaid) { + assert(get_default_stream_cta(ctaid)); + CUstream_st * stream = new CUstream_st(); + g_stream_manager->add_stream(stream); + assert(m_cta_streams.find(ctaid) != m_cta_streams.end()); + assert(m_cta_streams[ctaid].size() >= 1); //must have default stream + m_cta_streams[ctaid].push_back(stream); + + return stream; +} + +CUstream_st * kernel_info_t::get_default_stream_cta(dim3 ctaid) { + if(m_cta_streams.find(ctaid) != m_cta_streams.end()) { + assert(m_cta_streams[ctaid].size() >= 1); //already created, must have default stream + return *(m_cta_streams[ctaid].begin()); + } + else { + m_cta_streams[ctaid] = std::list(); + CUstream_st * stream = new CUstream_st(); + g_stream_manager->add_stream(stream); + m_cta_streams[ctaid].push_back(stream); + return stream; + } +} + +bool kernel_info_t::cta_has_stream(dim3 ctaid, CUstream_st* stream) { + if(m_cta_streams.find(ctaid) == m_cta_streams.end()) + return false; + + std::list &stream_list = m_cta_streams[ctaid]; + if(std::find(stream_list.begin(), stream_list.end(), stream) + == stream_list.end()) + return false; + else + return true; +} + + simt_stack::simt_stack( unsigned wid, unsigned warpSize) { m_warp_id=wid; diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index ba4ea29..16f4b31 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -154,15 +154,35 @@ enum _memory_op_t { #include #include #include +#include #if !defined(__VECTOR_TYPES_H__) struct dim3 { unsigned int x, y, z; }; #endif +struct dim3comp { + bool operator() (const dim3 & a, const dim3 & b) const + { + if(a.z < b.z) + return true; + else if(a.y < b.y) + return true; + else if (a.x < b.x) + return true; + else + return false; + } +}; void increment_x_then_y_then_z( dim3 &i, const dim3 &bound); +//Jin: child kernel information for CDP +#include "stream_manager.h" +class stream_manager; +struct CUstream_st; +extern stream_manager * g_stream_manager; + class kernel_info_t { public: // kernel_info_t() @@ -250,6 +270,25 @@ private: std::list m_active_threads; class memory_space *m_param_mem; + +public: + //Jin: parent and child kernel management for CDP + void set_parent(kernel_info_t * parent, dim3 parent_ctaid, dim3 parent_tid); + void set_child(kernel_info_t * child); + bool is_finished(); + bool children_all_finished(); + void notify_parent_finished(); + CUstream_st * create_stream_cta(dim3 ctaid); + CUstream_st * get_default_stream_cta(dim3 ctaid); + bool cta_has_stream(dim3 ctaid, CUstream_st* stream); + +private: + kernel_info_t * m_parent_kernel; + dim3 m_parent_ctaid; + dim3 m_parent_tid; + std::list m_child_kernels; //child kernel launched + std::map< dim3, std::list, dim3comp > m_cta_streams; //streams created in each CTA + }; struct core_config { diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 4933029..482e3f6 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1762,8 +1762,13 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) extern stream_manager *g_stream_manager; //openCL kernel simulation calls don't register the kernel so we don't register its exit - if(!openCL) - g_stream_manager->register_finished_kernel(kernel.get_uid()); + if(!openCL) { + + //Jin: for CDP, children should be finished first + if(kernel.is_finished()) { + g_stream_manager->register_finished_kernel(kernel.get_uid()); + } + } //******PRINTING******* printf( "GPGPU-Sim: Done functional simulation (%u instructions simulated).\n", g_ptx_sim_num_insn ); diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 937eec8..4d8d1e1 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -21,8 +21,7 @@ } std::map g_cuda_device_launch_map; -struct CUstream_st * g_device_default_stream = NULL; -extern stream_manager *g_stream_manager; +extern stream_manager * g_stream_manager; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -78,7 +77,13 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); //create child kernel_info_t and index it with parameter_buffer address - kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); + kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); + kernel_info_t & parent_grid = thread->get_kernel(); + DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << + thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << + "), thread (" << thread->get_tid().x << ", " << thread->get_tid().y << ", " << thread->get_tid().z << + ")"); + child_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid()); assert(g_cuda_device_launch_map.find(param_buffer) == g_cuda_device_launch_map.end()); g_cuda_device_launch_map[param_buffer] = child_grid; @@ -137,28 +142,30 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * for(unsigned n = 0; n < child_kernel_arg_size; n++) { unsigned char one_byte; thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte); + std::cout << "one byte " << std::hex << one_byte << "\n"; child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); } } else if(arg == 1) { //cudaStream for the child kernel assert(size == sizeof(cudaStream_t)); thread->m_local_mem->read(from_addr, size, &child_stream); - if(child_stream == 0) { //default stream on device - if(!g_device_default_stream) { - //g_device_default_stream = new struct CUstream_st(); - //g_stream_manager->add_stream(g_device_default_stream); - } - child_stream = g_device_default_stream; + + kernel_info_t & parent_kernel = thread->get_kernel(); + if(child_stream == 0) { //default stream on device for current CTA + child_stream = parent_kernel.get_default_stream_cta(thread->get_ctaid()); + } + else { + assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); } -// DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid()); + DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << " " << child_stream); } } //launch child kernel stream_operation op(child_grid, g_ptx_sim_mode, child_stream); -// g_stream_manager->push(op); -// g_cuda_device_launch_map.erase(parameter_buffer); + g_stream_manager->push(op); + g_cuda_device_launch_map.erase(parameter_buffer); //set retval0 const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h index f926e6d..ea171c5 100644 --- a/src/cuda-sim/ptx_sim.h +++ b/src/cuda-sim/ptx_sim.h @@ -418,6 +418,9 @@ public: void or_reduction(unsigned ctaid, unsigned barid, bool value) {m_core->or_reduction(ctaid,barid,value);} void popc_reduction(unsigned ctaid, unsigned barid, bool value) {m_core->popc_reduction(ctaid,barid,value);} + //Jin: get corresponding kernel grid for CDP purpose + kernel_info_t & get_kernel() { return m_kernel; } + public: addr_t m_last_effective_address; bool m_branch_taken; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index eafb909..47dbc89 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -622,6 +622,10 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) *active_sms=0; last_liveness_message_time = 0; + + //Jin: functional simulation for CDP + m_functional_sim = false; + m_functional_sim_kernel = NULL; } int gpgpu_sim::shared_mem_size() const diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 33fffd3..a2d1b9b 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -492,6 +492,7 @@ private: std::string executed_kernel_info_string(); //< format the kernel information into a string for stat printout void clear_executed_kernel_info(); //< clear the kernel information after stat printout + public: unsigned long long gpu_sim_insn; unsigned long long gpu_tot_sim_insn; @@ -504,6 +505,25 @@ public: void change_cache_config(FuncCache cache_config); void set_cache_config(std::string kernel_name); + //Jin: functional simulation for CDP +private: + //set by stream operation every time a functoinal simulation is done + bool m_functional_sim; + kernel_info_t * m_functional_sim_kernel; + +public: + bool is_functional_sim() { return m_functional_sim; } + kernel_info_t * get_functional_kernel() { return m_functional_sim_kernel; } + void functional_launch(kernel_info_t * k) { + m_functional_sim = true; + m_functional_sim_kernel = k; + } + void finish_functional_sim(kernel_info_t * k) { + assert(m_functional_sim); + assert(m_functional_sim_kernel == k); + m_functional_sim = false; + m_functional_sim_kernel = NULL; + } }; #endif diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 6ba38eb..31f3a41 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -127,6 +127,15 @@ void *gpgpu_sim_thread_concurrent(void*) if(g_stream_manager->operation(&sim_cycles) && !g_the_gpu->active()) break; + //functional simulation + if( g_the_gpu->is_functional_sim()) { + kernel_info_t * kernel = g_the_gpu->get_functional_kernel(); + assert(kernel); + gpgpu_cuda_ptx_sim_main_func(*kernel); + g_the_gpu->finish_functional_sim(kernel); + } + + //performance simulation if( g_the_gpu->active() ) { g_the_gpu->cycle(); sim_cycles = true; diff --git a/src/stream_manager.cc b/src/stream_manager.cc index dd42f0a..1f93c03 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -155,7 +155,7 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) gpu->set_cache_config(m_kernel->name()); printf("kernel \'%s\' transfer to GPU hardware scheduler\n", m_kernel->name().c_str() ); if( m_sim_mode ) - gpgpu_cuda_ptx_sim_main_func( *m_kernel ); + gpu->functional_launch( m_kernel ); else gpu->launch( m_kernel ); } @@ -212,11 +212,9 @@ bool stream_manager::operation( bool * sim) bool stream_manager::check_finished_kernel() { - - unsigned grid_uid = m_gpu->finished_kernel(); - bool check=register_finished_kernel(grid_uid); - return check; - + unsigned grid_uid = m_gpu->finished_kernel(); + bool check=register_finished_kernel(grid_uid); + return check; } bool stream_manager::register_finished_kernel(unsigned grid_uid) @@ -226,13 +224,17 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) CUstream_st *stream = m_grid_id_to_stream[grid_uid]; kernel_info_t *kernel = stream->front().get_kernel(); assert( grid_uid == kernel->get_uid() ); - stream->record_next_done(); - m_grid_id_to_stream.erase(grid_uid); - delete kernel; - return true; - }else{ - return false; + + //Jin: should check children kernels for CDP + if(kernel->is_finished()) { + stream->record_next_done(); + m_grid_id_to_stream.erase(grid_uid); + kernel->notify_parent_finished(); + delete kernel; + return true; + } } + return false; } @@ -259,21 +261,22 @@ stream_operation stream_manager::front() { // called by gpu simulation thread stream_operation result; - if( concurrent_streams_empty() ) - m_service_stream_zero = true; +// if( concurrent_streams_empty() ) + m_service_stream_zero = true; if( m_service_stream_zero ) { - if( !m_stream_zero.empty() ) { - if( !m_stream_zero.busy() ) { + if( !m_stream_zero.empty() && !m_stream_zero.busy() ) { result = m_stream_zero.next(); if( result.is_kernel() ) { unsigned grid_id = result.get_kernel()->get_uid(); m_grid_id_to_stream[grid_id] = &m_stream_zero; } - } } else { m_service_stream_zero = false; } - } else { + } + + if(!m_service_stream_zero) + { std::list::iterator s; for( s=m_streams.begin(); s != m_streams.end(); s++) { CUstream_st *stream = *s; -- cgit v1.3 From d31c2da2b34d3d7a63d9980335c85d5e1a19ad02 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sun, 19 Oct 2014 01:02:20 -0400 Subject: BUG: parameter alignment --- src/cuda-sim/cuda-sim.cc | 26 +++++++++++++++++++++----- src/cuda-sim/cuda_device_runtime.cc | 1 - 2 files changed, 21 insertions(+), 6 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 482e3f6..4281913 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1066,12 +1066,24 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg } unsigned function_info::get_args_aligned_size() { - unsigned int align_size = 4; // a word + + unsigned param_address = 0; unsigned int total_size = 0; - for(unsigned int i = 0; i < num_args(); i++) { - total_size += ((m_args[i]->get_size_in_bytes() + align_size - 1) / align_size) * align_size; + 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(); + symbol *param = m_symtab->lookup(name.c_str()); + + size_t arg_size = p.get_size(); // size of param in bytes + total_size = (total_size + arg_size - 1) / arg_size * arg_size; //aligned + p.add_offset(total_size); + param_address += total_size; + param->set_address(param_address); + total_size += arg_size; } - return total_size; + + return (total_size + 3) / 4; //final size aligned to word + } @@ -1097,13 +1109,17 @@ void function_info::finalize( memory_space *param_mem ) size = (size<(p.get_size()/8))?size:(p.get_size()/8); } // 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 size_t word_size = 4; + 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(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); param->set_address(param_address); - param_address += size; + param_address += size; } } diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 4d8d1e1..148471b 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -142,7 +142,6 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * for(unsigned n = 0; n < child_kernel_arg_size; n++) { unsigned char one_byte; thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte); - std::cout << "one byte " << std::hex << one_byte << "\n"; child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); } } -- cgit v1.3 From 5fc7bf5872aad126a09cad4b385054c4b3a094aa Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 20 Oct 2014 17:19:32 -0400 Subject: BUG: do not handle cudaGetParameterBufferV2 and cudaLaunchDeviceV2 as call.uni in reconvergence --- src/abstract_hardware_model.h | 2 ++ src/cuda-sim/cuda-sim.cc | 6 ++++-- src/cuda-sim/ptx_ir.cc | 3 +++ 3 files changed, 9 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 16f4b31..11fee10 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -864,6 +864,7 @@ public: m_mem_accesses_created=false; m_cache_hit=false; m_is_printf=false; + m_is_cdp = false; } virtual ~warp_inst_t(){ } @@ -1016,6 +1017,7 @@ protected: unsigned cycles; // used for implementing initiation interval delay bool m_isatomic; bool m_is_printf; + bool m_is_cdp; unsigned m_warp_id; unsigned m_dynamic_warp_id; const core_config *m_config; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 4281913..58bd4e0 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -638,7 +638,7 @@ void ptx_instruction::set_opcode_and_latency() case MEMBAR_OP: op = MEMORY_BARRIER_OP; break; case CALL_OP: { - if(m_is_printf) + if(m_is_printf || m_is_cdp) op = ALU_OP; else op = CALL_OPS; @@ -646,7 +646,7 @@ void ptx_instruction::set_opcode_and_latency() } case CALLP_OP: { - if(m_is_printf) + if(m_is_printf || m_is_cdp) op = ALU_OP; else op = CALL_OPS; @@ -1219,6 +1219,8 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) bool skip = false; int op_classification = 0; addr_t pc = next_instr(); + if(pc == 440) + pc = 440; 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() ); diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 915c623..861f0dc 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1241,6 +1241,9 @@ ptx_instruction::ptx_instruction( int opcode, if (fname =="vprintf"){ m_is_printf = true; } + if (fname == "cudaGetParameterBufferV2" + || fname == "cudaLaunchDeviceV2") + m_is_cdp = true; } } -- cgit v1.3 From 6286547cfdc5d14c84568505db267f5e8dd9841f Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 20 Oct 2014 23:23:00 -0400 Subject: BUG: multiple child kernels finish --- src/abstract_hardware_model.cc | 22 ++++++++++++++++++---- src/abstract_hardware_model.h | 2 ++ src/cuda-sim/cuda-sim.cc | 6 +----- src/stream_manager.cc | 3 ++- 4 files changed, 23 insertions(+), 10 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index b19b9d1..f3c5b21 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -595,6 +595,12 @@ void kernel_info_t::set_child(kernel_info_t * child) { m_child_kernels.push_back(child); } +void kernel_info_t::remove_child(kernel_info_t * child) { + assert(std::find(m_child_kernels.begin(), m_child_kernels.end(), child) + != m_child_kernels.end()); + m_child_kernels.remove(child); +} + bool kernel_info_t::is_finished() { if(done() && children_all_finished()) return true; @@ -603,16 +609,15 @@ bool kernel_info_t::is_finished() { } bool kernel_info_t::children_all_finished() { - for(auto child = m_child_kernels.begin(); child != m_child_kernels.end(); child++) - { - if(!(*child)->is_finished()) + if(!m_child_kernels.empty()) return false; - } + return true; } void kernel_info_t::notify_parent_finished() { if(m_parent_kernel) { + m_parent_kernel->remove_child(this); g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); } } @@ -654,6 +659,15 @@ bool kernel_info_t::cta_has_stream(dim3 ctaid, CUstream_st* stream) { return true; } +void kernel_info_t::print_parent_info() { + if(m_parent_kernel) { + printf("Parent %d: \'%s\', Block (%d, %d, %d), Thread (%d, %d, %d)\n", + m_parent_kernel->get_uid(), m_parent_kernel->name().c_str(), + m_parent_ctaid.x, m_parent_ctaid.y, m_parent_ctaid.z, + m_parent_tid.x, m_parent_tid.y, m_parent_tid.z); + } +} + simt_stack::simt_stack( unsigned wid, unsigned warpSize) { diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 11fee10..a0bf80d 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -275,12 +275,14 @@ public: //Jin: parent and child kernel management for CDP void set_parent(kernel_info_t * parent, dim3 parent_ctaid, dim3 parent_tid); void set_child(kernel_info_t * child); + void remove_child(kernel_info_t * child); bool is_finished(); bool children_all_finished(); void notify_parent_finished(); CUstream_st * create_stream_cta(dim3 ctaid); CUstream_st * get_default_stream_cta(dim3 ctaid); bool cta_has_stream(dim3 ctaid, CUstream_st* stream); + void print_parent_info(); private: kernel_info_t * m_parent_kernel; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 58bd4e0..8dd1078 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1781,11 +1781,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) //openCL kernel simulation calls don't register the kernel so we don't register its exit if(!openCL) { - - //Jin: for CDP, children should be finished first - if(kernel.is_finished()) { - g_stream_manager->register_finished_kernel(kernel.get_uid()); - } + g_stream_manager->register_finished_kernel(kernel.get_uid()); } //******PRINTING******* diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 1f93c03..687d544 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -153,7 +153,8 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) case stream_kernel_launch: if( gpu->can_start_kernel() ) { gpu->set_cache_config(m_kernel->name()); - printf("kernel \'%s\' transfer to GPU hardware scheduler\n", m_kernel->name().c_str() ); + printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); + m_kernel->print_parent_info(); if( m_sim_mode ) gpu->functional_launch( m_kernel ); else -- cgit v1.3 From cf507ddb207337bc7a67ded7c4438f0cb0bed26f Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sun, 2 Nov 2014 19:24:37 -0500 Subject: BUG: PTX section id. ADD: cudaDeviceSetLimit. BUG: parameter addresses for child kernels in CDP. BUG: .weak .entry and .weak .global directives in ptx file. BUG: empty_protected() for stream manager causes deadlock, change to empty() --- libcuda/cuda_runtime_api.cc | 5 +++++ src/cuda-sim/cuda-sim.cc | 9 +++------ src/cuda-sim/cuda_device_runtime.cc | 9 +++++---- src/cuda-sim/ptx.y | 2 ++ src/gpgpusim_entrypoint.cc | 2 +- 5 files changed, 16 insertions(+), 11 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 5310a52..30bf823 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2018,6 +2018,11 @@ __host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfig(const char *func, enum cud context->get_device()->get_gpgpu()->set_cache_config(context->get_kernel(func)->get_name(), (FuncCache)cacheConfig); return g_last_cudaError = cudaSuccess; } + +//Jin: hack for cdp +__host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t value) { + return g_last_cudaError = cudaSuccess; +} #endif #endif diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 8dd1078..980afc8 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1074,15 +1074,14 @@ unsigned function_info::get_args_aligned_size() { std::string name = p.get_name(); symbol *param = m_symtab->lookup(name.c_str()); - size_t arg_size = p.get_size(); // size of param in bytes + size_t arg_size = p.get_size() / 8; // size of param in bytes total_size = (total_size + arg_size - 1) / arg_size * arg_size; //aligned p.add_offset(total_size); - param_address += total_size; - param->set_address(param_address); + param->set_address(param_address + total_size); total_size += arg_size; } - return (total_size + 3) / 4; //final size aligned to word + return (total_size + 3) / 4 * 4; //final size aligned to word } @@ -1219,8 +1218,6 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) bool skip = false; int op_classification = 0; addr_t pc = next_instr(); - if(pc == 440) - pc = 440; 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() ); diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index c53ea04..2a90cba 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -139,10 +139,11 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * DEV_RUNTIME_REPORT("child_kernel_arg_size " << child_kernel_arg_size); memory_space *child_kernel_param_mem = child_grid->get_param_memory(); size_t param_start_address = 0; - for(unsigned n = 0; n < child_kernel_arg_size; n++) { - unsigned char one_byte; - thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte); - child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); + //copy in word + for(unsigned n = 0; n < child_kernel_arg_size; n += 4) { + unsigned int oneword; + thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 4, &oneword); + child_kernel_param_mem->write(param_start_address + n, 4, &oneword, NULL, NULL); } } else if(arg == 1) { //cudaStream for the child kernel diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index c8208ea..e29b973 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -241,6 +241,7 @@ function_ident_param: IDENTIFIER { add_function_name($1); } LEFT_PAREN {func_hea function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } | VISIBLE_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } + | WEAK_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } | FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } | VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } | EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; func_header(".func"); } @@ -323,6 +324,7 @@ var_spec: space_spec | type_spec | align_spec | EXTERN_DIRECTIVE { add_extern_spec(); } + | WEAK_DIRECTIVE ; align_spec: ALIGN_DIRECTIVE INT_OPERAND { add_alignment_spec($2); } diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 31f3a41..fb17eed 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -100,7 +100,7 @@ void *gpgpu_sim_thread_concurrent(void*) printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n"); fflush(stdout); } - while( g_stream_manager->empty_protected() && !g_sim_done ) + while( g_stream_manager->empty() && !g_sim_done ) ; if(g_debug_execution >= 3) { printf("GPGPU-Sim: ** START simulation thread (detected work) **\n"); -- cgit v1.3 From dafeb411265dbc0228889fe97d85b00f71363f10 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 1 Jul 2016 06:54:46 -0400 Subject: MOD: compute child parameter size --- src/cuda-sim/cuda-sim.cc | 9 +++++++-- src/cuda-sim/ptx_ir.cc | 1 + src/cuda-sim/ptx_ir.h | 3 +++ 3 files changed, 11 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 980afc8..f5d8a88 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1066,7 +1066,10 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg } unsigned function_info::get_args_aligned_size() { - + + if(m_args_aligned_size >= 0) + return m_args_aligned_size; + unsigned param_address = 0; unsigned int total_size = 0; for( std::map::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { @@ -1081,7 +1084,9 @@ unsigned function_info::get_args_aligned_size() { total_size += arg_size; } - return (total_size + 3) / 4 * 4; //final size aligned to word + m_args_aligned_size = (total_size + 3) / 4 * 4; //final size aligned to word + + return m_args_aligned_size; } diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index bdc8381..176eb14 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1292,6 +1292,7 @@ function_info::function_info(int entry_point ) m_kernel_info.regs = 0; m_kernel_info.smem = 0; m_local_mem_framesize = 0; + m_args_aligned_size = -1; } unsigned function_info::print_insn( unsigned pc, FILE * fp ) const diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index a7ca27e..a54ae41 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -291,6 +291,9 @@ private: std::list m_initializer; static unsigned sm_next_uid; + + //parameter size for device kernels + int m_args_aligned_size; }; class symbol_table { -- cgit v1.3 From 0b37da2434beb66713754869a1de775b82a72283 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 1 Jul 2016 07:00:20 -0400 Subject: ADD: launch all device kernels at once in functional simulator --- src/cuda-sim/cuda-sim.cc | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index f5d8a88..276cb9d 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -48,6 +48,7 @@ #include "../gpgpusim_entrypoint.h" #include "decuda_pred_table/decuda_pred_table.h" #include "../stream_manager.h" +#include "cuda_device_runtime.h" int gpgpu_ptx_instruction_classification; void ** g_inst_classification_stat = NULL; @@ -1776,13 +1777,15 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) g_the_gpu->getShaderCoreConfig()->warp_size ); cta.execute(); + + launch_all_device_kernels(); } //registering this kernel as done - extern stream_manager *g_stream_manager; //openCL kernel simulation calls don't register the kernel so we don't register its exit if(!openCL) { + extern stream_manager *g_stream_manager; g_stream_manager->register_finished_kernel(kernel.get_uid()); } -- cgit v1.3 From 8ef2e4eb13093c59190439800fdd0cc552a3779e Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 14 Nov 2014 18:45:46 -0500 Subject: ADD: add cdp latency --- src/abstract_hardware_model.h | 8 ++++++-- src/cuda-sim/cuda-sim.cc | 21 ++++++++++++++++----- src/cuda-sim/ptx_ir.cc | 10 ++++++---- src/gpgpu-sim/shader.cc | 26 ++++++++++++++++++++++++++ src/gpgpu-sim/shader.h | 13 +++++++++++++ 5 files changed, 67 insertions(+), 11 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 3a268ad..45334b6 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -867,7 +867,7 @@ public: m_mem_accesses_created=false; m_cache_hit=false; m_is_printf=false; - m_is_cdp = false; + m_is_cdp = 0; } virtual ~warp_inst_t(){ } @@ -1020,7 +1020,6 @@ protected: unsigned cycles; // used for implementing initiation interval delay bool m_isatomic; bool m_is_printf; - bool m_is_cdp; unsigned m_warp_id; unsigned m_dynamic_warp_id; const core_config *m_config; @@ -1041,6 +1040,11 @@ protected: std::list m_accessq; static unsigned sm_next_uid; + + //Jin: cdp support +public: + int m_is_cdp; + }; void move_warp( warp_inst_t *&dst, warp_inst_t *&src ); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 276cb9d..9ecd92b 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -64,6 +64,8 @@ unsigned gpgpu_param_num_shaders = 0; char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp; char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp; +char *cdp_latency_str; +unsigned cdp_latency[4]; void ptx_opcocde_latency_options (option_parser_t opp) { option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int, @@ -90,6 +92,11 @@ void ptx_opcocde_latency_options (option_parser_t opp) { "Opcode initiation intervals for double precision floating points " "Default 8,8,8,8,130", "8,8,8,8,130"); + option_parser_register(opp, "-cdp_latency", OPT_CSTR, &cdp_latency_str, + "CDP API latency " + "Default 1,7200,19320,1680", + "1,7200,19320,1680"); } static address_type get_converge_point(address_type pc); @@ -609,6 +616,8 @@ void ptx_instruction::set_opcode_and_latency() sscanf(opcode_initiation_dp, "%u,%u,%u,%u,%u", &dp_init[0],&dp_init[1],&dp_init[2], &dp_init[3],&dp_init[4]); + sscanf(cdp_latency_str, "%u,%u,%u,%u", + &cdp_latency[0],&cdp_latency[1],&cdp_latency[2], &cdp_latency[3]); if(!m_operands.empty()){ std::vector::iterator it; @@ -639,19 +648,21 @@ void ptx_instruction::set_opcode_and_latency() case MEMBAR_OP: op = MEMORY_BARRIER_OP; break; case CALL_OP: { - if(m_is_printf || m_is_cdp) + if(m_is_printf || m_is_cdp) { op = ALU_OP; + } else op = CALL_OPS; break; } case CALLP_OP: { - if(m_is_printf || m_is_cdp) + if(m_is_printf || m_is_cdp) { op = ALU_OP; - else - op = CALL_OPS; - break; + } + else + op = CALL_OPS; + break; } case RET_OP: case RETP_OP: op = RET_OPS;break; case ADD_OP: case ADDP_OP: case ADDC_OP: case SUB_OP: case SUBC_OP: diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 176eb14..4931213 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1241,10 +1241,12 @@ ptx_instruction::ptx_instruction( int opcode, if (fname =="vprintf"){ m_is_printf = true; } - if (fname == "cudaGetParameterBufferV2" - || fname == "cudaLaunchDeviceV2" - || fname == "cudaStreamCreateWithFlags") - m_is_cdp = true; + if(fname == "cudaGetParameterBufferV2") + m_is_cdp = 1; + if(fname == "cudaStreamCreateWithFlags") + m_is_cdp = 2; + if(fname == "cudaLaunchDeviceV2") + m_is_cdp = 3; } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index cd38cb7..e85c4a8 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -841,6 +841,13 @@ void scheduler_unit::cycle() unsigned max_issue = m_shader->m_config->gpgpu_max_insn_issue_per_warp; while( !warp(warp_id).waiting() && !warp(warp_id).ibuffer_empty() && (checked < max_issue) && (checked <= issued) && (issued < max_issue) ) { const warp_inst_t *pI = warp(warp_id).ibuffer_next_inst(); + //Jin: handle cdp latency; + if(pI->m_is_cdp && warp(warp_id).m_cdp_latency > 0) { + assert(warp(warp_id).m_cdp_dummy); + warp(warp_id).m_cdp_latency--; + break; + } + bool valid = warp(warp_id).ibuffer_next_valid(); bool warp_inst_issued = false; unsigned pc,rpc; @@ -875,6 +882,25 @@ void scheduler_unit::cycle() bool sp_pipe_avail = m_sp_out->has_free(); bool sfu_pipe_avail = m_sfu_out->has_free(); if( sp_pipe_avail && (pI->op != SFU_OP) ) { + + //Jin: special for CDP api + if(pI->m_is_cdp && !warp(warp_id).m_cdp_dummy) { + assert(warp(warp_id).m_cdp_latency == 0); + + extern unsigned cdp_latency[3]; + if(pI->m_is_cdp != 3) + warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1]; + else //cudaLaunchDeviceV2 + warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1] + + cdp_latency[pI->m_is_cdp] * active_mask.count(); + warp(warp_id).m_cdp_dummy = true; + break; + } + else if(pI->m_is_cdp && warp(warp_id).m_cdp_dummy) { + assert(warp(warp_id).m_cdp_latency == 0); + warp(warp_id).m_cdp_dummy = false; + } + // always prefer SP pipe for operations that can use both SP and SFU pipelines m_shader->issue_warp(*m_sp_out,pI,active_mask,warp_id); issued++; diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index fcbc8aa..882868e 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -108,6 +108,10 @@ public: m_last_fetch=0; m_next=0; m_inst_at_barrier=NULL; + + //Jin: cdp support + m_cdp_latency = 0; + m_cdp_dummy = false; } void init( address_type start_pc, unsigned cta_id, @@ -124,6 +128,10 @@ public: n_completed -= active.count(); // active threads are not yet completed m_active_threads = active; m_done_exit=false; + + //Jin: cdp support + m_cdp_latency = 0; + m_cdp_dummy = false; } bool functional_done() const; @@ -260,6 +268,11 @@ private: unsigned m_stores_outstanding; // number of store requests sent but not yet acknowledged unsigned m_inst_in_pipeline; + + //Jin: cdp support +public: + unsigned int m_cdp_latency; + bool m_cdp_dummy; }; -- cgit v1.3 From bbcb492c0f6d887c4034bd15adf57420dd735c5e Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 14 Nov 2014 20:01:10 -0500 Subject: ADD: add separate cdp latency --- src/cuda-sim/cuda-sim.cc | 16 +++++++++------- src/cuda-sim/ptx_ir.cc | 6 +++--- src/gpgpu-sim/shader.cc | 7 ++++--- 3 files changed, 16 insertions(+), 13 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 9ecd92b..c87e3e4 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -65,7 +65,7 @@ unsigned gpgpu_param_num_shaders = 0; char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp; char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp; char *cdp_latency_str; -unsigned cdp_latency[4]; +unsigned cdp_latency[5]; void ptx_opcocde_latency_options (option_parser_t opp) { option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int, @@ -93,10 +93,11 @@ void ptx_opcocde_latency_options (option_parser_t opp) { "Default 8,8,8,8,130", "8,8,8,8,130"); option_parser_register(opp, "-cdp_latency", OPT_CSTR, &cdp_latency_str, - "CDP API latency " - "Default 1,7200,19320,1680", - "1,7200,19320,1680"); + "CDP API latency " + "Default 7200,8000,100,12000,1600", + "7200,8000,100,12000,1600"); } static address_type get_converge_point(address_type pc); @@ -616,8 +617,9 @@ void ptx_instruction::set_opcode_and_latency() sscanf(opcode_initiation_dp, "%u,%u,%u,%u,%u", &dp_init[0],&dp_init[1],&dp_init[2], &dp_init[3],&dp_init[4]); - sscanf(cdp_latency_str, "%u,%u,%u,%u", - &cdp_latency[0],&cdp_latency[1],&cdp_latency[2], &cdp_latency[3]); + sscanf(cdp_latency_str, "%u,%u,%u,%u,%u", + &cdp_latency[0],&cdp_latency[1],&cdp_latency[2], + &cdp_latency[3],&cdp_latency[4]); if(!m_operands.empty()){ std::vector::iterator it; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 4931213..783c885 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1241,12 +1241,12 @@ ptx_instruction::ptx_instruction( int opcode, if (fname =="vprintf"){ m_is_printf = true; } - if(fname == "cudaGetParameterBufferV2") - m_is_cdp = 1; if(fname == "cudaStreamCreateWithFlags") + m_is_cdp = 1; + if(fname == "cudaGetParameterBufferV2") m_is_cdp = 2; if(fname == "cudaLaunchDeviceV2") - m_is_cdp = 3; + m_is_cdp = 4; } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index e85c4a8..8ce2146 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -887,12 +887,13 @@ void scheduler_unit::cycle() if(pI->m_is_cdp && !warp(warp_id).m_cdp_dummy) { assert(warp(warp_id).m_cdp_latency == 0); - extern unsigned cdp_latency[3]; - if(pI->m_is_cdp != 3) + extern unsigned cdp_latency[5]; + if(pI->m_is_cdp == 1) warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1]; - else //cudaLaunchDeviceV2 + else //cudaLaunchDeviceV2 and cudaGetParameterBufferV2 warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1] + cdp_latency[pI->m_is_cdp] * active_mask.count(); + printf("set latency %d\n", warp(warp_id).m_cdp_latency); warp(warp_id).m_cdp_dummy = true; break; } -- cgit v1.3 From b70a5a69fa14cc01f707b910f8f021e36067922f Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sat, 15 Nov 2014 05:44:03 -0500 Subject: ADD: print kernel parameter size footprint. BUG: concurrent kernels on same shader, should use hw_cta_id to store shared mem info --- src/abstract_hardware_model.cc | 2 ++ src/cuda-sim/cuda-sim.cc | 3 ++- src/cuda-sim/cuda_device_runtime.cc | 7 ++++++- src/gpgpu-sim/gpu-sim.cc | 23 +++++++++++++++-------- src/stream_manager.cc | 6 ++++-- 5 files changed, 29 insertions(+), 12 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 4db5f2f..be5c5b9 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -618,6 +618,8 @@ bool kernel_info_t::children_all_finished() { void kernel_info_t::notify_parent_finished() { if(m_parent_kernel) { + extern unsigned long long g_total_param_size; + g_total_param_size -= ((m_kernel_entry->get_args_aligned_size() + 255)/256*256); m_parent_kernel->remove_child(this); g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index c87e3e4..3f5af7e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1499,7 +1499,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, unsigned max_cta_per_sm = num_threads/cta_size; // e.g., 256 / 48 = 5 assert( max_cta_per_sm > 0 ); - unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid; + //unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid; + unsigned sm_idx = hw_cta_id*gpgpu_param_num_shaders + sid; if ( shared_memory_lookup.find(sm_idx) == shared_memory_lookup.end() ) { if ( g_debug_execution >= 1 ) { diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 1b8c8d9..66cd063 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -59,6 +59,8 @@ public: std::map g_cuda_device_launch_param_map; std::list g_cuda_device_launch_op; extern stream_manager *g_stream_manager; +unsigned long long g_total_param_size = 0; +unsigned long long g_max_total_param_size = 0; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -111,7 +113,10 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i //get total child kernel argument size and malloc buffer in global memory unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); + g_total_param_size += ((child_kernel_arg_size + 255) / 256 * 256); DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); + if(g_total_param_size > g_max_total_param_size) + g_max_total_param_size = g_total_param_size; //store param buffer address and launch config device_launch_config_t device_launch_config(grid_dim, block_dim, shared_mem, child_kernel_entry); @@ -176,7 +181,7 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * //create child kernel_info_t and index it with parameter_buffer address device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry); - device_grid->launch_cycle = gpu_sim_cycle; + device_grid->launch_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; kernel_info_t & parent_grid = thread->get_kernel(); DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 27362cf..a9da1c9 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -524,20 +524,26 @@ bool gpgpu_sim::get_more_cta_left() const kernel_info_t *gpgpu_sim::select_kernel() { if(m_running_kernels[m_last_issued_kernel] && - !m_running_kernels[m_last_issued_kernel]->no_more_ctas_to_run()) + !m_running_kernels[m_last_issued_kernel]->no_more_ctas_to_run()) { + unsigned launch_uid = m_running_kernels[m_last_issued_kernel]->get_uid(); + if(std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()) { + m_running_kernels[m_last_issued_kernel]->start_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; + m_executed_kernel_uids.push_back(launch_uid); + m_executed_kernel_names.push_back(m_running_kernels[m_last_issued_kernel]->name()); + } return m_running_kernels[m_last_issued_kernel]; + } for(unsigned n=0; n < m_running_kernels.size(); n++ ) { unsigned idx = (n+m_last_issued_kernel+1)%m_config.max_concurrent_kernel; if( kernel_more_cta_left(m_running_kernels[idx]) ){ m_last_issued_kernel=idx; - m_running_kernels[idx]->start_cycle = gpu_sim_cycle; + m_running_kernels[idx]->start_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; // record this kernel for stat print if it is the first time this kernel is selected for execution unsigned launch_uid = m_running_kernels[idx]->get_uid(); - if (std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()) { - m_executed_kernel_uids.push_back(launch_uid); - m_executed_kernel_names.push_back(m_running_kernels[idx]->name()); - } + assert(std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()); + m_executed_kernel_uids.push_back(launch_uid); + m_executed_kernel_names.push_back(m_running_kernels[idx]->name()); return m_running_kernels[idx]; } @@ -561,7 +567,7 @@ void gpgpu_sim::set_kernel_done( kernel_info_t *kernel ) std::vector::iterator k; for( k=m_running_kernels.begin(); k!=m_running_kernels.end(); k++ ) { if( *k == kernel ) { - kernel->end_cycle = gpu_sim_cycle; + kernel->end_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; *k = NULL; break; } @@ -942,7 +948,8 @@ void gpgpu_sim::gpu_print_stat() printf("gpu_tot_ipc = %12.4f\n", (float)(gpu_tot_sim_insn+gpu_sim_insn) / (gpu_tot_sim_cycle+gpu_sim_cycle)); printf("gpu_tot_issued_cta = %lld\n", gpu_tot_issued_cta + m_total_cta_launched); - + extern unsigned long long g_max_total_param_size; + fprintf(statfout, "max_total_param_size = %llu\n", g_max_total_param_size); // performance counter for stalls due to congestion. printf("gpu_stall_dramfull = %d\n", gpu_stall_dramfull); diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 5bd7737..3459c6d 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -219,7 +219,7 @@ bool stream_manager::operation( bool * sim) { bool check=check_finished_kernel(); pthread_mutex_lock(&m_lock); - if(check)m_gpu->print_stats(); +// if(check)m_gpu->print_stats(); stream_operation op =front(); if(!op.do_operation( m_gpu )) //not ready to execute { @@ -260,8 +260,10 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) // kernel_stat << ", parent " << kernel->get_parent()->get_uid() << // ", launch " << kernel->launch_cycle; // kernel_stat<< ", start " << kernel->start_cycle << -// ", end " << kernel->end_cycle << ", retire " << gpu_sim_cycle << "\n"; +// ", end " << kernel->end_cycle << ", retire " << gpu_sim_cycle + gpu_tot_sim_cycle << "\n"; // printf("kernel %d finishes, retires from stream %d\n", grid_uid, stream->get_uid()); +// kernel_stat.flush(); +// kernel_stat.close(); stream->record_next_done(); m_grid_id_to_stream.erase(grid_uid); kernel->notify_parent_finished(); -- cgit v1.3 From 68336f112117bcef5b943650819a6764e9ebf4ce Mon Sep 17 00:00:00 2001 From: sspenst Date: Wed, 24 Aug 2016 15:24:19 -0700 Subject: Added shfl instruction --- src/abstract_hardware_model.h | 1 + src/cuda-sim/cuda-sim.cc | 20 +++++++++++- src/cuda-sim/instructions.cc | 75 +++++++++++++++++++++++++++++++++++++++++++ src/cuda-sim/opcodes.def | 1 + src/cuda-sim/opcodes.h | 4 ++- src/cuda-sim/ptx.l | 6 ++++ src/cuda-sim/ptx.y | 8 +++++ src/cuda-sim/ptx_ir.cc | 6 ++++ src/cuda-sim/ptx_ir.h | 2 ++ src/cuda-sim/ptx_sim.cc | 21 ++++++++++++ src/cuda-sim/ptx_sim.h | 12 +++++++ 11 files changed, 154 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index b29f918..d0c807d 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -1051,6 +1051,7 @@ class core_t { warp_inst_t getExecuteWarp(unsigned warpId); void get_pdom_stack_top_info( unsigned warpId, unsigned *pc, unsigned *rpc ) const; kernel_info_t * get_kernel_info(){ return m_kernel;} + class ptx_thread_info ** get_thread_info() { return m_thread; } unsigned get_warp_size() const { return m_warp_size; } void and_reduction(unsigned ctaid, unsigned barid, bool value) { reduction_storage[ctaid][barid] &= value; } void or_reduction(unsigned ctaid, unsigned barid, bool value) { reduction_storage[ctaid][barid] |= value; } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 09e9a81..8bf4ec8 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -769,6 +769,10 @@ void ptx_instruction::set_opcode_and_latency() initiation_interval = dp_init[2]; op = SFU_OP; break; + case SHFL_OP: + latency = 32; + initiation_interval = 15; + break; default: break; } @@ -845,8 +849,10 @@ void ptx_instruction::pre_decode() switch ( get_opcode() ) { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: has_dst = (DST!=0); break; +#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: has_dst = (DST!=0); break; #include "opcodes.def" #undef OP_DEF +#undef OP_W_DEF default: printf( "Execution error: Invalid opcode (0x%x)\n", get_opcode() ); break; @@ -1240,8 +1246,10 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) } 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; } delete pJ; @@ -1408,6 +1416,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, static std::map shared_memory_lookup; static std::map ptx_cta_lookup; + static std::map ptx_warp_lookup; static std::map > local_memory_lookup; if ( *thread_info != NULL ) { @@ -1486,7 +1495,16 @@ 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(); // num_threads should be threads in the warp + ptx_warp_lookup[hw_warp_id] = warp_info; + } else { + warp_info = ptx_warp_lookup[hw_warp_id]; + } + thd->m_warp_info = warp_info; + memory_space *local_mem = NULL; std::map::iterator l = local_mem_lookup.find(new_tid); if ( l != local_mem_lookup.end() ) { diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 7b0f4fa..05ba78f 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -47,8 +47,10 @@ unsigned ptx_instruction::g_num_ptx_inst_uid=0; const char *g_opcode_string[NUM_OPCODES] = { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) STR, +#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) STR, #include "opcodes.def" #undef OP_DEF +#undef OP_W_DEF }; void inst_not_implemented( const ptx_instruction * pI ) ; @@ -3516,6 +3518,79 @@ void set_impl( const ptx_instruction *pI, ptx_thread_info *thread ) } +void shfl_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) +{ + unsigned i_type = pI->get_type(); + int tid = inst.warp_id() * core->get_warp_size(); + ptx_thread_info *thread = core->get_thread_info()[tid]; + ptx_warp_info *warp_info = thread->m_warp_info; + int lane = warp_info->get_done_threads(); + thread = core->get_thread_info()[tid + lane]; + + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + const operand_info &src3 = pI->src3(); + int bval = (thread->get_operand_value(src2, dst, i_type, thread, 1)).u32; + int cval = (thread->get_operand_value(src3, dst, i_type, thread, 1)).u32; + int mask = cval >> 8; + cval &= 0x1F; + + int maxLane = (lane & mask) | (cval & ~mask); + int minLane = lane & mask; + + int src_idx; + int p; + switch(pI->shfl_op()) { + case UP_OPTION: + src_idx = lane - bval; + p = (src_idx >= maxLane); + break; + case DOWN_OPTION: + src_idx = lane + bval; + p = (src_idx <= maxLane); + break; + case BFLY_OPTION: + src_idx = lane ^ bval; + p = (src_idx <= maxLane); + break; + case IDX_OPTION: + src_idx = minLane | (bval & ~mask); + p = (src_idx <= maxLane); + break; + default: + printf("GPGPU-Sim PTX: ERROR: Unrecognized shfl option\n"); + assert(0); + break; + } + // copy from own lane + if (!p) src_idx = lane; + // copy input from lane src_idx + ptx_reg_t data; + /*if (inst.active(src_idx) && i_type == PRED_TYPE) { + ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; + data = source->get_operand_value(src1, dst, i_type, source, 1); + data.pred = p; + } else { + printf("GPGPU-Sim PTX: WARNING: shfl input value unpredictable for inactive/predicated-off threads in a warp\n"); + data.u32 = 0; + }*/ + if (inst.active(src_idx)) { + ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; + data = source->get_operand_value(src1, dst, i_type, source, 1); + } + if (i_type == PRED_TYPE) { + data.pred = p; + } + thread->set_operand_value(dst, data, i_type, thread, pI); + + // keep track of the number of threads that have executed in the warp + warp_info->inc_done_threads(); + if (warp_info->get_done_threads() == inst.active_count()) { + warp_info->reset_done_threads(); + } +} + void shl_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t a, b, d; diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 2ee6976..e1b1422 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -98,6 +98,7 @@ OP_DEF(SAD_OP,sad_impl,"sad",1,1) OP_DEF(SELP_OP,selp_impl,"selp",1,1) OP_DEF(SETP_OP,setp_impl,"setp",1,1) OP_DEF(SET_OP,set_impl,"set",1,1) +OP_W_DEF(SHFL_OP,shfl_impl,"shfl",1,10) OP_DEF(SHL_OP,shl_impl,"shl",1,1) OP_DEF(SHR_OP,shr_impl,"shr",1,1) OP_DEF(SIN_OP,sin_impl,"sin",1,4) diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h index 871091c..aa133da 100644 --- a/src/cuda-sim/opcodes.h +++ b/src/cuda-sim/opcodes.h @@ -30,9 +30,11 @@ enum opcode_t { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) OP, +#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) OP, #include "opcodes.def" - NUM_OPCODES + NUM_OPCODES #undef OP_DEF +#undef OP_W_DEF }; enum special_regs { diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 88ccf6a..8fac4ac 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -115,6 +115,7 @@ sad TC; ptx_lval.int_value = SAD_OP; return OPCODE; selp TC; ptx_lval.int_value = SELP_OP; return OPCODE; setp TC; ptx_lval.int_value = SETP_OP; return OPCODE; set TC; ptx_lval.int_value = SET_OP; return OPCODE; +shfl TC; ptx_lval.int_value = SHFL_OP; return OPCODE; shl TC; ptx_lval.int_value = SHL_OP; return OPCODE; shr TC; ptx_lval.int_value = SHR_OP; return OPCODE; sin TC; ptx_lval.int_value = SIN_OP; return OPCODE; @@ -329,6 +330,11 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.nc TC; return NC_OPTION; +\.up TC; return UP_OPTION; +\.down TC; return DOWN_OPTION; +\.bfly TC; return BFLY_OPTION; +\.idx TC; return IDX_OPTION; + \.popc TC; return ATOMIC_POPC; \.and TC; return ATOMIC_AND; \.or TC; return ATOMIC_OR; diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 4de39d1..166b15d 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -194,6 +194,10 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %token WB_OPTION; %token WT_OPTION; %token NC_OPTION; +%token UP_OPTION; +%token DOWN_OPTION; +%token BFLY_OPTION; +%token IDX_OPTION; %type function_decl_header %type function_decl @@ -451,6 +455,10 @@ option: type_spec | WB_OPTION { add_option(WB_OPTION); } | WT_OPTION { add_option(WT_OPTION); } | NC_OPTION { add_option(NC_OPTION); } + | UP_OPTION { add_option(UP_OPTION); } + | DOWN_OPTION { add_option(DOWN_OPTION); } + | BFLY_OPTION { add_option(BFLY_OPTION); } + | IDX_OPTION { add_option(IDX_OPTION); } ; atomic_operation_spec: ATOMIC_AND { add_option(ATOMIC_AND); } diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 2eccabc..4cfe1b9 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1171,6 +1171,12 @@ ptx_instruction::ptx_instruction( int opcode, break; case NC_OPTION: break; + case UP_OPTION: + case DOWN_OPTION: + case BFLY_OPTION: + case IDX_OPTION: + m_shfl_op = last_ptx_inst_option; + break; default: assert(0); break; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 601a13d..0abbc83 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -993,6 +993,7 @@ public: unsigned saturation_mode() const { return m_saturation_mode;} unsigned dimension() const { return m_geom_spec;} unsigned barrier_op() const {return m_barrier_op;} + unsigned shfl_op() const {return m_shfl_op;} enum vote_mode_t { vote_any, vote_all, vote_uni, vote_ballot }; enum vote_mode_t vote_mode() const { return m_vote_mode; } @@ -1058,6 +1059,7 @@ private: unsigned m_compare_op; unsigned m_saturation_mode; unsigned m_barrier_op; + unsigned m_shfl_op; std::list m_scalar_type; memory_space_t m_space_spec; diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc index 09844ae..a3e43aa 100644 --- a/src/cuda-sim/ptx_sim.cc +++ b/src/cuda-sim/ptx_sim.cc @@ -128,6 +128,26 @@ unsigned ptx_cta_info::get_sm_idx() const return m_sm_idx; } +ptx_warp_info::ptx_warp_info() +{ + reset_done_threads(); +} + +unsigned ptx_warp_info::get_done_threads() const +{ + return m_done_threads; +} + +void ptx_warp_info::inc_done_threads() +{ + m_done_threads++; +} + +void ptx_warp_info::reset_done_threads() +{ + m_done_threads = 0; +} + unsigned g_ptx_thread_info_uid_next=1; unsigned g_ptx_thread_info_delete_count=0; @@ -153,6 +173,7 @@ ptx_thread_info::ptx_thread_info( kernel_info_t &kernel ) m_last_memory_space = undefined_space; m_branch_taken = 0; m_shared_mem = NULL; + m_warp_info = NULL; m_cta_info = NULL; m_local_mem = NULL; m_symbol_table = NULL; diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h index f926e6d..449511f 100644 --- a/src/cuda-sim/ptx_sim.h +++ b/src/cuda-sim/ptx_sim.h @@ -167,6 +167,17 @@ private: std::set m_dangling_pointers; }; +class ptx_warp_info { +public: + ptx_warp_info(); + unsigned get_done_threads() const; + void inc_done_threads(); + void reset_done_threads(); + +private: + unsigned m_done_threads; +}; + class symbol; struct stack_entry { @@ -425,6 +436,7 @@ public: dram_callback_t m_last_dram_callback; memory_space *m_shared_mem; memory_space *m_local_mem; + ptx_warp_info *m_warp_info; ptx_cta_info *m_cta_info; ptx_reg_t m_last_set_operand_value; -- cgit v1.3 From bffac032402b35acb81200540f6f808bd6f851d8 Mon Sep 17 00:00:00 2001 From: sspenst Date: Wed, 24 Aug 2016 15:30:16 -0700 Subject: Cleanup --- src/cuda-sim/cuda-sim.cc | 2 +- src/cuda-sim/instructions.cc | 8 -------- 2 files changed, 1 insertion(+), 9 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 8bf4ec8..d3c928d 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1498,7 +1498,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &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(); // num_threads should be threads in the warp + warp_info = new ptx_warp_info(); ptx_warp_lookup[hw_warp_id] = warp_info; } else { warp_info = ptx_warp_lookup[hw_warp_id]; diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 05ba78f..4c9392b 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -3567,14 +3567,6 @@ void shfl_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) if (!p) src_idx = lane; // copy input from lane src_idx ptx_reg_t data; - /*if (inst.active(src_idx) && i_type == PRED_TYPE) { - ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; - data = source->get_operand_value(src1, dst, i_type, source, 1); - data.pred = p; - } else { - printf("GPGPU-Sim PTX: WARNING: shfl input value unpredictable for inactive/predicated-off threads in a warp\n"); - data.u32 = 0; - }*/ if (inst.active(src_idx)) { ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; data = source->get_operand_value(src1, dst, i_type, source, 1); -- cgit v1.3 From 623a88e5d5c6c3edb94404ef6e5ea100caec9deb Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 29 Aug 2016 18:10:00 -0700 Subject: MOD: Add macros to turn off cuda_device_runtime for CUDA < 5.0 --- src/cuda-sim/cuda-sim.cc | 2 ++ src/cuda-sim/cuda_device_runtime.cc | 10 +++++++--- src/cuda-sim/cuda_device_runtime.h | 3 ++- src/cuda-sim/instructions.cc | 2 ++ src/gpgpu-sim/gpu-sim.cc | 2 ++ src/gpgpu-sim/shader.cc | 2 +- 6 files changed, 16 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 3f5af7e..ec51779 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1792,7 +1792,9 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) ); cta.execute(); +#if (CUDART_VERSION >= 5000) launch_all_device_kernels(); +#endif } //registering this kernel as done diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 12c83d2..4a8ffe5 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -1,9 +1,15 @@ //Jin: cuda_device_runtime.cc //Defines CUDA device runtime APIs for CDP support + #include #include +unsigned long long g_total_param_size = 0; +unsigned long long g_max_total_param_size = 0; + + +#if (CUDART_VERSION >= 5000) #define __CUDA_RUNTIME_API_H__ #include @@ -59,8 +65,6 @@ public: std::map g_cuda_device_launch_param_map; std::list g_cuda_device_launch_op; extern stream_manager *g_stream_manager; -unsigned long long g_total_param_size = 0; -unsigned long long g_max_total_param_size = 0; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -313,4 +317,4 @@ void launch_all_device_kernels() { launch_one_device_kernel(); } } - +#endif diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 385d605..6dbcd71 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -1,6 +1,6 @@ //Jin: cuda_device_runtime.h //Defines CUDA device runtime APIs for CDP support - +#if (CUDART_VERSION >= 5000) #pragma once void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); @@ -8,3 +8,4 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void launch_all_device_kernels(); void launch_one_device_kernel(); +#endif diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 5d909d3..e68f9fd 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1414,6 +1414,7 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) return; } +#if (CUDART_VERSION >= 5000) //Jin: handle device runtime apis for CDP else if(fname == "cudaGetParameterBufferV2") { gpgpusim_cuda_getParameterBufferV2(pI, thread, target_func); @@ -1427,6 +1428,7 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpusim_cuda_streamCreateWithFlags(pI, thread, target_func); return; } +#endif // read source arguements into register specified in declaration of function arg_buffer_list_t arg_values; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 0b4b2f6..363fe5a 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1513,8 +1513,10 @@ void gpgpu_sim::cycle() try_snap_shot(gpu_sim_cycle); spill_log_to_file (stdout, 0, gpu_sim_cycle); +#if (CUDART_VERSION >= 5000) //launch device kernel launch_one_device_kernel(); +#endif } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index b9caf18..59a2d8b 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -846,7 +846,7 @@ void scheduler_unit::cycle() while( !warp(warp_id).waiting() && !warp(warp_id).ibuffer_empty() && (checked < max_issue) && (checked <= issued) && (issued < max_issue) ) { const warp_inst_t *pI = warp(warp_id).ibuffer_next_inst(); //Jin: handle cdp latency; - if(pI->m_is_cdp && warp(warp_id).m_cdp_latency > 0) { + if(pI && pI->m_is_cdp && warp(warp_id).m_cdp_latency > 0) { assert(warp(warp_id).m_cdp_dummy); warp(warp_id).m_cdp_latency--; break; -- cgit v1.3