-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathCHANGES
198 lines (192 loc) · 10.2 KB
/
CHANGES
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
CHANGE LOG:
Version 3.1.0 versus 3.0.2
- Support for CUDA 4.0 for both PTX and PTXPlus.
- Support for using cuobjdump to extract PTX/SASS from binaries.
- Complete rework for the SASS to PTXPlus conversion code including support for
code generated by cuobjdump.
- Support for configurable pipeline widths.
- Support for a configurable number of functional units within an SM.
- Model for GDDR5 bank groups.
- GDDR5 configuration.
- Extended the result bus model to allow issuing/retiring more than one
instructions from multiple function unit per cycle.
- Support for configurable instruction latency and initiation interval
- Added corresponding configuration to Fermi and Quadro config files according
to benchmark results
- Support for 64-bit atomicAdd, atomicExch, and atomicCAS.
- Support for 32-bit floating point atomicAdd.
- New installation instructions for AerialVision, removed deprecated install
script.
- Rework for some makefiles
- Removed the no longer supported/used decuda_to_ptxplus
- Removed all dependencies on the boost libraries. Now GPGPU-Sim can run
without boost.
- Added missing support for -gpgpu_perfect_mem. When enabled this option models
a memory system with single cycle latency per memory request for accesses that
miss in the L1 cache. Bandwidth is limited to one memory request issued per
SIMT core cycle. Note this means uncoalesced accessess will be slower than
coalesced accesses.
- Bug Fixes:
- Added a missed call to operand collector writeback which fixes some
writebacks that were not considered by the operand collector.
- Fix for a bug that caused global symbol tables to be re-initialized after
parsing each ptx file.
- Fix for a rare race condition that prevented kernel stats from being
printed
- Fix for miscalculation of the gpu_total_sim_rate value printed at the end
of a kernel launch
- Fix for a bug that caused atomic accesses not to have a payload to the
Atomic Operation Unit; atomic accesses now have a payload both to and
from the Atomic Operation Unit
- Fixed a bug that caused deadlock check to be omitted
- Updated the Fermi config files such that when an access misses the L1 data
cache, it allocates a line immediately before sending a data fetch request
out to the memory partition.
- Changed the writeback arbitration among multiple clients in the LDST unit
to round-robin.
- Fixed a rare pathological case that caused LD instructions with divergent
memory accesses to be overcounted.
- Fixed linear_to_raw_address_translation::partition_address() so that it
works for non-power-of-two number of memory partitions. Before, it just
returns the input memory address without any change.
- Fixed warp_inst_t::memory_coalescing_arch_13() to use the
gpgpu_shmem_warp_parts option to divide up the warps in the coalescing
logic. It was hardcoded to 2, which causes coalescing to always operate
on half-warps.
Version 3.0.2 versus 3.0.1
- Added Fermi configuration
- Revived functional simulation mode
- AerialVision: Revived support for PC-Histogram
- AerialVision: Revived collection of source code view metrics (all except
exposed pipeline latency)
- Added missing function implementations (caused static linking with GPGPU-Sim
to fail).
- Fixed div_impl with 32-bit and smaller integers (Uninitialized upper bits in
the operands was causing the instruction to behave incorrectly).
- Updated setup_environment script to handle host names with empty domains.
- Forced L2 cache line size to be greater than or equal to L1 cache line size
- L2 caches can now be disabled
- Increased Quadro config's L2 line size to 256B and total L2 size to 256kB
- Fixed generation of memory requests for local.s8 load instructions (fixed a
bug in local memory address translation)
- Conflicting atomic accesses now generate multiple requests during coalescing
- Updated Doxygen scripts
- Fixed dump pipeline in case of L1 disabled.
- Fixed a bug in the debug printing system by making the enum and string array
use the same list
- Fixed a bug where the ldst_unit::cycle and ldst_unit::writeback were being
processed at different rates
- Changes to setup_environment to prevent path explosion and crashing due to
empty host domain name
- Changed the remote OpenCL compilation to use rsync instead of scp (scp had
issues with concurrent copying).
- Fixed the data collection code for the following metrics (in 3.0.0 and 3.0.1
they were only counting memory instructions that generate off-core traffic):
gpgpu_n_load_insn
gpgpu_n_store_insn
gpgpu_n_shmem_insn
gpgpu_n_tex_insn
gpgpu_n_const_mem_insn
gpgpu_n_param_mem_insn
- Fixed a bug where the PTX load and store instructions' input address register
dependencies were not being registered with the scoreboard.
- Updated how pending hits in data cache are reported (3.0.0 and 3.0.1 reports
the difference between the number of cache misses and pending hits).
- Fixed a bug where a configuration with non power-of-two number of memory
partitions can cause two different linear address to be aliased into the
same intra-partition address at the same memory partition.
Version 3.0.1 versus 3.0.0b
- Updated README and setup_environment for more streamlined setup.
Including install instructions for Ubuntu 10.04.3 LTS
- Changed Quadro FX 5800 config to use PTX by default (not SASS/decuda)
- Added Doxygen generated documentation
- Removed ISPASS 2009 configuration (use Quadro FX 5800 config instead)
- Removed common.mk (run native CUDA/OpenCL applications)
- Removed out-of-date GPGPU-Sim documentation. This wil be updated in a
later release. Refer to Doxygen documentation and/or source code.
- Added prints for L1 data cache statistics
- SIMD width option removed from shader_core_pipeline_opt (was not used)
- Added read-to-precharge constraint in DRAM
- Disabled Stream Manager's verbose output by default, use debug level 3 to
enable it
- Addresses returned by memory allocation are now 256 bytes aligned
- Ejection from the clock domain interface buffer between interconnection
network and L2 cache happens in the L2 clock domain instead of interconnect
clock domain.
- Update OpenCL support to work with AMD OpenCL sample applications
- Bug fixes
- Fixed the variation in instruction count seen under different cache
configurations on the same workload
- Fixed unnecessary flushing of instruction buffer
- Fixed mislabeling of stall cycles recorded by scheduler
- Fixed compilation issues when using gcc 4.5.1
- Cleanup and bugfixes to build scripts
- Fixed gpu_sim_cycle with sequential kernel launches
- Fixed memory corruption bug in decuda_to_ptxplus causing parse errors
Version 3.0.0b versus 2.1.2b
- massive refactoring of code to C++ (most global variables eliminated)
- support for executing native SASS code via decuda
- added to timing model:
- in order scoreboard allowing multiple instructions per warp in
pipeline at a time
- operand collector for modeling access to banked register files
- prefetching texture cache model (Igehy et al., Graphics Hardware 1998)
- updated data and constant cache models with updated MSHR model
- other changes I don't remember right now
Version 2.1.2b versus 2.1.1b:
- Better 1D texture support
- bar.sync timing change for compatibility with NVIDIA hardware
(release barrier when all warps reach barrier, irrespective of
divergence state). Functional simulation mode disabled to
implement the expected barrier behavior under branch divergence
- Refactored memory stage
- Limited miss status holding registers
- Experimental support for running on Mac OS X
- Register bank conflict modeling (see MICRO-42 tutorial)
- Changed files to C++
- Updated makefiles
- Various improvements to aerialvision
- Simulation speed improvement
- Various minor fixes
Version 2.1.1b versus 2.1.0b:
- Added support for OpenCL (requires OpenCL driver from NVIDIA)
- Added performance visualizer tool
- Added manual documenting usage of the simulator (see doc directory)
- Added configuration file for Quadro FX5800 (see configs directory)
- Simultaneous release of benchmarks from ISPASS 2009 paper (separate download)
- Improved shared memory bank conflict modeling
- Improved default mapping of local memory accesses to global address space
- Added interconnect concentration modeling (to approximate effect of a TPC)
- Added support for GPU-to-GPU memory copies (same GPU), timing not modeled
- SimpleScalar code (for command line parsing and functional simulation of
memory) removed.
- PTX support for vector operands in mov instruction
- Bug fixes
- stack overwrite bug fix (with thanks to Jake Adriaens)
- interconnect clock domain crossing "bottleneck" bug (was not removing
traffic quickly enough going from memory to shader cores)
- immediate post dominator detection (affected eigenvalues SDK benchmark)
- CTA/core calculation error with PTX containing calls and returns
- global symbol lookup error when building with multiple .cu files
- Other general improvements and bugfixes
Version 2.1.0b versus 2.0b:
- Added support for parsing and functionally simulating up to CUDA 2.2
generated PTX
- Added dynamic warp formation timing model (see MICRO'07, ACM TACO'09 papers)
- Updated gpgpusim.config and mesh in benchmark subdirectories to be similar to
ISPASS 2009 paper baseline configurations
- Added OpenGL interoperability support
- Added support for parsing embedded PTX files without requiring recompilation
(suggested by Gregory Diamos)
- Improved support for texture filtering (linear mode for 2D, closer agreement
to hardware for 2D point sampling where sample points are close to texel
boundaries)
- Benchmark examples updated to version from CUDA 2.2 SDK (NOTE: these will not
compile with earlier CUDA installations. However, GPGPU-Sim should still work
with applications written for older versions of CUDA.)
- Fixed splash screen: was reporting incorrect version
- Fixed bug comparing strings in basic block formation code
- Fixed bug with additional basic block (which lead to incorrect postdominator
analysis on some benchmarks with recent versions of CUDA)
- Fixed "warning: "/*" within comment" when building.
- "Other general improvements"