summaryrefslogtreecommitdiff
path: root/libopencl/opencl_runtime_api.cc
AgeCommit message (Collapse)Author
2014-08-14Fix for bug 51 and 55: Declaration of clCreateContextFromType() is updated ↵Wilson Fung
to match OpenCL 1.1 and later. The API call now accepts device types besides CL_DEVICE_TYPE_GPU. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 15796]
2014-08-14 Fixing a slew of compiler warningsTim Rogers
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 15218]
2014-08-14Apply patch from Kito Cheng to update libopencl for checking NULL error code ↵Wilson Fung
pointer. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 15043]
2014-08-14OpenCL newer diver fixHadi Jooybar
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 14554]
2014-08-14Integrated in CL14390.Inderpreet Singh
Fixed clCreateProgramWithSource to accept NULL or 0 as string length parameter Fixed replacement of printf modifiers when compiling OpenCL code (e.g. '%f') [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 14391]
2014-08-14Fixing bad memory references in opencl based on some disconnect between what ↵Tim Rogers
the nvidia opencl wraper expects and what our opencl_runtime_api provides. This was causing us to print grabage in the dumped ptx files generated by calls to oclPtxLog and valgrind complains about it. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 13837]
2014-08-14Integration from tm-311 branch. Adding support for atomic operation with ↵Wilson Fung
generic memory space. Also adding a define to allow support for OpenCL 1.0 commands that are deprecated in OpenCL 1.1. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 13722]
2014-08-14Fixing a bug in the opencl prebuilt ptx flow.Tim Rogers
If a kernel made modidifcations to the .cl code at runtime (like changed the block size or some other constants) then recompiled itself the saved_embedded_ptx would just get overwritten. Same on using the prebuilt ptx file - I always wanted to load the same ptx file even though there should have been more than one. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 13217]
2014-08-14Adding stub function implementation to get EA_Ropa running on fermi branch.Wilson Fung
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 11311]
2014-08-14using rsync because of ssh having issues with concurreny copyingTim Rogers
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 11065]
2014-08-14Fixing the copy directoriesTim Rogers
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 11053]
2014-08-14Adding support to change the directory on the remote opencl compilationTim Rogers
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 11052]
2014-08-14Warnings cleanupTim Rogers
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10753]
2014-08-14Fix for Bug 109 - memory alignment should be 256 bytes.Inderpreet Singh
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10090]
2014-08-14Initial modifications to the opencl runtime API to allow the AMD sample ↵Tayler Hetherington
applications to run. Some of these modifications are temporary or incomplete and will be modified as necessary. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 9942]
2014-08-14Integration change. Bringing in the remote opencl changes from mem_divergenceTim Rogers
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 9764]
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-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-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:45:10 +0000