| Age | Commit message (Collapse) | Author |
|---|
|
(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]
|
|
- 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 = 7183]
|
|
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 = 6907]
|
|
example included in this CL)
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6886]
|
|
- 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]
|
|
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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6857]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829]
|
