summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libcuda/cuda_runtime_api.cc22
-rw-r--r--src/abstract_hardware_model.cc13
-rw-r--r--src/abstract_hardware_model.h36
-rw-r--r--src/cuda-sim/cuda-sim.cc7
-rw-r--r--src/cuda-sim/instructions.cc8
5 files changed, 57 insertions, 29 deletions
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<std::string, const struct cudaArray*> NameToCudaArray, std::map<std::string, const struct textureInfo*> 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<std::string, const struct cudaArray*> nameToCudaArray, std::map<std::string, const struct textureInfo*> 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<std::string, const struct cudaArray*> m_NameToCudaArray, std::map<std::string, const struct textureInfo*> m_NameToTexureInfo);
+ 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);
~kernel_info_t();
void inc_running() { m_num_cores_running++; }
@@ -275,19 +275,22 @@ public:
std::list<class ptx_thread_info *> &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<std::string,const struct cudaArray*>::const_iterator t=m_NameToCudaArray.find(texname);
- assert(t != m_NameToCudaArray.end());
- return t->second;
- }
+ {
+ 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;
- }
+ 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
@@ -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<std::string, const struct cudaArray*> m_NameToCudaArray;
std::map<std::string, const struct textureInfo*> 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<std::string, const struct cudaArray*> getNameArrayMapping() {return m_NameToCudaArray;}
std::map<std::string, const struct textureInfo*> 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<std::string, std::set<const struct textureReference*> > m_NameToTextureRef;
std::map<const struct textureReference*, std::string> m_TextureRefToName;
std::map<std::string, const struct cudaArray*> 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);
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 22:41:38 +0000