| Age | Commit message (Collapse) | Author |
|
|
|
|
|
|
|
__cudaPopCallConfiguration, neither of which are documented in the CUDA API but they preceed any cudaLaunchKernel() call and push the call configuration on a stack which is then popped when cudaLaunchKernel() is called. Currently the max stack size is 1, an assert will be triggered when this is violated.
|
|
|
|
|
|
|
|
|
|
stuck in deadlock
|
|
added mention of a couple of recent changes
|
|
specifically, changed the number of clusters, controllers, and interconnect nodes. Also modified parse_and_compare to fix an occasional division by zero bug.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 21812]
|
|
architecture. Modified shader.h to allow for larger CTA per warp, to accomodate Maxwell specs.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 21798]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 21797]
|
|
|
|
|
|
|
|
|
|
interconnection node number
Signed-off-by: Myrice <[email protected]>
|
|
Signed-off-by: Myrice <[email protected]>
|
|
initial support for CUDA 5.0, 5.5, 6.0 to get template from SDK running
|
|
|
|
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 18452]
|
|
- Added a parameter to the cache configuration to configure the set index function.
- Added a hash set index function to the Fermi L1 data cache for the two default cache sizes, 16KB/48KB with 32/64 sets.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 18202]
|
|
http://www.gpgpu-sim.org/bugs/show_bug.cgi?id=81
Review ID: 173001
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 18179]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 17739]
|
|
http://www.gpgpu-sim.org/bugs/show_bug.cgi?id=78
review: 123001, LGTM: 2
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 17692]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 17315]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 17183]
|
|
fix dram clock energy scaling
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 17124]
|
|
Review ID: 89001
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 17086]
|
|
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]
|
|
write-allocation policy. Also added description for the write-allocation fix implemented by Tayler.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16914]
|
|
is used.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16904]
|
|
Review ID: 75001 lgtm: 2
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16758]
|
|
Details: See Review 80001 https://gpgpu-sim-code-review.appspot.com/80001/
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16747]
|
|
and the data port in each bank is limited to 32B/cycle.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16700]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16671]
|
|
keep the DRAM utilized. Also extended the state printing function to print out mem_fetch entries inside the DRAM delay queue.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16639]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16635]
|
|
partition. Each L2 cache banks has its own connection to the interconnection network to allow L2 bandwidth to scale without increase the number of memory parttiion units.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16613]
|
|
still caching data from local memory space.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16601]
|
|
Main contribution is a static warp limiting scheduler.
There is also some minor cleanup to the heirarchy of the cache code and removal some excessively long lines
Review ID: 36001 lgtm: 1
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16580]
|
|
contains no data. Instead of crashing, it prints out an warning and proceeds.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16569]
|
|
Also changed some initialization code when cores are created in both the funcational and perfromance simulator
review:3001 lgtm:5
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16506]
|
|
generated go in a directory independent of the source directory
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16503]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16495]
|
|
simulation liveness printout (default to 1 per second in wall clock time).
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16482]
|
|
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16480]
|
|
Not sure if this is the only issue, but this was definitely a source of non-determinism (checked with many of the <rodinia>-FT benchmarks). In gpgpu_sim_thread_concurrent() (gpgpusim_entrypoint.cc), the inner do{...}while(active); only breaks when the gpu is not active. As a result, the gpu is only initialized when the gpu becomes momentarily inactive. If one kernel completes while another kernel is currently pending in the stream's queue, the next kernel will start running immediately causing the GPU to not be reinitialized (which includes resetting per-kernel stats). The fix simply recognizes that a kernel has completed and breaks from the loop prior to starting the next operation.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 16463]
|