| Age | Commit message (Collapse) | Author |
|---|
|
Kernel launch to stream not yet implemented
|
|
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 17048]
|
|
Fixing the alignment of all types of memory so that it's aligned to size or 128byte whichever is smaller
Refactoring out the padding code
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 17047]
|
|
external C linkage
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 15912]
|
|
//depot/gpgpu_sim_research/fermi_tim/...
to //depot/gpgpu_sim_research/fermi/...
Integrating CLs up to 15295. Descriptions of these CL's are included.
***
A couple changes to aeriel-vision for warp issue plot support
***
More arielvision changes to support the variable-entry length stacked bar chart
***
Properly printing the right resolution of dynamic warp ids
***.
Generalized the scheduler code and added detailed statistics for which warps issue each cycle.
Verified the execution of the LRR scheduler - still have to get the two level scheduler to work.
***
Implementing the 2lvl scehduler has it has been originally coded.
LRR on both the inner and outer levels
***
Adding in a debug tracing system to GPGPU-Sim.
I am sick of writing debug code - then having to comment out, ifdef out or delete it to checkin.
This also allows for print streams so the user can decided which traces they would like to see.
Every print in GPGPU-Sim should go through this system - then it will be really easy to only get the information you want and more importantly people will (a) write and (b) checkin code that actually profiles what they are building.
Reading tracefiles is superiour in many ways to single stepping since you can print the world and just vet the logfile for what you need.
This also fascilitates advice from the Debugging Rules! book which states that you should never throw away a debugging tool. Having debug prints that don't get thrown away is big.
***
Allowing the trace to be specified in the Make.
Run Make TRACE=0 to compile the code without any traces
***
Allowing prints from the performance sim to get the actual ptx instruction text
***
Getting the two level scheduler to actaully work...
What is released in fermi does not work at all - it effectively performs "static warp limit" from my CCWS paper.
Warps are never demoted from the active list since the functionality checking to see if they are waiting on a longop is completly broken.
Maybe if the original author had access to the tracing functions this would not have happened.
The islongop test was completely broken. It did not mark the register as used, it marked the register number in the instruction as used.
For example if this instruction was creating a long op:
ld r6 [r1]
It would mark register 0 as waiting for a long op (since it is register 0 of the two registers in this instruction), not register 6.
Additionally, whenever ANY instruction from a warp releases registers, ALL the longops being tracked for this warp get cleared....
The only way anyone ever thought this worked is if they did not test it....
***
Reworking the warp schedulers to share common code.
Making the GTX480 use gto by default. I am not sure wht they really use, but it really can't be LRR.
Also adding in a new file for custom shared trace defines. These are useful when you want a print that has some
additional criteria or information printed.
Verified that the schedulers all work to a first order based on traces.
***
Making it so you can run the stats collection scripts from any directory.
Also allow the caller to specify a stats file instead of just assume its always the same one
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 15296]
|
|
shared memory buffer to those pointers. This is required to support OpenCL local memorywith the newer NVIDIA driver.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 14565]
|
|
4.1 and later. KMN-FT should pass with this changelist.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 13833]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 13761]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 13703]
|
|
Adding support for double destination to mad instruction
Fixing broken madp instruction
Adding a patch to cuobjdump_to_ptxplus to work around the C3 problem (Documented in bug #154 internal).
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 13349]
|
|
1. Adds support for using cuobjdump for both ptx and ptxplus execution. This has been tested with CUDA 4.0
. Ptxplus is no longer supported through decuda/decuda_to_ptxplus
2. Adds support for converting the SASS output by cuobjdump to ptxplus. This has been tested with CUDA 4.0
. The old path that extracts ptx from cubin files is still preserved
3. Adds a bank group model. (WARNING: memory config has changed, please adapt yours). To disable the bank groups model, set nbkgrp to 1 and tCCDL and tRTPL to 0
Diff the configuration files to learn about how to use those new options.
Merging
//depot/gpgpu_sim_research/fermi-test/distribution/...
to //depot/gpgpu_sim_research/fermi/distribution/...
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12023]
|
|
(expected). The key is that the missing function is now in place. Also removed reference to print_shader_cycle_distro() (this is deprecated by AerialVision).
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 11346]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 9872]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 9687]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 8119]
|
|
passing CUDA 3.1 and ptxplus correlation
correlation back to around 0.89 on ptxplus vs quadro fx5800
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7898]
|
|
(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]
|
|
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]
|
|
(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]
|
|
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]
|
|
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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7167]
|
