summaryrefslogtreecommitdiff
path: root/src/gpgpu-sim/gpu-sim.h
AgeCommit message (Collapse)Author
2010-10-161. creating cache_config object to encapsulate cache configuration informationTor Aamodt
(and parse it before creating the simulator objects). 2. creating core_config to hold only features of a shader_core that are high level enough either (a) the functional simulator needs to know about them, or (b) they affect memory *access* generation. 3. in config files only (so far) separate out notion of write-{through,back}, from notion of when a line is allocated... will use this to distinguish different types of caches. passing CUDA 3.1 regression [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7870]
2010-10-161. refactoring histogram/logger so that classes are in header filesTor Aamodt
2. starting to redo cache_t 3. deleting more perf counters 4. other minor cleaning [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7869]
2010-10-161. moving address decoding into a class (and out of cache entirely)Tor Aamodt
2. moving DRAM timing parameters into memory_config class 3. removing some likely useless statistics 4. other cleaning up of code passing CUDA 3.1 regression [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7868]
2010-10-121. adding simt_core_cluster, which models a TPC or (for fermi) GPC...Tor Aamodt
this gives us a place to stick caches shared among shader cores but on the shader side of the interconnect... maybe move the clock boundary code here? after integrating booksim 2 code? 2. added a pending write table to ldst_unit rather than scoreboard ... rationale is that ld/st unit needs to process register writes once it is done it can notify scoreboard once. 3. re-enabled shared memory delay (use pipeline within ldst_unit) 4. re-enabling operand collector writeback for all instruction types 5. disable MSHRs in this change list passing CUDA 3.1 regression next? texture cache, then redo mshrs? [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7845]
2010-10-101. create function unit classes for SP, SFU, LD/ST.Tor Aamodt
2. refactor memory stage into a ld/st function unit 3. refactor memory access generation (moved into warp_inst_t class) the above should make supporting fermi uarch much easier passing CUDA 3.1 regression still need to... (a) update scoreboard to keep count of outstanding memory requests and use operand collector for writebacks into register file (b) add back shared memory pipeline delay (c) remove use of MSHR's for non-cached global/local accesses (d) replace texture cache with a split tag/data array pipe (e) re-implement memory_partition stuff so it makes more sense [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7844]
2010-10-09Refactoring:Tor Aamodt
1. Moving mem_access_t to abstract_hardware_model and making set (queue) of accesses part of warp_inst_t. I.e., treat set of accesses as an ISA concept rather than a hardware organization concept. This is only partly "done"... logic for computing accesses is still part of shader_core_ctx in this CL. Given number of warp_inst_t accessors for accessq, now seems like we might even want to move some memory stage code into warp_inst_t class. How those accesses make it to memory system is the hardware concept. 2. Making warp_inst_t an explicit arguement of subroutines used in memory stage... The eventual goal here is (likely) to refactor memory into a hardware block... i.e., have function units be a class that contains some set of pipeline stages internally and some set of input/output "ports". 3. Moving accessor functions is_load, is_store; is_const, is_local into class declaration (where they belong). 4. Removing code for selecting pipeline uarch (might add it back later, but first want a clean GT200 organization). In particular, removing option to have an operand collector -- now you MUST have the operand collector. 5. Removing more deadcode from prior changes (fixed delay queue related) Scripts/configs: 6. Correlation script not printing out exit condition when hardware launch fails 7. Update config files to have proper compute model selected [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7834]
2010-10-08some fixes for ptxplus (correlation test now running)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7828]
2010-10-081. refactoring cuda api code for loading PTX (removing external PTX loading ↵Tor Aamodt
entirely) 2. some bug fixes for warp_inst_t 3. creating a new class, gpgpu_t, which contains the functional "memory" state visible to all threads running on a GPU (doing this as part of my continuing effort to hunt down and eradicate every global variable that is not the top level "the gpu") 4. other misc. changes Almost passing CUDA 3.1 regression? oclHistogram keeps failing under torque, but does not fail when run on the command line from the same directory. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7827]
2010-10-05simulator "working" on ptxplus, simulated IPC "too low" by about a factor of ↵Tor Aamodt
up to 50% [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7816]
2010-10-05broken change list: builds, but does not run, yetTor Aamodt
refactoring: create warp_inst_t that provides notion of a group of scalar instructions traveling down the pipeline. delete DWF delete MIMD delete warp_tracker delete old writeback stage, replace it with a stub that just writes back everything delete old pipeline model current status: MSHR's need to change to deal with the new structure [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7814]
2010-10-03moving some CUDA API centric code to cuda_runtime_api.ccTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7811]
2010-10-031. enable L2 cache as a texture cache (also some bug fixes for L2 as regular ↵Tor Aamodt
cache) 2. update gpgpusim.config for Quadro to use L1 cache geometry from Henry's ISPASS paper 3. minor edit to CUDA api : add notion of fat_cubin_handle (currently not used for anything) 4. minor edits to deadlock detection message (more accurate reporting of source of deadlock) 5. other minor edits [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7809]
2010-10-02refactor: mem_fetch now a classTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7807]
2010-10-01integrating recent changes from fermi-test into fermiTor Aamodt
(i'll use "fermi" for more disruptive changes to the pipeline model such as updating the MSHRs and getting rid of the warp tracker, ripping out DWF, etc...) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7805]
2010-08-29- integrate changes from fermi-test (CL's under that path in range 7261-7418).Tor Aamodt
(add scoreboard logic from ptxplus branch and modified operand collector with parallel ALU/SFU pipelines) passing regressions [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7419]
2010-08-24some minor cleaning upTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7259]
2010-08-23- add '-keep' option to keep intermediate files (used for OpenCL only right now)Tor Aamodt
- remove redundant initialization code of options (which confuses new users) - make bool option variables bool type - remove some more extern decls [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7247]
2010-08-10refactor: mostly finished getting rid of extern declTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7183]
2010-08-09refactor: remove extern decl up to gpgpusim_entrypoint.ccTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7178]
2010-08-08refactor: shader.cc free of extern declarationsTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7176]
2010-08-08refactoring: continued, buildsTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7174]
2010-08-08refactoring: mem_latency_stat in impl, etc...Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7173]
2010-08-08refactor: moving extern variable declarations into headersTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7172]
2010-08-08** fully decouple kernel compilation from kernel launch for openclTor Aamodt
changes required to do this (this enables oclReduction SDK 3.1 to pass and is a prerequisite to Fermi style concurrent kernel launching) - separate kernel launch for cuda and opencl (and init_grid functions) - add pgm_info structure for holding kernel function_info pointer in _cl_program object after ptx compilation ** added support for clEnqueueNDRangeKernel with local_work_size==NULL which is part of OpenCL spec and used in oclSortingNetworks (however, this is a braindead implementation that only handles the case where global_work_size[0] is smaller than the max number of threads per shader; moreover oclSortingNetworks is still not working but for what looks like another reason) ** refactoring / cleanup - g_global_symbol_table made static to ptx_parser.cc - remove g_kernel_name_to_symtab_lookup (not really being used) - moving various function prototypes into headers (e.g., if foo() defined in bar.cc, then place prototype in bar.h) - adding icnt_reg_options() [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7168]
2010-08-07watchpoint facility for ptx debuggingTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7160]
2010-07-15creating branch for adding support for CUDA 3.x and FermiTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829]
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 09:47:23 +0000