summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libcuda/cuda_runtime_api.cc28
-rw-r--r--src/abstract_hardware_model.cc4
-rw-r--r--src/abstract_hardware_model.h22
-rw-r--r--src/cuda-sim/cuda-sim.cc5
-rw-r--r--src/cuda-sim/cuda_device_runtime.cc3
-rw-r--r--src/cuda-sim/instructions.cc5
6 files changed, 56 insertions, 11 deletions
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<std::string, const struct cudaArray*> NameToCudaArray, std::map<std::string, const struct textureInfo*> 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<std::string, const struct cudaArray*> m_NameToCudaArray, std::map<std::string, const struct textureInfo*> m_NameToTexureInfo);
~kernel_info_t();
void inc_running() { m_num_cores_running++; }
@@ -275,6 +275,20 @@ public:
std::list<class ptx_thread_info *> &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<std::string,const struct cudaArray*>::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<std::string, const struct textureInfo*>::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<std::string, const struct cudaArray*> m_NameToCudaArray;
+ std::map<std::string, const struct textureInfo*> 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<std::string, const struct cudaArray*> getNameArrayMapping() {return m_NameToCudaArray;}
+ std::map<std::string, const struct textureInfo*> 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
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 01:23:44 +0000