summaryrefslogtreecommitdiff
path: root/src
AgeCommit message (Collapse)Author
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-05bug fixesTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7815]
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-03bug fix for last check in (problem w/ CUDA 3.1 regressions)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7812]
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. integrating Inder's changesTor Aamodt
2. edit to 3dfd to enable execution on Quadro 3. update script to not pass device directly to 3dfd on command line (which it does not support) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7810]
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-02refactoring: make shd_cache_t into a class (cache_t), plus some other ↵Tor Aamodt
cleaning up [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7806]
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-30integrating CL 7421-7434 from fermi-test Tor Aamodt
this integrates the ptxplus functional model changes, but has only been verified with CUDA 3.1 PTX regressions (which pass) still need to verify ptxplus itself is working in this branch [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7435]
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-24integrate wilson's bug fix for cycle countingTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7257]
2010-08-23- oclMatrixMul passing (added support for dynamically allocating shared ↵Tor Aamodt
memory on kernel launch for OpenCL) - added some debugging printfs to oclMatrixMul [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7255]
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-08refactoring: moving PTX parsing functions into ptx_parser.*Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7167]
2010-08-08refactoring: moving 'loading' operations into ptx_loader.*Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7166]
2010-08-08refactoring: making g_current_symbol_table local to ptx_ir.ccTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7165]
2010-08-08refactoring: start moving prototypes into headersTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7164]
2010-08-08check if debug thread selectedTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7163]
2010-08-08add listing function for ptx debuggingTor Aamodt
print out compilation errors encountered during OpenCL to PTX conversion (still a bit cryptic) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7162]
2010-08-07watchpoint facility for ptx debuggingTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7160]
2010-08-04intial support for integrated interactive debugger (for debugging PTX assembly)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7137]
2010-07-30fix stack pointer update on returnTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7113]
2010-07-30missing definitionTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7112]
2010-07-29compiling (not operand collector probably won't work for anything using tex ↵Tor Aamodt
instructions currently) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7106]
2010-07-29integrate operand collector (not compiling yet)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7105]
2010-07-23debug printing stuffTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6964]
2010-07-21- reduce 3.1 sdk regression test runtimes (as per 2.3 sdk versions)Tor Aamodt
- add ldu support (just use ld for now) - get rid of asserts for .const[n] - status: cuda sdk 3.1 dct8x8, eigenvalues, MersenneTwister, reduction, scan still running (too slow) ocl sdk 3.1 oclQuasirandomGenerator still running (too slow) ocl sdk 3.1 oclReduction fail comparison (cuda 3.1) ocl sdk 3.1 oclMatrixMul exits mid computation ocl sdk 3.1 oclTranspose hits "clCreateBuffer - buffer already created for this host variable" ocl sdk 3.1 oclSortingNetworks fails w/ message about min workgroup size remaining 30 tests pass [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6927]
2010-07-20- adding CL_DEVICE_SINGLE_FP_CONFIG (need to get info on what it should be ↵Tor Aamodt
set to) - bug fix for shader_core object allocation/initialization - bug fix for ret instruction w/ OpenCL + new PTX ABI - some formatting of output [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6910]
2010-07-19OpenCL running on simulator w/ CUDA 3.1 and nvidia driver 256.35Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6907]
2010-07-18finished refactor+implement call/return passing (now working on simple ↵Tor Aamodt
example included in this CL) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6886]
2010-07-18more refactoring of param passingTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6885]
2010-07-18continuing refactoring Tor Aamodt
modify template to actually do a function call (for testing real function call, not printf) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6884]
2010-07-18refactoring call/return ABI code (want unified interface for call/return ↵Tor Aamodt
reg/param) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6883]
2010-07-18finished moving cuda printf stuff to seperate file (step 2)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6882]
2010-07-18moving cuda printf stuff to a seperate file (step 1)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6881]
2010-07-18refactoringTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6880]
2010-07-18- simple device printf support added (supports %u, %d, %f)Tor Aamodt
- add support to call_impl for passing values in through local param space (need to add support for return) - track local framesize during ptx parsing / update local stack-pointer during call/return - detect appropriate param space for ld.param (and st.param) depending upon address symbol's scope - putting back the crazy math for local mem allocation (seems to be required) - bug fix for global byte array initialization - force ld and st w/o space specifier to be generic_space - make type_decode method of type_info_key - adjust debug printing so listing of code and CFG information requires higher numbers (100 and 50) - local param ld/st accesses treated like local ld/st for timing model [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6878]
2010-07-17- added *implied* local memory stack pointer for functional execution Tor Aamodt
(still need to determine framesize during parsing, and need to support alternate mode were stack pointer is explicit register a la zev...) - define local param mapping to local memory for functional execution [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6875]
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 17:25:48 +0000