summaryrefslogtreecommitdiff
path: root/libcuda
AgeCommit message (Collapse)Author
2014-08-14Remove -gpgpu_ptx_use_cuobjdump from configs. Update error reporting to be ↵Tor Aamodt
more helpful. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12654]
2014-08-14Catching the problem where the embedded ptx cannot be found because the ↵Andrew M. B. Boktor
benchmark is compiled with CUDA4 and cuobjdump is not enabled [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12653]
2014-08-14supressing signed to unsigned comparison and unused warnings (flex generates ↵Andrew M. B. Boktor
them) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12636]
2014-08-14Removing warningsAndrew M. B. Boktor
A bit of cleanup [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12627]
2014-08-14Re-wrote the intermediate structure for the output of cuobjdump using C++ ↵Andrew M. B. Boktor
classes Some changes to remove warnings [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12609]
2014-08-14Removing a couple of warningsAndrew M. B. Boktor
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12564]
2014-08-14Adding copyright notices and updating CHANGESAndrew M. B. Boktor
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12552]
2014-08-14Fixing a small bugAndrew M. B. Boktor
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12523]
2014-08-14Reworking the mechanism that selects which sm versino to run.Andrew M. B. Boktor
Now, the highest sm version below forced_max_capability is selected for each cu file individually [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12516]
2014-08-14Supporting the option for removing temporary filesAndrew M. B. Boktor
. 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]
2014-08-14Handling the case where force_max_capability is set to 0 (means no restriction)Andrew M. B. Boktor
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12280]
2014-08-141. Fixing ptx with sm_20Andrew M. B. Boktor
2. Fixing the choice of files (since new we can choose) 3. Not allowing sm_20 with PTXPlus, forcing max capability it to 19 and printing a warning [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12279]
2014-08-14The generic makefile applies rules to files outside the dir and failsAndrew M. B. Boktor
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12277]
2014-08-14Moving things where they belongAndrew M. B. Boktor
cubojdump.l/.y are used to parse the output of cuobjdump, and thus should be part of libcuda not cuobjdump_to_ptxplus Included are also modifications to Makefiles [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12255]
2014-08-14Using nvcc present in CUDA_INSTALL_PATH instead of the one in PATHAndrew M. B. Boktor
Adding -DCUDART_VERSION to CXXFLAGS [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12253]
2014-08-14Removing all remaining traced of decudaAndrew M. B. Boktor
Stopping if someone tries to use PTXPlus without cuobjdump [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12138]
2014-08-14This changelist implements the following:Andrew M. B. Boktor
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]
2014-08-14Back out changelist 10951Hadi Jooybar
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10963]
2014-08-14Should be tested.Hadi Jooybar
Does not support sm_20 [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10951]
2014-08-14The simulator stops and outputs a message if it cannot find a cubin file ↵Andrew M. B. Boktor
while trying to convert to ptxplus [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10869]
2014-08-14Undoing the rest of files mistakly changed in CL10695, this change completes ↵Ayub Gubran
the work of CL10697 which undone the changes to one file (stream_manager.cc). Here I undo the changes for 4 make files and more importantly the setup_enviroment file where I uncommented the include path of openCL as the case before CL10695, and I kept Tim changes to the CUDA path done later in CL10723 which he made originally to undo the work I did but didn't uncommet the openCL path which I do here. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10730]
2014-08-14My last work added to the branch, wokring on the functional simulator in ↵Ayub Gubran
cudasim. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10695]
2014-08-14- Minor change to make things simpler, basically removing all instances ofAndrew M. B. Boktor
$CUDAHOME and replacing them with $CUDA_INSTALL_PATH. Removing all instances of $NVIDIA_CUDA_SDK_LOCATION with $NVIDIA_COMPUTE_SDK_LOCATION. - Some additions/changes to the README file to make it a little more intuitive. - Default values added to setup_environment to make things easier for the average user. - My first perforce submit :D [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10401]
2014-08-14Assigned debug level 3 to stream manager output.Inderpreet Singh
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10013]
2011-07-07Integration change. Bringing in some changes from mem_divergence that allow ↵Tim Rogers
for multiple configs of the built simulator to exist at one time. Now you no longer have to clean build when changing from debug to release configs it also eliminates the possibility of having a fraken-file where some objects are in debug and some are in release. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 9743]
2011-04-09initial edits to compile with CUDA 2.3 (passes regression with CUDA 3.1)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 8762]
2010-12-28- parameter memory and active threads now part of kernel_info_t:Tor Aamodt
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]
2010-12-28- Checkpointing new support for concurrent kernel execution (CUDA only, not ↵Tor Aamodt
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]
2010-10-241. adding top level configuration class and making shader and memory ↵Tor Aamodt
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]
2010-10-09refactoring: moving texture reference mappings into gpgpu_t Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7835]
2010-10-081. modify shader_core_ctx::execute_pipe() to model instruction throughput ↵Tor Aamodt
correctly (off-by-one error) 2. adding code to dump_pipeline to display reason for memory stage stalling 3. removing dead code resulting from prior changes correlation vs. GT200 is 0.95 (need to add back modeling of memory writebacks, shared memory latency) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7831]
2010-10-08some fixes for ptxplus (correlation test now running)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7828]
2010-10-081. refactoring cuda api code for loading PTX (removing external PTX loading ↵Tor Aamodt
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]
2010-10-03moving some CUDA API centric code to cuda_runtime_api.ccTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7811]
2010-10-031. integrating Inder's changesTor Aamodt
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]
2010-10-031. enable L2 cache as a texture cache (also some bug fixes for L2 as regular ↵Tor Aamodt
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]
2010-10-01integrating recent changes from fermi-test into fermiTor Aamodt
(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]
2010-08-08refactor: shader.cc free of extern declarationsTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7176]
2010-08-08refactoring: mem_latency_stat in impl, etc...Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7173]
2010-08-08** fully decouple kernel compilation from kernel launch for openclTor Aamodt
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]
2010-08-08refactoring: start moving prototypes into headersTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7164]
2010-07-18simple event support re-enabled (e.g., 3dfd wants this)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6887]
2010-07-16- ifdef out stream stuff (not compatible with CUDA 3.1 :)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6853]
2010-07-16- integrate Wilson's bug fixes from release Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6852]
2010-07-16- template runs/passes with CUDA 3.1, but requires that CUDA 2.3 Tor Aamodt
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]
2010-07-15compiles with CUDA 3.1 (but, nothing tested)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6832]
2010-07-15creating branch for adding support for CUDA 3.x and FermiTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829]
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 23:44:37 +0000