summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--CHANGES1
-rw-r--r--configs/tested-cfgs/SM2_GTX480/gpgpusim.config6
-rw-r--r--configs/tested-cfgs/SM6_TITANX/gpgpusim.config8
-rw-r--r--configs/tested-cfgs/SM7_QV100/gpgpusim.config4
-rw-r--r--configs/tested-cfgs/SM7_TITANV/gpgpusim.config6
-rw-r--r--libcuda/cuda_runtime_api.cc12
-rw-r--r--libopencl/opencl_runtime_api.cc54
-rw-r--r--setup_environment3
-rw-r--r--src/cuda-sim/cuda-sim.cc34
-rw-r--r--src/cuda-sim/ptx.y18
-rw-r--r--src/gpgpu-sim/shader.cc9
-rw-r--r--src/stream_manager.cc28
-rw-r--r--src/stream_manager.h73
13 files changed, 181 insertions, 75 deletions
diff --git a/CHANGES b/CHANGES
index 2a4222b..31ae672 100644
--- a/CHANGES
+++ b/CHANGES
@@ -37,6 +37,7 @@ Version 4.0.0 (development branch) versus 3.2.3
3- Addig new system stats: gpu occupancy, L2BW, etc
-Library:
1 Enabled CUTLASS Library on GPGPU-Sim
+2 Enabled CUDA 10
-Regression:
1- Added TensorCore Regression Kernel
-Configs:
diff --git a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
index 4a7a3c3..6f088ea 100644
--- a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
+++ b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
@@ -46,9 +46,9 @@
-gpgpu_num_dp_units 0
# Instruction latencies and initiation intervals
-# "ADD,MAX,MUL,MAD,DIV"
--ptx_opcode_latency_int 4,13,4,5,145
--ptx_opcode_initiation_int 1,2,2,1,8
+# "ADD,MAX,MUL,MAD,DIV,SHFL"
+-ptx_opcode_latency_int 4,13,4,5,145,32
+-ptx_opcode_initiation_int 1,2,2,1,8,4
-ptx_opcode_latency_fp 4,13,4,5,39
-ptx_opcode_initiation_fp 1,2,1,1,4
-ptx_opcode_latency_dp 8,19,8,8,330
diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
index e6d8f1d..391269e 100644
--- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
+++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
@@ -56,11 +56,11 @@
# Instruction latencies and initiation intervals
-# "ADD,MAX,MUL,MAD,DIV"
+# "ADD,MAX,MUL,MAD,DIV,SHFL"
# All Div operations are executed on SFU unit
# Throughput (initiation latency) are adopted from CUDA SDK document V8, section 5.4.1, Table 2
--ptx_opcode_latency_int 4,13,4,5,145
--ptx_opcode_initiation_int 1,1,1,1,4
+-ptx_opcode_latency_int 4,13,4,5,145,32
+-ptx_opcode_initiation_int 1,1,1,1,4,4
-ptx_opcode_latency_fp 4,13,4,5,39
-ptx_opcode_initiation_fp 1,2,1,1,4
-ptx_opcode_latency_dp 8,19,8,8,330
@@ -136,7 +136,7 @@
# interconnection
-network_mode 1
--inter_config_file config_fermi_islip.icnt
+-inter_config_file config_pascal_islip.icnt
# memory partition latency config
-rop_latency 120
diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config
index 23a57fa..91ff0b4 100644
--- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config
+++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config
@@ -69,8 +69,8 @@
# All Div operations are executed on SFU unit
# Throughput (initiation latency) are adopted from
# http://on-demand.gputechconf.com/gtc/2018/presentation/s8122-dissecting-the-volta-gpu-architecture-through-microbenchmarking.pdf
--ptx_opcode_latency_int 4,13,4,5,145
--ptx_opcode_initiation_int 2,2,2,2,8
+-ptx_opcode_latency_int 4,13,4,5,145,21
+-ptx_opcode_initiation_int 2,2,2,2,8,4
-ptx_opcode_latency_fp 4,13,4,5,39
-ptx_opcode_initiation_fp 2,2,2,2,4
-ptx_opcode_latency_dp 8,19,8,8,330
diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
index 0339b0d..a77ab74 100644
--- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
+++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
@@ -67,10 +67,10 @@
# Instruction latencies and initiation intervals
# "ADD,MAX,MUL,MAD,DIV"
# All Div operations are executed on SFU unit
-# Throughput (initiation latency) are adopted from
+# Throughput (initiation latency except shfl) are adopted from
# http://on-demand.gputechconf.com/gtc/2018/presentation/s8122-dissecting-the-volta-gpu-architecture-through-microbenchmarking.pdf
--ptx_opcode_latency_int 4,13,4,5,145
--ptx_opcode_initiation_int 2,2,2,2,8
+-ptx_opcode_latency_int 4,13,4,5,145,32
+-ptx_opcode_initiation_int 2,2,2,2,8,4
-ptx_opcode_latency_fp 4,13,4,5,39
-ptx_opcode_initiation_fp 2,2,2,2,4
-ptx_opcode_latency_dp 8,19,8,8,330
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index cc01f12..12f9636 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -3772,6 +3772,18 @@ cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr,
return cudaFuncGetAttributesInternal(attr, hostFun);
}
+cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, int flags)
+{
+ CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync);
+ g_timer_events[e->get_uid()] = e;
+#if CUDART_VERSION >= 3000
+ *event = e;
+#else
+ *event = e->get_uid();
+#endif
+ return g_last_cudaError = cudaSuccess;
+}
+
cudaError_t CUDARTAPI cudaDriverGetVersion(int *driverVersion) {
if (g_debug_execution >= 3) {
announce_call(__my_func__);
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index b032c05..30a50fc 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -264,15 +264,27 @@ void _cl_kernel::SetKernelArg(
cl_int _cl_kernel::bind_args( gpgpu_ptx_sim_arg_list_t &arg_list )
{
+ size_t offset = 0;
+
assert( arg_list.empty() );
unsigned k=0;
std::map<unsigned, arg_info>::iterator i;
for( i = m_args.begin(); i!=m_args.end(); i++ ) {
if( i->first != k )
return CL_INVALID_KERNEL_ARGS;
+
arg_info arg = i->second;
- gpgpu_ptx_sim_arg param( arg.m_arg_value, arg.m_arg_size, 0);
+ const symbol *sym = m_kernel_impl->get_arg(i->first);
+ const type_info_key &t = sym->type()->get_key();
+
+ int align = (t.get_alignment_spec() == -1) ? arg.m_arg_size : t.get_alignment_spec();
+ if( offset % align )
+ offset += (align - (offset % align));
+
+ gpgpu_ptx_sim_arg param( arg.m_arg_value, arg.m_arg_size, offset );
arg_list.push_front( param );
+
+ offset += arg.m_arg_size;
k++;
}
return CL_SUCCESS;
@@ -957,6 +969,17 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
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);
+
+ //do dynamic PDOM analysis for performance simulation scenario
+ std::string kname = grid->name();
+ function_info *kernel_func_info = grid->entry();
+ if (kernel_func_info->is_pdom_set()) {
+ printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() );
+ } else {
+ printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() );
+ kernel_func_info->do_pdom();
+ kernel_func_info->set_pdom();
+ }
if ( ctx->func_sim->g_ptx_sim_mode )
ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid );
else
@@ -1263,6 +1286,35 @@ clGetProgramInfo(cl_program program,
}
extern CL_API_ENTRY cl_int CL_API_CALL
+clGetProgramBuildInfo (cl_program program,
+ cl_device_id device,
+ cl_program_build_info param_name,
+ size_t param_value_size,
+ void * param_value,
+ size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+ char *buf = (char*)param_value;
+
+ switch( param_name ) {
+ case CL_PROGRAM_BUILD_STATUS:
+ CL_CASE( cl_build_status, CL_BUILD_SUCCESS );
+ break;
+ case CL_PROGRAM_BUILD_OPTIONS:
+ case CL_PROGRAM_BUILD_LOG:
+ CL_STRING_CASE( "" );
+ break;
+ case CL_PROGRAM_BINARY_TYPE:
+ CL_CASE( cl_program_binary_type, CL_PROGRAM_BINARY_TYPE_EXECUTABLE );
+ break;
+ default:
+ return CL_INVALID_VALUE;
+ break;
+ }
+
+ return CL_SUCCESS;
+}
+
+extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueCopyBuffer(cl_command_queue command_queue,
cl_mem src_buffer,
cl_mem dst_buffer,
diff --git a/setup_environment b/setup_environment
index ca60d6b..daf7fab 100644
--- a/setup_environment
+++ b/setup_environment
@@ -50,7 +50,8 @@ CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {pri
export CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
if [ $CUDA_VERSION_NUMBER -gt 10100 -o $CUDA_VERSION_NUMBER -lt 2030 ]; then
echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)";
- return
+ echo $CUDA_VERSION_NUMBER
+ return
fi
if [ $CUDA_VERSION_NUMBER -ge 6000 ]; then
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 39e2b7e..75dd3c8 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -57,11 +57,11 @@ int g_debug_execution = 0;
// Output debug information to file options
void cuda_sim::ptx_opcocde_latency_options(option_parser_t opp) {
- option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR,
- &opcode_latency_int,
- "Opcode latencies for integers <ADD,MAX,MUL,MAD,DIV>"
- "Default 1,1,19,25,145",
- "1,1,19,25,145");
+ option_parser_register(
+ opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int,
+ "Opcode latencies for integers <ADD,MAX,MUL,MAD,DIV,SHFL>"
+ "Default 1,1,19,25,145,32",
+ "1,1,19,25,145,32");
option_parser_register(opp, "-ptx_opcode_latency_fp", OPT_CSTR,
&opcode_latency_fp,
"Opcode latencies for single precision floating "
@@ -86,9 +86,9 @@ void cuda_sim::ptx_opcocde_latency_options(option_parser_t opp) {
"64");
option_parser_register(
opp, "-ptx_opcode_initiation_int", OPT_CSTR, &opcode_initiation_int,
- "Opcode initiation intervals for integers <ADD,MAX,MUL,MAD,DIV>"
- "Default 1,1,4,4,32",
- "1,1,4,4,32");
+ "Opcode initiation intervals for integers <ADD,MAX,MUL,MAD,DIV,SHFL>"
+ "Default 1,1,4,4,32,4",
+ "1,1,4,4,32,4");
option_parser_register(opp, "-ptx_opcode_initiation_fp", OPT_CSTR,
&opcode_initiation_fp,
"Opcode initiation intervals for single precision "
@@ -682,12 +682,12 @@ void ptx_instruction::set_bar_type() {
}
void ptx_instruction::set_opcode_and_latency() {
- unsigned int_latency[5];
+ unsigned int_latency[6];
unsigned fp_latency[5];
unsigned dp_latency[5];
unsigned sfu_latency;
unsigned tensor_latency;
- unsigned int_init[5];
+ unsigned int_init[6];
unsigned fp_init[5];
unsigned dp_init[5];
unsigned sfu_init;
@@ -698,10 +698,11 @@ void ptx_instruction::set_opcode_and_latency() {
* [2] MUL
* [3] MAD
* [4] DIV
+ * [5] SHFL
*/
- sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u",
+ sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u,%u",
&int_latency[0], &int_latency[1], &int_latency[2], &int_latency[3],
- &int_latency[4]);
+ &int_latency[4], &int_latency[5]);
sscanf(gpgpu_ctx->func_sim->opcode_latency_fp, "%u,%u,%u,%u,%u",
&fp_latency[0], &fp_latency[1], &fp_latency[2], &fp_latency[3],
&fp_latency[4]);
@@ -710,8 +711,9 @@ void ptx_instruction::set_opcode_and_latency() {
&dp_latency[4]);
sscanf(gpgpu_ctx->func_sim->opcode_latency_sfu, "%u", &sfu_latency);
sscanf(gpgpu_ctx->func_sim->opcode_latency_tensor, "%u", &tensor_latency);
- sscanf(gpgpu_ctx->func_sim->opcode_initiation_int, "%u,%u,%u,%u,%u",
- &int_init[0], &int_init[1], &int_init[2], &int_init[3], &int_init[4]);
+ sscanf(gpgpu_ctx->func_sim->opcode_initiation_int, "%u,%u,%u,%u,%u,%u",
+ &int_init[0], &int_init[1], &int_init[2], &int_init[3], &int_init[4],
+ &int_init[5]);
sscanf(gpgpu_ctx->func_sim->opcode_initiation_fp, "%u,%u,%u,%u,%u",
&fp_init[0], &fp_init[1], &fp_init[2], &fp_init[3], &fp_init[4]);
sscanf(gpgpu_ctx->func_sim->opcode_initiation_dp, "%u,%u,%u,%u,%u",
@@ -940,8 +942,8 @@ void ptx_instruction::set_opcode_and_latency() {
op = TENSOR_CORE_OP;
break;
case SHFL_OP:
- latency = 4;
- initiation_interval = 4;
+ latency = int_latency[5];
+ initiation_interval = int_init[5];
break;
default:
break;
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index a01c3c6..b38f783 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -298,6 +298,7 @@ ptr_align_spec: ALIGN_DIRECTIVE INT_OPERAND
statement_block: LEFT_BRACE statement_list RIGHT_BRACE
statement_list: directive_statement { recognizer->add_directive(); }
+ | statement_list prototype_block {printf("Prototype statement detected. WARNING: this is not supported yet on GPGPU-SIM\n"); }
| instruction_statement { recognizer->add_instruction(); }
| statement_list directive_statement { recognizer->add_directive(); }
| statement_list instruction_statement { recognizer->add_instruction(); }
@@ -414,6 +415,23 @@ initializer_list: LEFT_BRACE literal_list RIGHT_BRACE { recognizer->add_array_in
literal_list: literal_operand
| literal_list COMMA literal_operand;
+// TODO: This is currently hardcoded to handle and ignore one specific case
+// that all prototype statements follow in the PTX from Pytorch. As a
+// workaround, this parses and ignores both the prototype declaration
+// and calling of the prototype (which conveniently comes right after the
+// declaration for all cases.) This should be changed to handle both
+// declaring the prototype, and actually calling it.
+prototype_block: prototype_decl prototype_call
+
+prototype_decl: IDENTIFIER COLON CALLPROTOTYPE_DIRECTIVE LEFT_PAREN prototype_param RIGHT_PAREN IDENTIFIER LEFT_PAREN prototype_param RIGHT_PAREN SEMI_COLON
+
+prototype_call: OPCODE LEFT_PAREN IDENTIFIER RIGHT_PAREN COMMA operand COMMA LEFT_PAREN IDENTIFIER RIGHT_PAREN COMMA IDENTIFIER SEMI_COLON
+ | OPCODE IDENTIFIER COMMA LEFT_PAREN IDENTIFIER RIGHT_PAREN COMMA IDENTIFIER SEMI_COLON
+
+prototype_param: /* empty */
+ | PARAM_DIRECTIVE B64_TYPE IDENTIFIER
+ | PARAM_DIRECTIVE B32_TYPE IDENTIFIER
+
instruction_statement: instruction SEMI_COLON
| IDENTIFIER COLON { recognizer->add_label($1); }
| pred_spec instruction SEMI_COLON;
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index a638f5c..b596c0d 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -3140,7 +3140,7 @@ unsigned int shader_core_config::max_cta(const kernel_info_t &k) const {
void shader_core_config::set_pipeline_latency() {
// calculate the max latency based on the input
- unsigned int_latency[5];
+ unsigned int_latency[6];
unsigned fp_latency[5];
unsigned dp_latency[5];
unsigned sfu_latency;
@@ -3152,10 +3152,11 @@ void shader_core_config::set_pipeline_latency() {
* [2] MUL
* [3] MAD
* [4] DIV
+ * [5] SHFL
*/
- sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u",
+ sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u,%u",
&int_latency[0], &int_latency[1], &int_latency[2], &int_latency[3],
- &int_latency[4]);
+ &int_latency[4], &int_latency[5]);
sscanf(gpgpu_ctx->func_sim->opcode_latency_fp, "%u,%u,%u,%u,%u",
&fp_latency[0], &fp_latency[1], &fp_latency[2], &fp_latency[3],
&fp_latency[4]);
@@ -3170,7 +3171,7 @@ void shader_core_config::set_pipeline_latency() {
max_sfu_latency = std::max(dp_latency[4], sfu_latency);
// assume that the max operation has the max latency
max_sp_latency = fp_latency[1];
- max_int_latency = int_latency[1];
+ max_int_latency = std::max(int_latency[1], int_latency[5]);
max_dp_latency = dp_latency[1];
max_tensor_core_latency = tensor_latency;
}
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index 48a1edc..e99bf87 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -181,13 +181,17 @@ bool stream_operation::do_operation(gpgpu_sim *gpu) {
m_event->update(gpu->gpu_tot_sim_cycle, wallclock);
m_stream->record_next_done();
} break;
- case stream_wait_event: {
+ case stream_wait_event:
// only allows next op to go if event is done
// otherwise stays in the stream queue
printf("stream wait event processing...\n");
- if (m_event->done()) printf("stream wait event done\n");
- m_stream->record_next_done();
- } break;
+ if (m_event->num_updates() >= m_cnt) {
+ printf("stream wait event done\n");
+ m_stream->record_next_done();
+ } else {
+ return false;
+ }
+ break;
default:
abort();
}
@@ -231,6 +235,7 @@ stream_manager::stream_manager(gpgpu_sim *gpu, bool cuda_launch_blocking) {
m_service_stream_zero = false;
m_cuda_launch_blocking = cuda_launch_blocking;
pthread_mutex_init(&m_lock, NULL);
+ m_last_stream = m_streams.begin();
}
bool stream_manager::operation(bool *sim) {
@@ -326,10 +331,18 @@ stream_operation stream_manager::front() {
m_service_stream_zero = false;
}
}
-
if (!m_service_stream_zero) {
- std::list<struct CUstream_st *>::iterator s;
- for (s = m_streams.begin(); s != m_streams.end(); s++) {
+ std::list<struct CUstream_st *>::iterator s = m_last_stream;
+ if (m_last_stream == m_streams.end()) {
+ s = m_streams.begin();
+ } else {
+ s++;
+ }
+ for (size_t ii = 0; ii < m_streams.size(); ii++, s++) {
+ if (s == m_streams.end()) {
+ s = m_streams.begin();
+ }
+ m_last_stream = s;
CUstream_st *stream = *s;
if (!stream->busy() && !stream->empty()) {
result = stream->next();
@@ -364,6 +377,7 @@ void stream_manager::destroy_stream(CUstream_st *stream) {
}
}
delete stream;
+ m_last_stream = m_streams.begin();
pthread_mutex_unlock(&m_lock);
}
diff --git a/src/stream_manager.h b/src/stream_manager.h
index d543e68..afcbb0e 100644
--- a/src/stream_manager.h
+++ b/src/stream_manager.h
@@ -44,6 +44,43 @@
// unsigned m_pending_streams;
//};
+struct CUevent_st {
+ public:
+ CUevent_st(bool blocking) {
+ m_uid = ++m_next_event_uid;
+ m_blocking = blocking;
+ m_updates = 0;
+ m_wallclock = 0;
+ m_gpu_tot_sim_cycle = 0;
+ m_issued = 0;
+ m_done = false;
+ }
+ void update(double cycle, time_t clk) {
+ m_updates++;
+ m_wallclock = clk;
+ m_gpu_tot_sim_cycle = cycle;
+ m_done = true;
+ }
+ // void set_done() { assert(!m_done); m_done=true; }
+ int get_uid() const { return m_uid; }
+ unsigned num_updates() const { return m_updates; }
+ bool done() const { return m_updates == m_issued; }
+ time_t clock() const { return m_wallclock; }
+ void issue() { m_issued++; }
+ unsigned int num_issued() const { return m_issued; }
+
+ private:
+ int m_uid;
+ bool m_blocking;
+ bool m_done;
+ int m_updates;
+ unsigned int m_issued;
+ time_t m_wallclock;
+ double m_gpu_tot_sim_cycle;
+
+ static int m_next_event_uid;
+};
+
enum stream_operation_type {
stream_no_op,
stream_memcpy_host_to_device,
@@ -106,6 +143,7 @@ class stream_operation {
m_kernel = NULL;
m_type = stream_wait_event;
m_event = e;
+ m_cnt = m_event->num_issued();
m_stream = stream;
m_done = false;
}
@@ -184,40 +222,6 @@ class stream_operation {
kernel_info_t *m_kernel;
struct CUevent_st *m_event;
};
-
-struct CUevent_st {
- public:
- CUevent_st(bool blocking) {
- m_uid = ++m_next_event_uid;
- m_blocking = blocking;
- m_updates = 0;
- m_wallclock = 0;
- m_gpu_tot_sim_cycle = 0;
- m_done = false;
- }
- void update(double cycle, time_t clk) {
- m_updates++;
- m_wallclock = clk;
- m_gpu_tot_sim_cycle = cycle;
- m_done = true;
- }
- // void set_done() { assert(!m_done); m_done=true; }
- int get_uid() const { return m_uid; }
- unsigned num_updates() const { return m_updates; }
- bool done() const { return m_done; }
- time_t clock() const { return m_wallclock; }
-
- private:
- int m_uid;
- bool m_blocking;
- bool m_done;
- int m_updates;
- time_t m_wallclock;
- double m_gpu_tot_sim_cycle;
-
- static int m_next_event_uid;
-};
-
struct CUstream_st {
public:
CUstream_st();
@@ -272,6 +276,7 @@ class stream_manager {
CUstream_st m_stream_zero;
bool m_service_stream_zero;
pthread_mutex_t m_lock;
+ std::list<struct CUstream_st *>::iterator m_last_stream;
};
#endif
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:44:10 +0000