From dbc5076945d08959c0106ff580f42adf3f86d7d9 Mon Sep 17 00:00:00 2001 From: tgrogers Date: Tue, 20 Nov 2018 15:34:03 -0500 Subject: Fixing the compile in gcc 4.4.x --- src/cuda-sim/cuda-sim.cc | 2 +- src/cuda-sim/instructions.cc | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 3773f6f..ec68b5b 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1308,7 +1308,7 @@ void function_info::list_param( FILE *fout ) const void function_info::ptx_jit_config(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) { - static unsigned long counter = 0; + static unsigned long long counter = 0; std::vector< std::pair > param_data; std::vector offsets; std::vector paramIsPointer; diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 31a33c6..c85654c 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -195,7 +195,7 @@ void ptx_thread_info::print_reg_thread(char * fname) { reg_map_t reg = m_regs.back(); - typename reg_map_t::const_iterator it; + reg_map_t::const_iterator it; for (it = reg.begin(); it != reg.end(); ++it) { const std::string &name = it->first->name(); -- cgit v1.3 From 8b8e35b9e8cb7b1e18ea2faf06450f8052e48b7e Mon Sep 17 00:00:00 2001 From: Deval Shah Date: Sun, 24 Mar 2019 15:05:10 -0700 Subject: Parse bar.warp.sync as NOP operation ( which is acceptable in Stack based reconvergence ) --- src/cuda-sim/ptx.l | 1 + 1 file changed, 1 insertion(+) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index d4bf631..3232361 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -61,6 +61,7 @@ addc TC; ptx_lval.int_value = ADDC_OP; return OPCODE; and TC; ptx_lval.int_value = AND_OP; return OPCODE; andn TC; ptx_lval.int_value = ANDN_OP; return OPCODE; atom TC; ptx_lval.int_value = ATOM_OP; return OPCODE; +bar.warp TC; ptx_lval.int_value = NOP_OP; return OPCODE; bar TC; ptx_lval.int_value = BAR_OP; return OPCODE; bfe TC; ptx_lval.int_value = BFE_OP; return OPCODE; bfi TC; ptx_lval.int_value = BFI_OP; return OPCODE; -- cgit v1.3 From 73050b730b31e9f58f8c2a0bfc4b3a8cd92be7ff Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Wed, 10 Apr 2019 23:27:42 -0500 Subject: Adding texture Fix --- libcuda/cuda_runtime_api.cc | 28 ++++++++++++++++++++++++---- src/abstract_hardware_model.cc | 4 +++- src/abstract_hardware_model.h | 22 +++++++++++++++++++++- src/cuda-sim/cuda-sim.cc | 5 +++-- src/cuda-sim/cuda_device_runtime.cc | 3 ++- src/cuda-sim/instructions.cc | 5 +++-- 6 files changed, 56 insertions(+), 11 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index cf16d9b..7569b20 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1204,8 +1204,24 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li if(g_debug_execution >= 3){ announce_call(__my_func__); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaSuccess; + switch(limit) { + case 0: + *pValue=1024; + break; + case 2: + *pValue=8388608; + break; + case 3: + *pValue=2; + break; + case 4: + *pValue=2048; + break; + default: + printf("ERROR:Limit %s unimplemented \n",limit); + abort(); + } + return g_last_cudaError = cudaSuccess; } @@ -1571,7 +1587,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) #if (CUDART_VERSION >= 3000) //synchronization required for application using external libraries without explicit synchronization in the code to //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. - synchronize(); + stream->synchronize(); g_stream_manager->destroy_stream(stream); #endif return g_last_cudaError = cudaSuccess; @@ -2702,6 +2718,7 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ if(g_debug_execution >= 3){ announce_call(__my_func__); } + synchronize(); return g_last_cudaError = cudaSuccess; } @@ -3135,6 +3152,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v } return g_last_cudaError = cudaSuccess; } + + #endif #endif @@ -3313,7 +3332,8 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, announce_call(__my_func__); } function_info *entry = context->get_kernel(hostFun); - kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry); + gpgpu_t* gpu= context->get_device()->get_gpgpu(); + kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping()); if( entry == NULL ) { printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun); abort(); diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 307df40..3db87c9 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -688,7 +688,7 @@ unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry ) +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map NameToCudaArray, std::map NameToTexureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -708,6 +708,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_launch_latency = g_kernel_launch_latency; volta_cache_config_set=false; + m_NameToCudaArray = NameToCudaArray; + m_NameToTexureInfo = NameToTexureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 28a22e0..bbc31e5 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,7 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry ); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map m_NameToCudaArray, std::map m_NameToTexureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -275,6 +275,20 @@ public: std::list &active_threads() { return m_active_threads; } class memory_space *get_param_memory() { return m_param_mem; } + const struct cudaArray* get_texarray( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToCudaArray.find(texname); + assert(t != m_NameToCudaArray.end()); + return t->second; + } + + const struct textureInfo* get_texinfo( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToTexureInfo.find(texname); + assert(t != m_NameToTexureInfo.end()); + return t->second; + } + private: kernel_info_t( const kernel_info_t & ); // disable copy constructor void operator=( const kernel_info_t & ); // disable copy operator @@ -284,6 +298,9 @@ private: unsigned m_uid; static unsigned m_next_uid; + std::map m_NameToCudaArray; + std::map m_NameToTexureInfo; + dim3 m_grid_dim; dim3 m_block_dim; dim3 m_next_cta; @@ -595,6 +612,9 @@ public: const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; } FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } + std::map getNameArrayMapping() {return m_NameToCudaArray;} + std::map getNameInfoMapping() {return m_NameToTexureInfo;} + protected: const gpgpu_functional_sim_config &m_function_model_config; FILE* ptx_inst_debug_file; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index ec68b5b..a51ea3e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1507,7 +1507,8 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th std::string texname = src1.name(); gpgpu_t *gpu = thread->get_gpu(); - const struct textureInfo* texInfo = gpu->get_texinfo(texname); + kernel_info_t& k = thread->get_kernel(); + const struct textureInfo* texInfo = k.get_texinfo(texname); unsigned data_size = texInfo->texel_size; return data_size; @@ -1918,7 +1919,7 @@ kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, struct dim3 blockDim, gpgpu_t *gpu ) { - kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry); + kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping()); unsigned argcount=args.size(); unsigned argn=1; for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index b399133..917e7a8 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -198,7 +198,8 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * 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); + gpgpu_t* gpu=thread->get_gpu(); + device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry, gpu->getNameArrayMapping(), gpu->getNameInfoMapping()); 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 (" << diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index c85654c..d29fc8b 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -5260,9 +5260,10 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture gpgpu_t *gpu = thread->get_gpu(); + kernel_info_t &k = thread->get_kernel(); const struct textureReference* texref = gpu->get_texref(texname); - const struct cudaArray* cuArray = gpu->get_texarray(texname); - const struct textureInfo* texInfo = gpu->get_texinfo(texname); + const struct cudaArray* cuArray = k.get_texarray(texname); + const struct textureInfo* texInfo = k.get_texinfo(texname); const struct textureReferenceAttr* texAttr = gpu->get_texattr(texname); //assume always 2D f32 input -- cgit v1.3 From 69af057daf1999f17b81d761c7a616f67f56adbb Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Tue, 16 Apr 2019 23:36:02 -0500 Subject: Adding additional fixes for the texture bugs --- libcuda/cuda_runtime_api.cc | 22 ++++++++++++---------- src/abstract_hardware_model.cc | 13 ++++++++++--- src/abstract_hardware_model.h | 36 +++++++++++++++++++++--------------- src/cuda-sim/cuda-sim.cc | 7 +++++++ src/cuda-sim/instructions.cc | 8 +++++++- 5 files changed, 57 insertions(+), 29 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 7569b20..f6e5e6f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1121,10 +1121,10 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic *value= 0; break; case 75: - *value= 9 ; + *value= 7 ; break; case 76: - *value= 3 ; + *value= 0 ; break; case 78: *value= 0 ; //TODO: as of now, we dont support stream priorities. @@ -1199,22 +1199,23 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device) *device = g_active_device; return g_last_cudaError = cudaSuccess; } + __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit ) { if(g_debug_execution >= 3){ announce_call(__my_func__); } switch(limit) { - case 0: + case 0: // cudaLimitStackSize *pValue=1024; break; - case 2: + case 2: // cudaLimitMallocHeapSize *pValue=8388608; break; - case 3: + case 3: // cudaLimitDevRuntimeSyncDepth *pValue=2; break; - case 4: + case 4: // cudaLimitDevRuntimePendingLaunchCount *pValue=2048; break; default: @@ -1225,7 +1226,6 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li } - __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority ) { if(g_debug_execution >= 3){ @@ -1585,7 +1585,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) announce_call(__my_func__); } #if (CUDART_VERSION >= 3000) - //synchronization required for application using external libraries without explicit synchronization in the code to + //per-stream synchronization required for application using external libraries without explicit synchronization in the code to //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. stream->synchronize(); g_stream_manager->destroy_stream(stream); @@ -2714,7 +2714,6 @@ cudaError_t cudaDeviceReset ( void ) { return g_last_cudaError = cudaSuccess; } cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ - // I don't know what this should do if(g_debug_execution >= 3){ announce_call(__my_func__); } @@ -2722,7 +2721,6 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ return g_last_cudaError = cudaSuccess; } - void CUDARTAPI __cudaRegisterFunction( void **fatCubinHandle, const char *hostFun, @@ -3333,6 +3331,10 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, } function_info *entry = context->get_kernel(hostFun); gpgpu_t* gpu= context->get_device()->get_gpgpu(); + /* + Passing a snapshot of the GPU's current texture mapping to the kernel's info + as kernels should use texture bindings present at the time of their launch. + */ kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping()); if( entry == NULL ) { printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun); diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 3db87c9..820f5fd 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -192,6 +192,9 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) checkpoint_CTA_t = m_function_model_config.get_checkpoint_CTA_t(); checkpoint_insn_Y = m_function_model_config.get_checkpoint_insn_Y(); + // initialize texture mappings to empty + m_NameToTextureInfo.clear(); + m_NameToCudaArray.clear(); if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w"); @@ -688,7 +691,11 @@ unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map NameToCudaArray, std::map NameToTexureInfo) +/*A snapshot of the texture mappings needs to be stored in the kernel's info as +kernels should use the texture bindings seen at the time of launch and textures + can be bound/unbound asynchronously with respect to streams. */ + +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -708,8 +715,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_launch_latency = g_kernel_launch_latency; volta_cache_config_set=false; - m_NameToCudaArray = NameToCudaArray; - m_NameToTexureInfo = NameToTexureInfo; + m_NameToCudaArray = nameToCudaArray; + m_NameToTexureInfo = nameToTexureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index bbc31e5..201b21f 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,7 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map m_NameToCudaArray, std::map m_NameToTexureInfo); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -275,19 +275,22 @@ public: std::list &active_threads() { return m_active_threads; } class memory_space *get_param_memory() { return m_param_mem; } + + //The following functions access texture bindings present at the kernel's launch + const struct cudaArray* get_texarray( const std::string &texname ) const - { - std::map::const_iterator t=m_NameToCudaArray.find(texname); - assert(t != m_NameToCudaArray.end()); - return t->second; - } + { + std::map::const_iterator t=m_NameToCudaArray.find(texname); + assert(t != m_NameToCudaArray.end()); + return t->second; + } - const struct textureInfo* get_texinfo( const std::string &texname ) const - { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); - return t->second; - } + const struct textureInfo* get_texinfo( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToTexureInfo.find(texname); + assert(t != m_NameToTexureInfo.end()); + return t->second; + } private: kernel_info_t( const kernel_info_t & ); // disable copy constructor @@ -297,7 +300,8 @@ private: unsigned m_uid; static unsigned m_next_uid; - + + //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; std::map m_NameToTexureInfo; @@ -611,7 +615,9 @@ public: const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; } FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } - + + // These maps return the current texture mappings for the GPU at any given time. + std::map getNameArrayMapping() {return m_NameToCudaArray;} std::map getNameInfoMapping() {return m_NameToTexureInfo;} @@ -624,7 +630,7 @@ protected: class memory_space *m_surf_mem; unsigned long long m_dev_malloc; - + // These maps contain the current texture mappings for the GPU at any given time. std::map > m_NameToTextureRef; std::map m_TextureRefToName; std::map m_NameToCudaArray; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index a51ea3e..9d59411 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1507,6 +1507,13 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th std::string texname = src1.name(); gpgpu_t *gpu = thread->get_gpu(); + /* + For programs with many streams, textures can be bound and unbound + asynchronously. This means we need to use the kernel's "snapshot" of + the state of the texture mappings when it was launched (so that we + don't try to access the incorrect texture mapping if it's been updated, + or that we don't access a mapping that has been unbound). + */ kernel_info_t& k = thread->get_kernel(); const struct textureInfo* texInfo = k.get_texinfo(texname); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index d29fc8b..11001d7 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -5258,7 +5258,13 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) if (!ptx_tex_regs) ptx_tex_regs = new ptx_reg_t[4]; unsigned nelem = src2.get_vect_nelem(); thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture - + /* + For programs with many streams, textures can be bound and unbound + asynchronously. This means we need to use the kernel's "snapshot" of + the state of the texture mappings when it was launched (so that we + don't try to access the incorrect texture mapping if it's been updated, + or that we don't access a mapping that has been unbound). + */ gpgpu_t *gpu = thread->get_gpu(); kernel_info_t &k = thread->get_kernel(); const struct textureReference* texref = gpu->get_texref(texname); -- cgit v1.3 From 49c6eeff7752b5816fb08c0d6baf2b47c789bc18 Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Wed, 17 Apr 2019 22:09:17 -0500 Subject: Adding configs for device limits and fixed other misc bugs --- libcuda/cuda_runtime_api.cc | 11 +++++++---- src/abstract_hardware_model.cc | 9 ++++----- src/abstract_hardware_model.h | 17 ++++++++--------- src/cuda-sim/cuda-sim.cc | 4 ++-- src/gpgpu-sim/gpu-sim.cc | 8 ++++++++ src/gpgpu-sim/gpu-sim.h | 11 ++++++++++- 6 files changed, 39 insertions(+), 21 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f6e5e6f..86435b9 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1205,18 +1205,20 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li if(g_debug_execution >= 3){ announce_call(__my_func__); } + _cuda_device_id *dev = GPGPUSim_Init(); + const gpgpu_sim_config& config=dev->get_gpgpu()->get_config(); switch(limit) { case 0: // cudaLimitStackSize - *pValue=1024; + *pValue=config.stack_limit(); break; case 2: // cudaLimitMallocHeapSize - *pValue=8388608; + *pValue=config.heap_limit(); break; case 3: // cudaLimitDevRuntimeSyncDepth - *pValue=2; + *pValue=config.sync_depth_limit(); break; case 4: // cudaLimitDevRuntimePendingLaunchCount - *pValue=2048; + *pValue=config.pending_launch_count_limit(); break; default: printf("ERROR:Limit %s unimplemented \n",limit); @@ -2717,6 +2719,7 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ if(g_debug_execution >= 3){ announce_call(__my_func__); } + //Blocks until the device has completed all preceding requested tasks synchronize(); return g_last_cudaError = cudaSuccess; } diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 820f5fd..cebdb25 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -88,8 +88,6 @@ void checkpoint::load_global_mem(class memory_space *temp_mem, char * f1name) fclose ( fp2 ); } - - void checkpoint::store_global_mem(class memory_space * mem, char *fname, char * format) { @@ -195,6 +193,8 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) // initialize texture mappings to empty m_NameToTextureInfo.clear(); m_NameToCudaArray.clear(); + m_TextureRefToName.clear(); + m_NameToAttribute.clear(); if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w"); @@ -694,8 +694,7 @@ unsigned kernel_info_t::m_next_uid = 1; /*A snapshot of the texture mappings needs to be stored in the kernel's info as kernels should use the texture bindings seen at the time of launch and textures can be bound/unbound asynchronously with respect to streams. */ - -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo) +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -716,7 +715,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * volta_cache_config_set=false; m_NameToCudaArray = nameToCudaArray; - m_NameToTexureInfo = nameToTexureInfo; + m_NameToTextureInfo = nameToTextureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 201b21f..e8716ab 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,7 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -287,8 +287,8 @@ public: const struct textureInfo* get_texinfo( const std::string &texname ) const { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); + std::map::const_iterator t=m_NameToTextureInfo.find(texname); + assert(t != m_NameToTextureInfo.end()); return t->second; } @@ -303,7 +303,7 @@ private: //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; - std::map m_NameToTexureInfo; + std::map m_NameToTextureInfo; dim3 m_grid_dim; dim3 m_block_dim; @@ -601,8 +601,8 @@ public: const struct textureInfo* get_texinfo( const std::string &texname ) const { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); + std::map::const_iterator t=m_NameToTextureInfo.find(texname); + assert(t != m_NameToTextureInfo.end()); return t->second; } @@ -617,9 +617,8 @@ public: FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } // These maps return the current texture mappings for the GPU at any given time. - std::map getNameArrayMapping() {return m_NameToCudaArray;} - std::map getNameInfoMapping() {return m_NameToTexureInfo;} + std::map getNameInfoMapping() {return m_NameToTextureInfo;} protected: const gpgpu_functional_sim_config &m_function_model_config; @@ -634,7 +633,7 @@ protected: std::map > m_NameToTextureRef; std::map m_TextureRefToName; std::map m_NameToCudaArray; - std::map m_NameToTexureInfo; + std::map m_NameToTextureInfo; std::map m_NameToAttribute; }; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 9d59411..f7bb9cc 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -218,7 +218,7 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te texInfo->Ty_numbits = intLOGB2(Ty); texInfo->texel_size = texel_size; texInfo->texel_size_numbits = intLOGB2(texel_size); - m_NameToTexureInfo[texname] = texInfo; + m_NameToTextureInfo[texname] = texInfo; } void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) @@ -226,7 +226,7 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) //assumes bind-use-unbind-bind-use-unbind pattern std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref); m_NameToCudaArray.erase(texname); - m_NameToTexureInfo.erase(texname); + m_NameToTextureInfo.erase(texname); } unsigned g_assemble_code_next_pc=0; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index ec570bf..2a862a4 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -532,6 +532,14 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-visualizer_zlevel", OPT_INT32, &g_visualizer_zlevel, "Compression level of the visualizer output log (0=no comp, 9=highest)", "6"); + option_parser_register(opp, "-gpgpu_stack_size_limit", OPT_INT32, &stack_size_limit, + "GPU thread stack size", "1024" ); + option_parser_register(opp, "-gpgpu_heap_size_limit", OPT_INT32, &heap_size_limit, + "GPU malloc heap size ", "8388608" ); + option_parser_register(opp, "-gpgpu_runtime_sync_depth_limit", OPT_INT32, &runtime_sync_depth_limit, + "GPU device runtime synchronize depth", "2" ); + option_parser_register(opp, "-gpgpu_runtime_pending_launch_count_limit", OPT_INT32, &runtime_pending_launch_count_limit, + "GPU device runtime pending launch count", "2048" ); option_parser_register(opp, "-trace_enabled", OPT_BOOL, &Trace::enabled, "Turn on traces", "0"); diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 6ce5524..11ffe49 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -337,6 +337,11 @@ public: unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; } unsigned checkpoint_option; + size_t stack_limit() const {return stack_size_limit; } + size_t heap_limit() const {return heap_size_limit; } + size_t sync_depth_limit() const {return runtime_sync_depth_limit; } + size_t pending_launch_count_limit() const {return runtime_pending_launch_count_limit;} + private: void init_clock_domains(void ); @@ -377,7 +382,11 @@ private: int gpu_stat_sample_freq; int gpu_runtime_stat_flag; - + // Device Limits + size_t stack_size_limit; + size_t heap_size_limit; + size_t runtime_sync_depth_limit; + size_t runtime_pending_launch_count_limit; unsigned long long liveness_message_freq; -- cgit v1.3