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/cuda_device_runtime.h | 7 +++++++ 1 file changed, 7 insertions(+) create mode 100644 src/cuda-sim/cuda_device_runtime.h (limited to 'src/cuda-sim/cuda_device_runtime.h') 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); -- cgit v1.3 From 3bb32d87175d873e7089ad50a0069acc195edb34 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Tue, 21 Oct 2014 01:24:26 -0400 Subject: ADD: add support for cudaStreamCreateWithFlags --- src/cuda-sim/cuda_device_runtime.cc | 61 ++++++++++++++++++++++++++++++++++++- src/cuda-sim/cuda_device_runtime.h | 1 + src/cuda-sim/instructions.cc | 4 +++ src/cuda-sim/ptx_ir.cc | 3 +- 4 files changed, 67 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim/cuda_device_runtime.h') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 148471b..be87f6a 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -152,11 +152,12 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * 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()); + DEV_RUNTIME_REPORT("launching child kernel to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); } else { assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); + DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << ": " << child_stream); } - DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << " " << child_stream); } } @@ -179,3 +180,61 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); } + + +//Handling device runtime api: +//cudaError_t cudaStreamCreateWithFlags ( cudaStream_t* pStream, unsigned int flags) +//flags can only be cudaStreamNonBlocking +void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) { + DEV_RUNTIME_REPORT("Calling cudaStreamCreateWithFlags"); + + unsigned n_return = target_func->has_return(); + assert(n_return); + unsigned n_args = target_func->num_args(); + assert( n_args == 2 ); + + size_t generic_pStream_addr; + addr_t pStream_addr; + unsigned int flags; + 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); //cudaStreamCreateWithFlags_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) {//cudaStream_t * pStream, address of cudaStream_t + assert(size == sizeof(cudaStream_t *)); + thread->m_local_mem->read(from_addr, size, &generic_pStream_addr); + + //pStream should be non-zero address in local memory + pStream_addr = generic_to_local(thread->get_hw_sid(), thread->get_hw_tid(), generic_pStream_addr); + + DEV_RUNTIME_REPORT("pStream locating at local memory " << pStream_addr); + } + else if(arg == 1) { //unsigned int flags, should be cudaStreamNonBlocking + assert(size == sizeof(unsigned int)); + thread->m_local_mem->read(from_addr, size, &flags); + assert(flags == cudaStreamNonBlocking); + } + } + + //create stream and write back to param0 + CUstream_st * stream = thread->get_kernel().create_stream_cta(thread->get_ctaid()); + DEV_RUNTIME_REPORT("Create stream " << stream->get_uid() << ": " << stream); + thread->m_local_mem->write(pStream_addr, sizeof(cudaStream_t), &stream, NULL, NULL); + + //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("cudaStreamCreateWithFlags 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 index 1b10407..7447d5a 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -5,3 +5,4 @@ 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); +void gpgpusim_cuda_streamCreateWithFlags(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 1c47ad3..5d909d3 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1423,6 +1423,10 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpusim_cuda_launchDeviceV2(pI, thread, target_func); return; } + else if(fname == "cudaStreamCreateWithFlags") { + gpgpusim_cuda_streamCreateWithFlags(pI, thread, target_func); + return; + } // read source arguements into register specified in declaration of function arg_buffer_list_t arg_values; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 861f0dc..bdc8381 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1242,7 +1242,8 @@ ptx_instruction::ptx_instruction( int opcode, m_is_printf = true; } if (fname == "cudaGetParameterBufferV2" - || fname == "cudaLaunchDeviceV2") + || fname == "cudaLaunchDeviceV2" + || fname == "cudaStreamCreateWithFlags") m_is_cdp = true; } -- cgit v1.3 From 3656376ccd83d5c514389c4c3818f1969bc30e0c Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 1 Jul 2016 05:29:58 -0400 Subject: MOD: schedule one child kernel each cycle --- src/cuda-sim/cuda_device_runtime.cc | 214 ++++++++++++++++++++++++------------ src/cuda-sim/cuda_device_runtime.h | 2 + src/gpgpu-sim/gpu-sim.cc | 4 + src/stream_manager.cc | 2 +- 4 files changed, 148 insertions(+), 74 deletions(-) (limited to 'src/cuda-sim/cuda_device_runtime.h') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 2a90cba..df2653e 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -20,8 +20,45 @@ std::cout.flush(); \ } -std::map g_cuda_device_launch_map; -extern stream_manager * g_stream_manager; +class device_launch_config_t { + +public: + device_launch_config_t() {} + + device_launch_config_t(dim3 _grid_dim, + dim3 _block_dim, + unsigned int _shared_mem, + function_info * _entry): + grid_dim(_grid_dim), + block_dim(_block_dim), + shared_mem(_shared_mem), + entry(_entry) {} + + dim3 grid_dim; + dim3 block_dim; + unsigned int shared_mem; + function_info * entry; + +}; + +class device_launch_operation_t { + +public: + device_launch_operation_t() {} + device_launch_operation_t(kernel_info_t *_grid, + CUstream_st * _stream) : + grid(_grid), stream(_stream) {} + + kernel_info_t * grid; //a new child grid + + CUstream_st * stream; + +}; + + +std::map g_cuda_device_launch_param_map; +std::list g_cuda_device_launch_op; +extern stream_manager *g_stream_manager; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -35,8 +72,8 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i assert( n_args == 4 ); function_info * child_kernel_entry; - struct dim3 gridDim, blockDim; - unsigned int sharedMem; + struct dim3 grid_dim, block_dim; + unsigned int shared_mem; for( unsigned arg=0; arg < n_args; arg ++ ) { const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# @@ -48,54 +85,48 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i if(arg == 0) {//function_info* for the child kernel unsigned long long buf; - assert(size == sizeof(function_info *)); + 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 == 1) { //dim3 grid_dim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & grid_dim); + DEV_RUNTIME_REPORT("grid (" << grid_dim.x << ", " << grid_dim.y << ", " << grid_dim.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 == 2) { //dim3 block_dim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & block_dim); + DEV_RUNTIME_REPORT("block (" << block_dim.x << ", " << block_dim.y << ", " << block_dim.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); + else if(arg == 3) { //unsigned int shared_mem + assert(size == sizeof(unsigned int)); + thread->m_local_mem->read(from_addr, size, & shared_mem); + DEV_RUNTIME_REPORT("shared memory " << shared_mem); } } - //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); - 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; - - //copy the buffer address to retval0 + //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); + + //store param buffer address and launch config + device_launch_config_t device_launch_config(grid_dim, block_dim, shared_mem, child_kernel_entry); + assert(g_cuda_device_launch_param_map.find(param_buffer) == g_cuda_device_launch_param_map.end()); + g_cuda_device_launch_param_map[param_buffer] = device_launch_config; + + //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 *)); + 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); + thread->m_local_mem->write(ret_param_addr, return_size, ¶m_buffer, NULL, NULL); } @@ -109,10 +140,13 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * unsigned n_args = target_func->num_args(); assert( n_args == 2 ); - kernel_info_t * child_grid = NULL; - function_info * child_kernel_entry = NULL; + kernel_info_t * device_grid = NULL; + function_info * device_kernel_entry = NULL; void * parameter_buffer; struct CUstream_st * child_stream; + device_launch_config_t config; + device_launch_operation_t device_launch_op; + 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_# @@ -123,64 +157,80 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * if(arg == 0) {//paramter buffer for child kernel (in global memory) //get parameter_buffer from the cudaLaunchDeviceV2_param0 - assert(size == sizeof(void *)); + 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(); + assert(g_cuda_device_launch_param_map.find(parameter_buffer) != g_cuda_device_launch_param_map.end()); + config = g_cuda_device_launch_param_map[parameter_buffer]; + //device_grid = op.grid; + device_kernel_entry = config.entry; + DEV_RUNTIME_REPORT("find device kernel " << device_kernel_entry->get_name()); + + //copy data in parameter_buffer to device kernel param memory + unsigned device_kernel_arg_size = device_kernel_entry->get_args_aligned_size(); + DEV_RUNTIME_REPORT("device_kernel_arg_size " << device_kernel_arg_size); + memory_space *device_kernel_param_mem; + + //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); + 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 << + ")"); + device_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid()); + device_launch_op = device_launch_operation_t(device_grid, NULL); + device_kernel_param_mem = device_grid->get_param_memory(); //kernel param + size_t param_start_address = 0; //copy in word - for(unsigned n = 0; n < child_kernel_arg_size; n += 4) { + for(unsigned n = 0; n < device_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); + device_kernel_param_mem->write(param_start_address + n, 4, &oneword, 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); + assert(size == sizeof(cudaStream_t)); + thread->m_local_mem->read(from_addr, size, &child_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()); - DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + DEV_RUNTIME_REPORT("launching child kernel " << device_grid->get_uid() << " to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); } else { assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); - DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + DEV_RUNTIME_REPORT("launching child kernel " << device_grid->get_uid() << " to stream " << child_stream->get_uid() << ": " << child_stream); } + + device_launch_op.stream = 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_cuda_device_launch_op.push_back(device_launch_op); + g_cuda_device_launch_param_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 + 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); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); } @@ -208,7 +258,7 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ addr_t from_addr = actual_param_op.get_symbol()->get_address(); if(arg == 0) {//cudaStream_t * pStream, address of cudaStream_t - assert(size == sizeof(cudaStream_t *)); + assert(size == sizeof(cudaStream_t *)); thread->m_local_mem->read(from_addr, size, &generic_pStream_addr); //pStream should be non-zero address in local memory @@ -217,7 +267,7 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ DEV_RUNTIME_REPORT("pStream locating at local memory " << pStream_addr); } else if(arg == 1) { //unsigned int flags, should be cudaStreamNonBlocking - assert(size == sizeof(unsigned int)); + assert(size == sizeof(unsigned int)); thread->m_local_mem->read(from_addr, size, &flags); assert(flags == cudaStreamNonBlocking); } @@ -231,13 +281,31 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ //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("cudaStreamCreateWithFlags 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 + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaStreamCreateWithFlags 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); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); } + + +void launch_one_device_kernel() { + if(!g_cuda_device_launch_op.empty()) { + device_launch_operation_t &op = g_cuda_device_launch_op.front(); + + stream_operation stream_op = stream_operation(op.grid, g_ptx_sim_mode, op.stream); + g_stream_manager->push(stream_op); + g_cuda_device_launch_op.pop_front(); + } +} + +void launch_all_device_kernels() { + while(!g_cuda_device_launch_op.empty()) { + launch_one_device_kernel(); + } +} + diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 7447d5a..385d605 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -6,3 +6,5 @@ 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); 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(); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 47dbc89..126e007 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -61,6 +61,7 @@ #include "power_stat.h" #include "visualizer.h" #include "stats.h" +#include "../cuda-sim/cuda_device_runtime.h" #ifdef GPGPUSIM_POWER_MODEL #include "power_interface.h" @@ -1370,6 +1371,9 @@ void gpgpu_sim::cycle() } try_snap_shot(gpu_sim_cycle); spill_log_to_file (stdout, 0, gpu_sim_cycle); + + //launch device kernel + launch_one_device_kernel(); } } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 546ab3b..f90d9be 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -217,8 +217,8 @@ stream_manager::stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking ) bool stream_manager::operation( bool * sim) { - pthread_mutex_lock(&m_lock); bool check=check_finished_kernel(); + pthread_mutex_lock(&m_lock); if(check)m_gpu->print_stats(); stream_operation op =front(); if(!op.do_operation( m_gpu )) //not ready to execute -- 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_device_runtime.h') 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