summaryrefslogtreecommitdiff
path: root/libopencl
AgeCommit message (Collapse)Author
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-18Re-designed cache model:Tor Aamodt
- read only cache model with integrated mshrs (no L1D, yet); new cache interface should be easily extendable to support texture cache with latency fifo and separate tag/data arrays, though this is not yet added (currently tags and data arrays are not decoupled for texture) - new partition model using the above removes all old MSHRs, L1D etc... passing CUDA 3.1 regression [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7875]
2010-10-161. enable OpenCL to PTX compilation on remote host without NFSTor Aamodt
2. hopefully one-size-fits-all setup_environment script [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7866]
2010-10-09fixes problem with remote OpenCL PTX generationTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7843]
2010-10-09integrate CL 7836 (running regressions on cluster)Tor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7841]
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. 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-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-23- oclMatrixMul passing (added support for dynamically allocating shared ↵Tor Aamodt
memory on kernel launch for OpenCL) - added some debugging printfs to oclMatrixMul [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7255]
2010-08-23- add '-keep' option to keep intermediate files (used for OpenCL only right now)Tor Aamodt
- 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]
2010-08-12- allow multiple clCreateBuffer wth same host_ptr (undefined behavior in ↵Tor Aamodt
OpenCL spec, but required for oclTranspose). - undo earlier fix that wasn't a fix [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7194]
2010-08-10refactor: mostly finished getting rid of extern declTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7183]
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-08refactor: moving extern variable declarations into headersTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7172]
2010-08-08- remember to return CL_SUCCESS if we allocate memory successfullyTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7171]
2010-08-08add clFlush (empty implementation) - oclMatrixMul completes but is failing ↵Tor Aamodt
(differences) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7170]
2010-08-08oclSortingNetworks passing: update workgroup size calculation (now ignores Tor Aamodt
shared memory usage since I believe the assumption is at least a single CTA can run for this calculation) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7169]
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-08-08add listing function for ptx debuggingTor Aamodt
print out compilation errors encountered during OpenCL to PTX conversion (still a bit cryptic) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7162]
2010-07-23debug printing stuffTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6964]
2010-07-21- reduce 3.1 sdk regression test runtimes (as per 2.3 sdk versions)Tor Aamodt
- add ldu support (just use ld for now) - get rid of asserts for .const[n] - status: cuda sdk 3.1 dct8x8, eigenvalues, MersenneTwister, reduction, scan still running (too slow) ocl sdk 3.1 oclQuasirandomGenerator still running (too slow) ocl sdk 3.1 oclReduction fail comparison (cuda 3.1) ocl sdk 3.1 oclMatrixMul exits mid computation ocl sdk 3.1 oclTranspose hits "clCreateBuffer - buffer already created for this host variable" ocl sdk 3.1 oclSortingNetworks fails w/ message about min workgroup size remaining 30 tests pass [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6927]
2010-07-20- adding CL_DEVICE_SINGLE_FP_CONFIG (need to get info on what it should be ↵Tor Aamodt
set to) - bug fix for shader_core object allocation/initialization - bug fix for ret instruction w/ OpenCL + new PTX ABI - some formatting of output [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6910]
2010-07-19updating opencl wrapper supportTor Aamodt
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6892]
2010-07-17- add support for ld/st with generic addressesTor Aamodt
- 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]
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:55:30 +0000