Main Page - Revision history

Gpgpu-sim-admin at 22:53, 13 June 2017

2017-06-13T22:53:03Z

← Older revision		Revision as of 22:53, 13 June 2017
Line 30:		Line 30:
	To help reviewers you should indicate the version of GPGPU-Sim you used (e.g., "GPGPU-Sim version 3.1.0", "GPGPU-Sim version 3.0.2", "GPGPU-Sim version 2.1.2b", etc...).		To help reviewers you should indicate the version of GPGPU-Sim you used (e.g., "GPGPU-Sim version 3.1.0", "GPGPU-Sim version 3.0.2", "GPGPU-Sim version 2.1.2b", etc...).

	The GPGPU-Sim 3.x source is available under a BSD style copyright from ~~the UBC ECE department's~~ [https://~~dev~~.~~ece.ubc.ca/projects~~/gpgpu-sim~~/ GIT server~~].		The GPGPU-Sim 3.x source is available under a BSD style copyright from [https://github.com/gpgpu-sim GitHub].

	GPGPU-Sim version 3.1.0 running PTXPlus has a correlation of 98.3% and 97.3% versus GT200 and Fermi hardware on the [https://www.cs.virginia.edu/~skadron/wiki/rodinia/index.php/Main_Page RODINIA benchmark suite] with scaled down problem sizes (see <xr id="fig:correl_gt200"/> and <xr id="fig:correl_fermi"/>).		GPGPU-Sim version 3.1.0 running PTXPlus has a correlation of 98.3% and 97.3% versus GT200 and Fermi hardware on the [https://www.cs.virginia.edu/~skadron/wiki/rodinia/index.php/Main_Page RODINIA benchmark suite] with scaled down problem sizes (see <xr id="fig:correl_gt200"/> and <xr id="fig:correl_fermi"/>).

Gpgpu-sim-admin at 19:47, 13 June 2017

2017-06-13T19:47:13Z

← Older revision		Revision as of 19:47, 13 June 2017
Line 34:		Line 34:
	GPGPU-Sim version 3.1.0 running PTXPlus has a correlation of 98.3% and 97.3% versus GT200 and Fermi hardware on the [https://www.cs.virginia.edu/~skadron/wiki/rodinia/index.php/Main_Page RODINIA benchmark suite] with scaled down problem sizes (see <xr id="fig:correl_gt200"/> and <xr id="fig:correl_fermi"/>).		GPGPU-Sim version 3.1.0 running PTXPlus has a correlation of 98.3% and 97.3% versus GT200 and Fermi hardware on the [https://www.cs.virginia.edu/~skadron/wiki/rodinia/index.php/Main_Page RODINIA benchmark suite] with scaled down problem sizes (see <xr id="fig:correl_gt200"/> and <xr id="fig:correl_fermi"/>).

	Please submit bug reports through the [~~http~~://~~www~~.gpgpu-sim~~.org~~/~~bugs~~/ GPGPU-Sim Bug Tracking System]. If you have further questions after reading the manual and searching the bugs database, you may want to sign up to the [http://groups.google.com/group/gpgpu-sim/ GPGPU-Sim Google Group].		Please submit bug reports through the [https://github.com/gpgpu-sim/gpgpu-sim_distribution/issues GPGPU-Sim Bug Tracking System]. If you have further questions after reading the manual and searching the bugs database, you may want to sign up to the [http://groups.google.com/group/gpgpu-sim/ GPGPU-Sim Google Group].

	Besides this manual, you may also want to consult the [http://www.gpgpu-sim.org/isca2012-tutorial/GPGPU-Sim-Tutorial-ISCA2012.pdf slides] from our [http://www.gpgpu-sim.org/isca2012-tutorial/ tutorial at ISCA 2012]		Besides this manual, you may also want to consult the [http://www.gpgpu-sim.org/isca2012-tutorial/GPGPU-Sim-Tutorial-ISCA2012.pdf slides] from our [http://www.gpgpu-sim.org/isca2012-tutorial/ tutorial at ISCA 2012]
Line 1,296:		Line 1,296:
	This section describes strategies for debugging errors that crash or deadlock GPGPU-Sim while running in performance simulation mode.		This section describes strategies for debugging errors that crash or deadlock GPGPU-Sim while running in performance simulation mode.

	If you ran a benchmark we haven't tested and it crashed we encourage you to [~~http~~://~~www~~.gpgpu-sim~~.org~~/~~bugs~~ file a bug].		If you ran a benchmark we haven't tested and it crashed we encourage you to [https://github.com/gpgpu-sim/gpgpu-sim_distribution/issues file a bug].

	However, suppose GPGPU-Sim crashes or deadlocks after you have made your own changes? What should you do then? This section describes some generic techniques we use at UBC for debugging errors in the GPGPU-Sim timing model.		However, suppose GPGPU-Sim crashes or deadlocks after you have made your own changes? What should you do then? This section describes some generic techniques we use at UBC for debugging errors in the GPGPU-Sim timing model.

Gpgpu-sim-admin at 19:42, 13 June 2017

2017-06-13T19:42:55Z

← Older revision		Revision as of 19:42, 13 June 2017
Line 126:		Line 126:

	At the decode stage, the recent fetched instructions are decoded and stored in their corresponding entry in the I-Buffer waiting to be issued.		At the decode stage, the recent fetched instructions are decoded and stored in their corresponding entry in the I-Buffer waiting to be issued.

	~~The simulator software design for this stage is described in [[GPGPU-Sim_3.x_Manual#Fetch_and_Decode_Software_Model \| Fetch and Decode]].~~

	The simulator software design for this stage is described in [[#Fetch_and_Decode_Software_Model \| Fetch and Decode]].		The simulator software design for this stage is described in [[#Fetch_and_Decode_Software_Model \| Fetch and Decode]].
Line 133:		Line 131:
	==== Instruction Issue ====		==== Instruction Issue ====

	A second round robin arbiter chooses a warp to issue from the I-Buffer to rest of the pipeline. This round robin arbiter is decoupled from the round robin arbiter used to schedule instruction cache accesses. The issue scheduler can be configured to issue multiple instructions from the same warp per cycle. Each valid instruction (i.e. decoded and not issued) in the currently checked warp is eligible for issuing if (1) its warp is not waiting at a barrier, (2) it has valid instructions in its I-Buffer entries (valid bit is set), (3) the scoreboard check passes (see section [[~~GPGPU-Sim_3.x_Manual~~#Scoreboard \| Scoreboard]] for more details), and (4) the operand access stage of the instruction pipeline is not stalled.		A second round robin arbiter chooses a warp to issue from the I-Buffer to rest of the pipeline. This round robin arbiter is decoupled from the round robin arbiter used to schedule instruction cache accesses. The issue scheduler can be configured to issue multiple instructions from the same warp per cycle. Each valid instruction (i.e. decoded and not issued) in the currently checked warp is eligible for issuing if (1) its warp is not waiting at a barrier, (2) it has valid instructions in its I-Buffer entries (valid bit is set), (3) the scoreboard check passes (see section [[#Scoreboard \| Scoreboard]] for more details), and (4) the operand access stage of the instruction pipeline is not stalled.

	Memory instructions (Load, store, or memory barriers) are issued to the memory pipeline. For other instructions, it always prefers the SP pipe for operations that can use both SP and SFU pipelines. However, if a control hazard is detected then instructions in the I-Buffer corresponding to this warp are flushed. The warp's next pc is updated to point to the next instruction (assuming all branches as not-taken). For more information about handling control flow, refer to [[~~GPGPU-Sim_3.x_Manual~~#SIMT_Stack \| SIMT Stack]].		Memory instructions (Load, store, or memory barriers) are issued to the memory pipeline. For other instructions, it always prefers the SP pipe for operations that can use both SP and SFU pipelines. However, if a control hazard is detected then instructions in the I-Buffer corresponding to this warp are flushed. The warp's next pc is updated to point to the next instruction (assuming all branches as not-taken). For more information about handling control flow, refer to [[#SIMT_Stack \| SIMT Stack]].

	At the issue stage barrier operations are executed. Also, the SIMT stack is updated (refer to [[~~GPGPU-Sim_3.x_Manual~~#SIMT_Stack \| SIMT Stack]] for more details) and register dependencies are tracked (refer to [[~~GPGPU-Sim_3.x_Manual~~#Scoreboard \| Scoreboard]] for more details). Warps wait for barriers ("__syncthreads()") at the issue stage.		At the issue stage barrier operations are executed. Also, the SIMT stack is updated (refer to [[#SIMT_Stack \| SIMT Stack]] for more details) and register dependencies are tracked (refer to [[#Scoreboard \| Scoreboard]] for more details). Warps wait for barriers ("__syncthreads()") at the issue stage.

	==== SIMT Stack ====		==== SIMT Stack ====
Line 147:		Line 145:
	See [http://doi.acm.org/10.1145/1543753.1543756 Dynamic Warp Formation: Efficient MIMD Control Flow on SIMD Graphics Hardware] for more details.		See [http://doi.acm.org/10.1145/1543753.1543756 Dynamic Warp Formation: Efficient MIMD Control Flow on SIMD Graphics Hardware] for more details.

	Note that it is known that NVIDIA and AMD actually modify the contents of their divergence stack using special instructions. These divergence stack instructions are not exposed in PTX but are visible in the actual hardware SASS instruction set (visible using decuda or NVIDIA's cuobjdump). When the current version of GPGPU-Sim 3.x is configured to execute SASS via PTXPlus (see [[~~GPGPU-Sim_3.x_Manual~~#PTX_vs._PTXPlus\|PTX vs. PTXPlus]]) it ignores these low level instructions and instead a comparable control flow graph is created to identify immediate post-dominators. We plan to support execution of the low level branch instructions in a future version of GPGPU-Sim 3.x.		Note that it is known that NVIDIA and AMD actually modify the contents of their divergence stack using special instructions. These divergence stack instructions are not exposed in PTX but are visible in the actual hardware SASS instruction set (visible using decuda or NVIDIA's cuobjdump). When the current version of GPGPU-Sim 3.x is configured to execute SASS via PTXPlus (see [[#PTX_vs._PTXPlus \| PTX vs. PTXPlus]]) it ignores these low level instructions and instead a comparable control flow graph is created to identify immediate post-dominators. We plan to support execution of the low level branch instructions in a future version of GPGPU-Sim 3.x.

	==== Scoreboard ====		==== Scoreboard ====
Line 1,720:		Line 1,718:

	===== Fetch and Decode Software Model =====		===== Fetch and Decode Software Model =====
	This section describes the software modelling [[~~GPGPU-Sim_3.x_Manual~~#Fetch_and_Decode \| Fetch and Decode]].		This section describes the software modelling [[#Fetch_and_Decode \| Fetch and Decode]].

	The I-Buffer shown in <xr id="fig:overall_arch"/> is implemented as an array of shd_warp_t objects inside shader_core_ctx. Each shd_warp_t has a set m_ibuffer of I-Buffer entries (ibuffer_entry) holding a configurable number of instructions (the maximum allowable instructions to fetch in one cycle). Also, shd_warp_t has flags that are used by the schedulers to determine the eligibility of the warp for issue. The decoded instructions stored in an ibuffer_entry as a pointer to a warp_inst_t object. The warp_inst_t holds information about the type of the operation of this instruction and the operands used.		The I-Buffer shown in <xr id="fig:overall_arch"/> is implemented as an array of shd_warp_t objects inside shader_core_ctx. Each shd_warp_t has a set m_ibuffer of I-Buffer entries (ibuffer_entry) holding a configurable number of instructions (the maximum allowable instructions to fetch in one cycle). Also, shd_warp_t has flags that are used by the schedulers to determine the eligibility of the warp for issue. The decoded instructions stored in an ibuffer_entry as a pointer to a warp_inst_t object. The warp_inst_t holds information about the type of the operation of this instruction and the operands used.
Line 1,735:		Line 1,733:

	===== SIMT Stack Software Model =====		===== SIMT Stack Software Model =====
	For each scheduler unit there is an array of SIMT stacks. Each SIMT stack corresponds to one warp. In the scheduler_unit::cycle(), the top of the stack entry for the SIMT stack of the scheduled warp determines the issued instruction. The program counter of the top of the stack entry is normally consistent with the program counter of the next instruction in the I-Buffer that corresponds to the scheduled warp (Refer to [[~~GPGPU-Sim_3.x_Manual~~#SIMT_Stack \| SIMT Stack]]). Otherwise, in case of control hazard, they will not be matched and the instructions within the I-Buffer are flushed.		For each scheduler unit there is an array of SIMT stacks. Each SIMT stack corresponds to one warp. In the scheduler_unit::cycle(), the top of the stack entry for the SIMT stack of the scheduled warp determines the issued instruction. The program counter of the top of the stack entry is normally consistent with the program counter of the next instruction in the I-Buffer that corresponds to the scheduled warp (Refer to [[#SIMT_Stack \| SIMT Stack]]). Otherwise, in case of control hazard, they will not be matched and the instructions within the I-Buffer are flushed.

	The implementation of the SIMT stack is in the simt_stack class in shader.h. The SIMT stack is updated after each issue using this function simt_stack::update(...). This function implements the algorithm required at divergence and reconvergence points. Functional execution (refer to [[#CUDA-sim_-_Functional_Simulation_Engine\|Instruction Execution]]) is performed at the issue stage before updating the SIMT stack. This allows the issue stage to have information of the next pc of each thread, hence, to update the SIMT stack as required.		The implementation of the SIMT stack is in the simt_stack class in shader.h. The SIMT stack is updated after each issue using this function simt_stack::update(...). This function implements the algorithm required at divergence and reconvergence points. Functional execution (refer to [[#CUDA-sim_-_Functional_Simulation_Engine\|Instruction Execution]]) is performed at the issue stage before updating the SIMT stack. This allows the issue stage to have information of the next pc of each thread, hence, to update the SIMT stack as required.
Line 1,745:		Line 1,743:


	The operand collector is modeled as one stage in the main pipeline executed by the function shader_core_ctx::cycle(). This stage is represented by the shader_core_ctx::read_operands() function. Refer to [[~~GPGPU-Sim_3.x_Manual~~#ALU_Pipelines\|ALU Pipeline]] for more details about the interfaces of the operand collector.		The operand collector is modeled as one stage in the main pipeline executed by the function shader_core_ctx::cycle(). This stage is represented by the shader_core_ctx::read_operands() function. Refer to [[#ALU_Pipelines\|ALU Pipeline]] for more details about the interfaces of the operand collector.

	The class opndcoll_rfu_t models the operand collector based register file unit. It contains classes that abstracts the collector unit sets, the arbiter and the dispatch units.		The class opndcoll_rfu_t models the operand collector based register file unit. It contains classes that abstracts the collector unit sets, the arbiter and the dispatch units.
Line 1,955:		Line 1,953:
	=== Dump Pipeline ===		=== Dump Pipeline ===

	See [[~~GPGPU-Sim_3.x_Manual~~#Visualizing_Cycle_by_Cycle_Microarchitecture_Behavior]] for how this is used.		See [[#Visualizing_Cycle_by_Cycle_Microarchitecture_Behavior]] for how this is used.

	The top level dump pipeline code is implemented in <tt>gpgpu_sim::pipeline(...)</tt> in gpgpu-sim/gpu-sim.cc. It calls <tt>shader_core_ctx::display_pipeline(...)</tt> in gpgpu-sim/shader.cc for the pipeline states in each SIMT core. For memory partition states, it calls <tt>memory_partition_unit::print(...)</tt> in gpgpu-sim/l2cache.cc.		The top level dump pipeline code is implemented in <tt>gpgpu_sim::pipeline(...)</tt> in gpgpu-sim/gpu-sim.cc. It calls <tt>shader_core_ctx::display_pipeline(...)</tt> in gpgpu-sim/shader.cc for the pipeline states in each SIMT core. For memory partition states, it calls <tt>memory_partition_unit::print(...)</tt> in gpgpu-sim/l2cache.cc.

Gpgpu-sim-admin at 19:40, 13 June 2017

2017-06-13T19:40:25Z

← Older revision		Revision as of 19:40, 13 June 2017
Line 128:		Line 128:

	The simulator software design for this stage is described in [[GPGPU-Sim_3.x_Manual#Fetch_and_Decode_Software_Model \| Fetch and Decode]].		The simulator software design for this stage is described in [[GPGPU-Sim_3.x_Manual#Fetch_and_Decode_Software_Model \| Fetch and Decode]].

			The simulator software design for this stage is described in [[#Fetch_and_Decode_Software_Model \| Fetch and Decode]].

	==== Instruction Issue ====		==== Instruction Issue ====

Gpgpu-sim-admin at 05:01, 5 March 2016

2016-03-05T05:01:01Z

← Older revision		Revision as of 05:01, 5 March 2016
Line 1:		Line 1:
	Revision 1.2 (GPGPU-Sim 3.1.1)		Revision 1.2 (GPGPU-Sim 3.1.1)

	Editors: Tor M. Aamodt, Wilson W.L. Fung		Editors: Tor M. Aamodt, Wilson W.L. Fung, Tayler H. Hetherington

	{{TOClimit\|4}}		{{TOClimit\|4}}

137.82.252.159 at 21:24, 3 March 2016

2016-03-03T21:24:57Z

Show changes

MediaWiki default at 05:10, 3 March 2016

2016-03-03T05:10:42Z

New page

'''MediaWiki has been successfully installed.'''

Consult the [//meta.wikimedia.org/wiki/Help:Contents User's Guide] for information on using the wiki software.

== Getting started ==
* [//www.mediawiki.org/wiki/Manual:Configuration_settings Configuration settings list]
* [//www.mediawiki.org/wiki/Manual:FAQ MediaWiki FAQ]
* [https://lists.wikimedia.org/mailman/listinfo/mediawiki-announce MediaWiki release mailing list]