| Age | Commit message (Collapse) | Author |
|---|
|
There are several problems with this:
1) To get proper occupancy, based on register file usage, you must querry ptxas with
the sm version that corresponds to the register usage in the config file. To enable this,
a new config option has been added that determines what sm version you pass to ptxas to compute reg-usage.
This configuration option is always required in the gpgpusim.config file
2) If you are running PTXPLUS with a modern card (i.e. volta/pascal), you need ptxas from CUDA 9.1.
However since PTXPLUS only supports sm_13 - you need a version of CUDA where cudaobjectdump
supports sm_1x. This ended at CUDA 5 - and PTXPLUS requires CUDA 4.2. Therefore, to run PTXPLUS
on a modern card, you need CUDA 4.2 + modern CUDA installed. To fascilitate this, a new envronment
varaible is added and the setup envrionment script prints an appropraite warning if you are using a newer CUDA.
We have tried to make this as fail-proof as possible - and die appropraitely when something is wrong.
|
|
|
|
|
|
classes
Some changes to remove warnings
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12609]
|
|
. the option -gpgpu_ptx_save_converted_ptxplus allows keeping the ptxplus file
. the option -gpgpu_keep allows keeping intermediate files used to communicate with other programs (e.g. cuobjdump)
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12304]
|
|
Stopping if someone tries to use PTXPlus without cuobjdump
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12138]
|
|
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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 9872]
|
|
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]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7811]
|
|
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]
|
|
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]
|
|
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 = 7166]
|
