| Age | Commit message (Collapse) | Author |
|---|
|
Parameters are finalized at kernel launch, which means the contents
of parameter memory are initialized. Kernel arguement names have a
fixed order, hence same address should be assigned on subsequent
kernel launches of same kernel in other streams provided the data size
param_t::size of arguments for each kernel launch is identical (an
assertion has been added to check this is true).
- passing regression
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 8303]
|
|
OpenCL)
This changelist adds full support for streams supported by a new class,
stream_manager and enables concurrent execution of kernels from different
streams.
- fast_regression.sh fails for simpleMultiCopy, simpleStreams (other tests
passing)
** Known issues **
- Kernel parameter passing is not done correctly for concurrent kernel execution
(somehow concurrentKernels is not affected by this): the parameters are
stored inside function_info, which is shared among parallel kernel launches
so that the values passed into the launch are likely to get overwritten if
multiple grids are launched in parallel streams.
- Statistics are printed out whenever the simulation thread runs out of
cuda commands (doesn't make sense to print out when a kernel ends during
concurrent kernel execution). This will probably require further tweaking
so as to be more compatible with data collection scripts.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 8302]
|
|
except it reads and stores from two adjacent registers instead of a single 64 bit register. ".ff64" instructions are now printed in decuda_to_ptxplus. Support in the simulator for ".ff64" has been added but it is untested.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 8278]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 8179]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 8153]
|
|
stalling to send four requests per warp into L1T tag lookup.
If L1T is really 32B blocks (as per Henry's paper), this suggests
banking of L1T needs to be modeled.
Other changes:
1. bug fix in memory access generation for texture/const cache access
2. adding back memory latency measurement for visualizer
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7913]
|
|
2. update texture to bypass ROP-delay queue... correlation now 0.9592
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7912]
|
|
configuration
components of this class.
2. clock memory pipeline no. subwarp times for each shader clock and increase
rob-size for texture cache (trying to improve correlation, currently at 0.9218)
3. start to modify shader stats to add back features for visualizer (warp
divergence distribution kind of working again)
passing cuda 3.1 regression and ptxplus correlation tests
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7909]
|
|
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]
|
|
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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7840]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7835]
|
|
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]
|
|
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]
|
|
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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7812]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7811]
|
|
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]
|
|
(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]
|
|
(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]
|
|
memory on kernel launch for OpenCL)
- added some debugging printfs to oclMatrixMul
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7255]
|
|
- 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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7178]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7176]
|
|
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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7166]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7165]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7164]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7160]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7137]
|
|
- 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]
|
|
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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6907]
|
|
- 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]
|
|
(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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6874]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6873]
|
|
local storage for local param memory)
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6871]
|
|
parameter spaces (kernel parameters are conceptually different from
function parameters, but PTX only has one .param keyword)
- removing more dead code
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6870]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6869]
|
|
- add support for parsing cache operators (still need to pass to timing model)
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6868]
|
|
a fixed address mapping between shared,local to generic that depends
upon hardware thread context used... might be interesting to explore
tradeoffs at some point)
- remove util.h... we don't need TRUE, FALSE anymore now that
everything is C++
- remove some dead code from shader_decode
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6867]
|
|
- modify template to use printf from kernel (this works on our Fermi card--want it working on simulator for debugging)
- add support for extern function declarations (so don't try to assemble them)... printf is extern
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6862]
|
|
is installed for math intrinsics used in instructions.cc (since
CUDA 3.1 does not seem to define __uint2float).
- currently cuda-math.h has hardcoded paths to my own install of
CUDA 2.3 (need to determine whether we can build simulator with
CUDA 2.3 and run CUDA 3.1 / compute_2x PTX, or need to make
current setup more portable)
- modified ptxinfo.l to print out meaningful error message if
error occurs on first line
- modified ptxinfo.l,y to understand "for 'compute_NN'" syntax
- updated template to 3.1 SDK version
- updated cuda_runtime_api.cc to select highest compute capability
version when multiple versions are present.
- updated template/BlackScholes Makefiles to generate a "DLL build"
instead of static linked build (still confusing... requires one to
copy config files... probably need to modify config scripts)
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6837]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829]
|
