summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libcuda/cuda_runtime_api.cc6
-rw-r--r--libcuda/gpgpu_context.h9
-rw-r--r--libopencl/opencl_runtime_api.cc16
-rw-r--r--src/abstract_hardware_model.cc3
-rw-r--r--src/abstract_hardware_model.h13
-rw-r--r--src/cuda-sim/cuda-sim.cc77
-rw-r--r--src/cuda-sim/cuda-sim.h49
-rw-r--r--src/cuda-sim/cuda_device_runtime.cc7
-rw-r--r--src/cuda-sim/cuda_device_runtime.h18
-rw-r--r--src/cuda-sim/ptx_ir.cc12
-rw-r--r--src/cuda-sim/ptx_ir.h9
-rw-r--r--src/cuda-sim/ptx_parser.cc9
-rw-r--r--src/debug.h2
-rw-r--r--src/gpgpu-sim/gpu-sim.cc16
-rw-r--r--src/gpgpu-sim/shader.cc11
-rw-r--r--src/gpgpu-sim/shader.h2
-rw-r--r--src/gpgpusim_entrypoint.cc6
-rw-r--r--src/stream_manager.cc5
18 files changed, 152 insertions, 118 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index bbbaf23..59d2a60 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1076,7 +1076,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
CUctx_st* context = GPGPUSim_Context();
char *mode = getenv("PTX_SIM_MODE_FUNC");
if( mode )
- sscanf(mode,"%u", &g_ptx_sim_mode);
+ sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode));
gpgpusim_ptx_assert( !ctx->api->g_cuda_launch_stack.empty(), "empty launch stack" );
kernel_config config = ctx->api->g_cuda_launch_stack.back();
{
@@ -1092,7 +1092,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
}
struct CUstream_st *stream = config.get_stream();
printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
- g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
+ (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
kernel_info_t *grid = ctx->api->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
//do dynamic PDOM analysis for performance simulation scenario
std::string kname = grid->name();
@@ -1143,7 +1143,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
}
printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z );
- stream_operation op(grid,g_ptx_sim_mode,stream);
+ stream_operation op(grid,ctx->func_sim->g_ptx_sim_mode,stream);
g_stream_manager()->push(op);
ctx->api->g_cuda_launch_stack.pop_back();
return g_last_cudaError = cudaSuccess;
diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index a2ae7b6..3c9f87c 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -5,6 +5,7 @@
#include "../src/cuda-sim/ptx_parser.h"
#include "../src/gpgpusim_entrypoint.h"
#include "../src/cuda-sim/cuda-sim.h"
+#include "../src/cuda-sim/cuda_device_runtime.h"
class gpgpu_context {
public:
@@ -14,7 +15,10 @@ class gpgpu_context {
ptxinfo = new ptxinfo_data(this);
ptx_parser = new ptx_recognizer(this);
the_gpgpusim = new GPGPUsim_ctx(this);
- func_sim = new cuda_sim();
+ func_sim = new cuda_sim(this);
+#if (CUDART_VERSION >= 5000)
+ device_runtime = new cuda_device_runtime(this);
+#endif
}
// global list
symbol_table *g_global_allfiles_symbol_table;
@@ -25,6 +29,9 @@ class gpgpu_context {
ptx_recognizer* ptx_parser;
GPGPUsim_ctx* the_gpgpusim;
cuda_sim* func_sim;
+#if (CUDART_VERSION >= 5000)
+ cuda_device_runtime* device_runtime;
+#endif
// member function list
void cuobjdumpParseBinary(unsigned int handle);
class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 0a6eb3e..7f029c7 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -884,9 +884,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
printf("\n\n\n");
char *mode = getenv("PTX_SIM_MODE_FUNC");
if ( mode )
- sscanf(mode,"%u", &g_ptx_sim_mode);
+ sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode));
printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel '%s' (mode=%s)\n", kernel->name().c_str(),
- g_ptx_sim_mode?"functional simulation":"performance simulation");
+ (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation");
if ( !work_dim || work_dim > 3 ) return CL_INVALID_WORK_DIMENSION;
size_t _local_size[3];
if( local_work_size != NULL ) {
@@ -950,14 +950,14 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
gpgpu_t *gpu = command_queue->get_device()->the_device();
if (kernel->get_implementation()->get_ptx_version().ver() <3.0){
- gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu );
- gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu );
- gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu );
- gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
- gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+ ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu );
+ ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu );
+ ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu );
+ ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+ ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
}
kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu);
- if ( g_ptx_sim_mode )
+ if ( ctx->func_sim->g_ptx_sim_mode )
ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid );
else
gpgpu_opencl_ptx_sim_main_perf( grid );
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 248e7a5..ef09051 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -39,6 +39,7 @@
#include <sys/stat.h>
#include <sstream>
#include <iostream>
+#include "../libcuda/gpgpu_context.h"
unsigned mem_access_t::sm_next_access_uid = 0;
unsigned warp_inst_t::sm_next_uid = 0;
@@ -945,7 +946,7 @@ void simt_stack::print (FILE *fout) const
} else {
fprintf(fout," " );
}
- ptx_print_insn( stack_entry.m_pc, fout );
+ m_gpu->gpgpu_ctx->func_sim->ptx_print_insn( stack_entry.m_pc, fout );
fprintf(fout,"\n");
}
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index da29a11..8ef8376 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -347,9 +347,11 @@ public:
mutable bool volta_cache_config_set;
};
-struct core_config {
- core_config()
- {
+class core_config {
+ public:
+ core_config(gpgpu_context* ctx)
+ {
+ gpgpu_ctx = ctx;
m_valid = false;
num_shmem_bank=16;
shmem_limited_broadcast = false;
@@ -361,6 +363,8 @@ struct core_config {
bool m_valid;
unsigned warp_size;
+ // backward pointer
+ class gpgpu_context* gpgpu_ctx;
// off-chip memory request architecture parameters
int gpgpu_coalesce_arch;
@@ -934,7 +938,7 @@ public:
m_empty=true;
m_config=NULL;
}
- warp_inst_t( const core_config *config )
+ warp_inst_t( const core_config *config )
{
m_uid=0;
assert(config->warp_size<=MAX_WARP_SIZE);
@@ -1105,7 +1109,6 @@ public:
unsigned get_uid() const { return m_uid; }
unsigned get_schd_id() const { return m_scheduler_id; }
-
protected:
unsigned m_uid;
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index b3e2965..7ff869c 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -53,17 +53,10 @@ typedef void * yyscan_t;
#include "cuda_device_runtime.h"
#include "../../libcuda/gpgpu_context.h"
-int gpgpu_ptx_instruction_classification;
-void ** g_inst_classification_stat = NULL;
-void ** g_inst_op_classification_stat= NULL;
int g_debug_execution = 0;
-int g_debug_thread_uid = 0;
addr_t g_debug_pc = 0xBEEF1518;
// Output debug information to file options
-unsigned gpgpu_param_num_shaders = 0;
-
-unsigned cdp_latency[5];
void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) {
option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int,
@@ -114,8 +107,6 @@ cudaLaunchDeviceV2_init_perWarp, cudaLaunchDevicV2_perKernel>"
"7200,8000,100,12000,1600");
}
-static address_type get_converge_point(address_type pc);
-
void gpgpu_t::gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref, int dim, int readmode, int ext)
{
std::string texname(name);
@@ -225,8 +216,6 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref)
m_NameToTextureInfo.erase(texname);
}
-unsigned g_assemble_code_next_pc=0;
-std::map<unsigned,function_info*> g_pc_to_finfo;
std::vector<ptx_instruction*> function_info::s_g_pc_to_insn;
#define MAX_INST_SIZE 8 /*bytes*/
@@ -246,7 +235,7 @@ void function_info::ptx_assemble()
fflush(stdout);
std::list<ptx_instruction*>::iterator i;
- addr_t PC = g_assemble_code_next_pc; // globally unique address (across functions)
+ addr_t PC = gpgpu_ctx->func_sim->g_assemble_code_next_pc; // globally unique address (across functions)
// start function on an aligned address
for( unsigned i=0; i < (PC%MAX_INST_SIZE); i++ )
s_g_pc_to_insn.push_back((ptx_instruction*)NULL);
@@ -263,7 +252,7 @@ void function_info::ptx_assemble()
const symbol *l = pI->get_label();
labels[l->name()] = n;
} else {
- g_pc_to_finfo[PC] = this;
+ gpgpu_ctx->func_sim->g_pc_to_finfo[PC] = this;
m_instr_mem[n] = pI;
s_g_pc_to_insn.push_back(pI);
assert(pI == s_g_pc_to_insn[PC]);
@@ -278,7 +267,7 @@ void function_info::ptx_assemble()
PC += pI->inst_size();
}
}
- g_assemble_code_next_pc=PC;
+ gpgpu_ctx->func_sim->g_assemble_code_next_pc=PC;
for ( unsigned ii=0; ii < n; ii += m_instr_mem[ii]->inst_size() ) { // handle branch instructions
ptx_instruction *pI = m_instr_mem[ii];
if ( pI->get_opcode() == BRA_OP || pI->get_opcode() == BREAKADDR_OP || pI->get_opcode() == CALLP_OP) {
@@ -503,7 +492,7 @@ void gpgpu_t::gpu_memset( size_t dst_start_addr, int c, size_t count )
}
}
-void ptx_print_insn( address_type pc, FILE *fp )
+void cuda_sim::ptx_print_insn( address_type pc, FILE *fp )
{
std::map<unsigned,function_info*>::iterator f = g_pc_to_finfo.find(pc);
if( f == g_pc_to_finfo.end() ) {
@@ -515,7 +504,7 @@ void ptx_print_insn( address_type pc, FILE *fp )
finfo->print_insn(pc,fp);
}
-std::string ptx_get_insn_str( address_type pc )
+std::string cuda_sim::ptx_get_insn_str( address_type pc )
{
std::map<unsigned,function_info*>::iterator f = g_pc_to_finfo.find(pc);
if( f == g_pc_to_finfo.end() ) {
@@ -688,8 +677,11 @@ void ptx_instruction::set_opcode_and_latency()
sscanf(gpgpu_ctx->func_sim->opcode_initiation_tensor, "%u",
&tensor_init);
sscanf(gpgpu_ctx->func_sim->cdp_latency_str, "%u,%u,%u,%u,%u",
- &cdp_latency[0],&cdp_latency[1],&cdp_latency[2],
- &cdp_latency[3],&cdp_latency[4]);
+ &gpgpu_ctx->func_sim->cdp_latency[0],
+ &gpgpu_ctx->func_sim->cdp_latency[1],
+ &gpgpu_ctx->func_sim->cdp_latency[2],
+ &gpgpu_ctx->func_sim->cdp_latency[3],
+ &gpgpu_ctx->func_sim->cdp_latency[4]);
if(!m_operands.empty()){
std::vector<operand_info>::iterator it;
@@ -1098,7 +1090,7 @@ void ptx_instruction::pre_decode()
}
// get reconvergence pc
- reconvergence_pc = get_converge_point(pc);
+ reconvergence_pc = gpgpu_ctx->func_sim->get_converge_point(pc);
m_decoded=true;
}
@@ -1462,7 +1454,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
}
template<int activate_level>
-bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc)
+bool cuda_sim::ptx_debug_exec_dump_cond(int thd_uid, addr_t pc)
{
if (g_debug_execution >= activate_level) {
// check each type of debug dump constraint to filter out dumps
@@ -1548,7 +1540,8 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
}
if ( g_debug_execution >= 6 || m_gpu->get_config().get_ptx_inst_debug_to_file()) {
- if ( (g_debug_thread_uid==0) || (get_uid() == (unsigned)g_debug_thread_uid) ) {
+ if ( (m_gpu->gpgpu_ctx->func_sim->g_debug_thread_uid==0)
+ || (get_uid() == (unsigned)(m_gpu->gpgpu_ctx->func_sim->g_debug_thread_uid)) ) {
clear_modifiedregs();
enable_debug_trace();
@@ -1620,7 +1613,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
fflush(m_gpu->get_ptx_inst_debug_file());
}
- if ( ptx_debug_exec_dump_cond<5>(get_uid(), pc) ) {
+ if ( m_gpu->gpgpu_ctx->func_sim->ptx_debug_exec_dump_cond<5>(get_uid(), pc) ) {
dim3 ctaid = get_ctaid();
dim3 tid = get_tid();
printf("%u [thd=%u][i=%u] : ctaid=(%u,%u,%u) tid=(%u,%u,%u) icount=%u [pc=%u] (%s:%u - %s) [0x%llx]\n",
@@ -1674,11 +1667,11 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
}
if ( g_debug_execution >= 6 ) {
- if ( ptx_debug_exec_dump_cond<6>(get_uid(), pc) )
+ if ( m_gpu->gpgpu_ctx->func_sim->ptx_debug_exec_dump_cond<6>(get_uid(), pc) )
dump_modifiedregs(stdout);
}
if ( g_debug_execution >= 10 ) {
- if ( ptx_debug_exec_dump_cond<10>(get_uid(), pc) )
+ if ( m_gpu->gpgpu_ctx->func_sim->ptx_debug_exec_dump_cond<10>(get_uid(), pc) )
dump_regs(stdout);
}
update_pc();
@@ -1688,7 +1681,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
if(!(this->m_functionalSimulationMode))
ptx_file_line_stats_add_exec_count(pI);
- if ( gpgpu_ptx_instruction_classification ) {
+ if ( m_gpu->gpgpu_ctx->func_sim->gpgpu_ptx_instruction_classification ) {
m_gpu->gpgpu_ctx->func_sim->init_inst_classification_stat();
unsigned space_type=0;
switch ( pI->get_space().get_type() ) {
@@ -1705,9 +1698,9 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
space_type = 0 ;
break;
}
- StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], op_classification);
- if (space_type) StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], ( int )space_type);
- StatAddSample( g_inst_op_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], (int) pI->get_opcode() );
+ StatAddSample( m_gpu->gpgpu_ctx->func_sim->g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], op_classification);
+ if (space_type) StatAddSample( m_gpu->gpgpu_ctx->func_sim->g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], ( int )space_type);
+ StatAddSample( m_gpu->gpgpu_ctx->func_sim->g_inst_op_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], (int) pI->get_opcode() );
}
if ( (m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn % 100000) == 0 ) {
dim3 ctaid = get_ctaid();
@@ -1736,7 +1729,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
}
-void set_param_gpgpu_num_shaders(int num_shaders)
+void cuda_sim::set_param_gpgpu_num_shaders(int num_shaders)
{
gpgpu_param_num_shaders = num_shaders;
}
@@ -1818,7 +1811,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
assert( max_cta_per_sm > 0 );
//unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid;
- unsigned sm_idx = hw_cta_id*gpgpu_param_num_shaders + sid;
+ unsigned sm_idx = hw_cta_id*gpu->gpgpu_ctx->func_sim->gpgpu_param_num_shaders + sid;
if ( shared_memory_lookup.find(sm_idx) == shared_memory_lookup.end() ) {
if ( g_debug_execution >= 1 ) {
@@ -1948,9 +1941,6 @@ void print_splash()
}
}
-std::set<std::string> g_globals;
-std::set<std::string> g_constants;
-
void cuda_sim::gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size )
{
printf("GPGPU-Sim PTX registering constant %s (%zu bytes) to name mapping\n", deviceName, size );
@@ -1963,7 +1953,7 @@ void cuda_sim::gpgpu_ptx_sim_register_global_variable(void *hostVar, const char
g_global_name_lookup[hostVar] = deviceName;
}
-void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu )
+void cuda_sim::gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu )
{
printf("GPGPU-Sim PTX: starting gpgpu_ptx_sim_memcpy_symbol with hostVar 0x%p\n", hostVar);
bool found_sym = false;
@@ -2033,13 +2023,9 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co
fflush(stdout);
}
-int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle)
-
extern int ptx_debug;
-bool g_cuda_launch_blocking = false;
-
-void read_sim_environment_variables()
+void cuda_sim::read_sim_environment_variables()
{
ptx_debug = 0;
g_debug_execution = 0;
@@ -2185,7 +2171,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL
cta.execute(cp_count,temp);
#if (CUDART_VERSION >= 5000)
- launch_all_device_kernels();
+ gpgpu_ctx->device_runtime->launch_all_device_kernels();
#endif
}
else
@@ -2522,14 +2508,7 @@ void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels )
clear_ptxinfo();
}
-struct rec_pts {
- gpgpu_recon_t *s_kernel_recon_points;
- int s_num_recon;
-};
-
-class std::map<function_info*,rec_pts> g_rpts;
-
-struct rec_pts find_reconvergence_points( function_info *finfo )
+struct rec_pts cuda_sim::find_reconvergence_points( function_info *finfo )
{
rec_pts tmp;
std::map<function_info*,rec_pts>::iterator r=g_rpts.find(finfo);
@@ -2568,7 +2547,7 @@ address_type get_return_pc( void *thd )
return the_thread->get_return_PC();
}
-address_type get_converge_point( address_type pc )
+address_type cuda_sim::get_converge_point( address_type pc )
{
// the branch could encode the reconvergence point and/or a bit that indicates the
// reconvergence point is the return PC on the call stack in the case the branch has
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index 25ebf7b..e4d34fe 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -36,21 +36,16 @@
#include <string>
#include"ptx_sim.h"
+class gpgpu_context;
class memory_space;
class function_info;
class symbol_table;
extern const char *g_gpgpusim_version_string;
-extern int g_ptx_sim_mode;
extern int g_debug_execution;
-extern int g_debug_thread_uid;
-extern void ** g_inst_classification_stat;
-extern void ** g_inst_op_classification_stat;
extern void print_splash();
-extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu );
-extern void read_sim_environment_variables();
extern void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels );
unsigned ptx_sim_init_thread( kernel_info_t &kernel,
class ptx_thread_info** thread_info,
@@ -65,9 +60,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
bool functionalSimulationMode = false);
const warp_inst_t *ptx_fetch_inst( address_type pc );
const struct gpgpu_ptx_sim_info* ptx_sim_kernel_info(const class function_info *kernel);
-void ptx_print_insn( address_type pc, FILE *fp );
-std::string ptx_get_insn_str( address_type pc );
-void set_param_gpgpu_num_shaders(int num_shaders);
/*!
@@ -122,11 +114,25 @@ void print_ptxinfo();
void clear_ptxinfo();
struct gpgpu_ptx_sim_info get_ptxinfo();
+class gpgpu_recon_t;
+struct rec_pts {
+ gpgpu_recon_t *s_kernel_recon_points;
+ int s_num_recon;
+};
+
+
class cuda_sim {
public:
- cuda_sim() {
+ cuda_sim( gpgpu_context* ctx ) {
g_ptx_sim_num_insn = 0;
g_ptx_kernel_count = -1; // used for classification stat collection purposes
+ gpgpu_param_num_shaders = 0;
+ g_cuda_launch_blocking = false;
+ g_inst_classification_stat = NULL;
+ g_inst_op_classification_stat= NULL;
+ g_assemble_code_next_pc=0;
+ g_debug_thread_uid = 0;
+ gpgpu_ctx = ctx;
}
//global variables
char *opcode_latency_int;
@@ -147,6 +153,21 @@ class cuda_sim {
int g_ptx_kernel_count; // used for classification stat collection purposes
std::map<const void*,std::string> g_global_name_lookup; // indexed by hostVar
std::map<const void*,std::string> g_const_name_lookup; // indexed by hostVar
+ int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle)
+ unsigned gpgpu_param_num_shaders;
+ class std::map<function_info*,rec_pts> g_rpts;
+ bool g_cuda_launch_blocking;
+ void ** g_inst_classification_stat;
+ void ** g_inst_op_classification_stat;
+ std::set<std::string> g_globals;
+ std::set<std::string> g_constants;
+ std::map<unsigned,function_info*> g_pc_to_finfo;
+ int gpgpu_ptx_instruction_classification;
+ unsigned cdp_latency[5];
+ unsigned g_assemble_code_next_pc;
+ int g_debug_thread_uid;
+ // backward pointer
+ class gpgpu_context* gpgpu_ctx;
//global functions
void ptx_opcocde_latency_options (option_parser_t opp);
void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false );
@@ -159,6 +180,14 @@ class cuda_sim {
gpgpu_t *gpu );
void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size );
void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size );
+ void read_sim_environment_variables();
+ void set_param_gpgpu_num_shaders(int num_shaders);
+ struct rec_pts find_reconvergence_points( function_info *finfo );
+ address_type get_converge_point( address_type pc );
+ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu );
+ void ptx_print_insn( address_type pc, FILE *fp );
+ std::string ptx_get_insn_str( address_type pc );
+ template<int activate_level> bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc);
};
#endif
diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc
index be8369f..354fa79 100644
--- a/src/cuda-sim/cuda_device_runtime.cc
+++ b/src/cuda-sim/cuda_device_runtime.cc
@@ -20,6 +20,7 @@ unsigned long long g_max_total_param_size = 0;
#include "../stream_manager.h"
#include "../gpgpusim_entrypoint.h"
#include "cuda_device_runtime.h"
+#include "../../libcuda/gpgpu_context.h"
#define DEV_RUNTIME_REPORT(a) \
if( g_debug_execution ) { \
@@ -318,17 +319,17 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_
}
-void launch_one_device_kernel() {
+void cuda_device_runtime::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);
+ stream_operation stream_op = stream_operation(op.grid, gpgpu_ctx->func_sim->g_ptx_sim_mode, op.stream);
g_stream_manager()->push(stream_op);
g_cuda_device_launch_op.pop_front();
}
}
-void launch_all_device_kernels() {
+void cuda_device_runtime::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 6dbcd71..851fed2 100644
--- a/src/cuda-sim/cuda_device_runtime.h
+++ b/src/cuda-sim/cuda_device_runtime.h
@@ -6,6 +6,20 @@
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();
+#endif
+#if (CUDART_VERSION >= 5000)
+
+class gpgpu_context;
+
+class cuda_device_runtime {
+ public:
+ cuda_device_runtime( gpgpu_context* ctx ) {
+ gpgpu_ctx = ctx;
+ }
+ // backward pointer
+ class gpgpu_context* gpgpu_ctx;
+ void launch_all_device_kernels();
+ void launch_one_device_kernel();
+};
+
#endif
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 1bd409e..c537091 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -83,8 +83,9 @@ symbol_table::symbol_table()
assert(0);
}
-symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent )
+symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx )
{
+ gpgpu_ctx = ctx;
m_scope_name = std::string(scope_name);
m_reg_allocator=0;
m_shared_next = 0;
@@ -183,7 +184,7 @@ symbol_table* symbol_table::start_inst_group() {
//previous added
assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end());
- symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this );
+ symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this, gpgpu_ctx );
sym_table->m_global_next = m_global_next;
sym_table->m_shared_next = m_shared_next;
@@ -221,7 +222,7 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio
*func_info = m_function_info_lookup[key];
prior_decl = true;
} else {
- *func_info = new function_info(entry_point);
+ *func_info = new function_info(entry_point, gpgpu_ctx);
(*func_info)->set_name(name);
(*func_info)->set_maxnt_id(0);
m_function_info_lookup[key] = *func_info;
@@ -232,7 +233,7 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio
*sym_table = m_function_symtab_lookup[key];
} else {
assert( !prior_decl );
- *sym_table = new symbol_table( "", entry_point, this );
+ *sym_table = new symbol_table( "", entry_point, this, gpgpu_ctx );
// Initial setup code to support a register represented as "_".
// This register is used when an instruction operand is
@@ -1373,8 +1374,9 @@ std::string ptx_instruction::to_string() const
unsigned function_info::sm_next_uid = 1;
-function_info::function_info(int entry_point )
+function_info::function_info(int entry_point, gpgpu_context* ctx )
{
+ gpgpu_ctx = ctx;
m_uid = sm_next_uid++;
m_entry_point = (entry_point==1)?true:false;
m_extern = (entry_point==2)?true:false;
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 1604551..babd54b 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -306,7 +306,7 @@ private:
class symbol_table {
public:
symbol_table();
- symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent );
+ symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx);
void set_name( const char *name );
const ptx_version &get_ptx_version() const;
unsigned get_sm_target() const;
@@ -348,6 +348,9 @@ public:
symbol_table* start_inst_group();
symbol_table* end_inst_group();
+ // backward pointer
+ class gpgpu_context* gpgpu_ctx;
+
private:
unsigned m_reg_allocator;
unsigned m_shared_next;
@@ -1233,7 +1236,7 @@ private:
class function_info {
public:
- function_info(int entry_point );
+ function_info(int entry_point, gpgpu_context* ctx );
const ptx_version &get_ptx_version() const { return m_symtab->get_ptx_version(); }
unsigned get_sm_target() const { return m_symtab->get_sm_target(); }
bool is_extern() const { return m_extern; }
@@ -1403,6 +1406,8 @@ public:
void set_maxnt_id(unsigned maxthreads) { maxnt_id = maxthreads;}
unsigned get_maxnt_id() { return maxnt_id;}
+ // backward pointer
+ class gpgpu_context* gpgpu_ctx;
private:
unsigned maxnt_id;
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 0139534..5a94679 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -112,7 +112,7 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename )
{
g_filename = strdup(ptx_filename);
if (g_global_allfiles_symbol_table == NULL) {
- g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL);
+ g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL, this);
ptx_parser->g_global_symbol_table = ptx_parser->g_current_symbol_table = g_global_allfiles_symbol_table;
}
/*else {
@@ -322,9 +322,6 @@ bool ptx_recognizer::check_for_duplicates( const char *identifier )
return ( s != NULL );
}
-extern std::set<std::string> g_globals;
-extern std::set<std::string> g_constants;
-
// Returns padding that needs to be inserted ahead of address to make it aligned to min(size, maxalign)
/*
* @param address the address in bytes
@@ -453,7 +450,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi
g_current_symbol_table->alloc_global( num_bits/8 + addr_pad );
}
if( g_current_symbol_table == g_global_symbol_table ) {
- g_constants.insert( identifier );
+ gpgpu_ctx->func_sim->g_constants.insert( identifier );
}
assert( g_current_symbol_table != NULL );
g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table;
@@ -471,7 +468,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi
fflush(stdout);
g_last_symbol->set_address( addr+addr_pad );
g_current_symbol_table->alloc_global( num_bits/8 + addr_pad );
- g_globals.insert( identifier );
+ gpgpu_ctx->func_sim->g_globals.insert( identifier );
assert( g_current_symbol_table != NULL );
g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table;
break;
diff --git a/src/debug.h b/src/debug.h
index 1277494..1005bd5 100644
--- a/src/debug.h
+++ b/src/debug.h
@@ -83,8 +83,6 @@ private:
unsigned m_value;
};
-extern int gpgpu_ptx_instruction_classification ;
-
class ptx_thread_info;
class ptx_instruction;
bool thread_at_brkpt( ptx_thread_info *thd_info, const class brk_pt &b );
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 9f47067..93f041a 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -499,11 +499,11 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
option_parser_register(opp, "-gpgpu_deadlock_detect", OPT_BOOL, &gpu_deadlock_detect,
"Stop the simulation at deadlock (1=on (default), 0=off)",
"1");
- option_parser_register(opp, "-gpgpu_ptx_instruction_classification", OPT_INT32,
- &gpgpu_ptx_instruction_classification,
+ option_parser_register(opp, "-gpgpu_ptx_instruction_classification", OPT_INT32,
+ &(gpgpu_ctx->func_sim->gpgpu_ptx_instruction_classification),
"if enabled will classify ptx instruction types per kernel (Max 255 kernels now)",
"0");
- option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &g_ptx_sim_mode,
+ option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &(gpgpu_ctx->func_sim->g_ptx_sim_mode),
"Select between Performance (default) or Functional simulation (1)",
"0");
option_parser_register(opp, "-gpgpu_clock_domains", OPT_CSTR, &gpgpu_clock_domains,
@@ -882,7 +882,7 @@ void gpgpu_sim::init()
gpu_sim_cycle_parition_util = 0;
reinit_clock_domains();
- set_param_gpgpu_num_shaders(m_config.num_shader());
+ gpgpu_ctx->func_sim->set_param_gpgpu_num_shaders(m_config.num_shader());
for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++)
m_cluster[i]->reinit();
m_shader_stats->new_grid();
@@ -1200,9 +1200,9 @@ void gpgpu_sim::gpu_print_stat()
spill_log_to_file (stdout, 1, gpu_sim_cycle);
insn_warp_occ_print(stdout);
}
- if ( gpgpu_ptx_instruction_classification ) {
- StatDisp( g_inst_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]);
- StatDisp( g_inst_op_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]);
+ if ( gpgpu_ctx->func_sim->gpgpu_ptx_instruction_classification ) {
+ StatDisp( gpgpu_ctx->func_sim->g_inst_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]);
+ StatDisp( gpgpu_ctx->func_sim->g_inst_op_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]);
}
#ifdef GPGPUSIM_POWER_MODEL
@@ -1753,7 +1753,7 @@ void gpgpu_sim::cycle()
#if (CUDART_VERSION >= 5000)
//launch device kernel
- launch_one_device_kernel();
+ gpgpu_ctx->device_runtime->launch_one_device_kernel();
#endif
}
}
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 69790fc..f380560 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -1025,7 +1025,7 @@ void scheduler_unit::cycle()
m_simt_stack[warp_id]->get_pdom_stack_top_info(&pc,&rpc);
SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) has valid instruction (%s)\n",
(*iter)->get_warp_id(), (*iter)->get_dynamic_warp_id(),
- ptx_get_insn_str( pc).c_str() );
+ m_shader->m_config->gpgpu_ctx->func_sim->ptx_get_insn_str( pc).c_str() );
if( pI ) {
assert(valid);
if( pc != pI->pc ) {
@@ -1084,12 +1084,11 @@ 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[5];
if(pI->m_is_cdp == 1)
- warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1];
+ warp(warp_id).m_cdp_latency = m_shader->m_config->gpgpu_ctx->func_sim->cdp_latency[pI->m_is_cdp - 1];
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();
+ warp(warp_id).m_cdp_latency = m_shader->m_config->gpgpu_ctx->func_sim->cdp_latency[pI->m_is_cdp - 1]
+ + m_shader->m_config->gpgpu_ctx->func_sim->cdp_latency[pI->m_is_cdp] * active_mask.count();
warp(warp_id).m_cdp_dummy = true;
break;
}
@@ -2690,7 +2689,7 @@ void warp_inst_t::print( FILE *fout ) const
for (unsigned j=0; j<m_config->warp_size; j++)
fprintf(fout, "%c", (active(j)?'1':'0') );
fprintf(fout, "]: ");
- ptx_print_insn( pc, fout );
+ m_config->gpgpu_ctx->func_sim->ptx_print_insn( pc, fout );
fprintf(fout, "\n");
}
void shader_core_ctx::incexecstat(warp_inst_t *&inst)
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index e0cefac..2837f1b 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -1367,7 +1367,7 @@ const char* const pipeline_stage_name_decode[] = {
class shader_core_config : public core_config
{
public:
- shader_core_config(gpgpu_context* ctx){
+ shader_core_config(gpgpu_context* ctx):core_config(ctx){
pipeline_widths_string = NULL;
gpgpu_ctx = ctx;
}
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index d9d1023..816159f 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -207,13 +207,11 @@ void exit_simulation()
fflush(stdout);
}
-extern bool g_cuda_launch_blocking;
-
gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf()
{
srand(1);
print_splash();
- read_sim_environment_variables();
+ func_sim->read_sim_environment_variables();
ptx_parser->read_parser_environment_variables();
option_parser_t opp = option_parser_create();
@@ -233,7 +231,7 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf()
GPGPUsim_ctx_ptr()->g_the_gpu_config->init();
GPGPUsim_ctx_ptr()->g_the_gpu = new gpgpu_sim(*(GPGPUsim_ctx_ptr()->g_the_gpu_config), this);
- GPGPUsim_ctx_ptr()->g_stream_manager = new stream_manager((GPGPUsim_ctx_ptr()->g_the_gpu),g_cuda_launch_blocking);
+ GPGPUsim_ctx_ptr()->g_stream_manager = new stream_manager((GPGPUsim_ctx_ptr()->g_the_gpu), func_sim->g_cuda_launch_blocking);
GPGPUsim_ctx_ptr()->g_simulation_starttime = time((time_t *)NULL);
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index e07f4e4..be3dd71 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -29,6 +29,7 @@
#include "gpgpusim_entrypoint.h"
#include "cuda-sim/cuda-sim.h"
#include "gpgpu-sim/gpu-sim.h"
+#include "../libcuda/gpgpu_context.h"
unsigned CUstream_st::sm_next_stream_uid = 0;
@@ -150,13 +151,13 @@ bool stream_operation::do_operation( gpgpu_sim *gpu )
case stream_memcpy_to_symbol:
if(g_debug_execution >= 3)
printf("memcpy to symbol\n");
- gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu);
+ gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu);
m_stream->record_next_done();
break;
case stream_memcpy_from_symbol:
if(g_debug_execution >= 3)
printf("memcpy from symbol\n");
- gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu);
+ gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu);
m_stream->record_next_done();
break;
case stream_kernel_launch:
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 06:10:58 +0000