Fix issues of missing kernel function symbols and missing argument list via
using __hipRegister* functions.
Then the following tests can pass,
directed_tests/runtimeApi/module/hipFuncGetAttributes
directed_tests/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice
directed_tests/gcc/LaunchKernel
Change-Id: I52135b61e8283eb4f9f10f77895151e4e55418d9
[ROCm/clr commit: 4a113bb669]
Implement _ihipGetGlobalVar() and ihipGetGlobalVar() to
get global variables.
Change-Id: I442ab6712e12306c3316f114f5dc42f6daefaad9
[ROCm/clr commit: 8e9e6a44a4]
Fix compilation error on hip-hcc+clang , hip-vdi+clang
Enabled hipExtLaunchMultiKernelMultiDevice test on hip-vdi path
hipExtLaunchMultiKernelMultiDevice common declaration for all paths
Change-Id: I76031840614fce8e12a8e845548fa43a389a741a
[ROCm/clr commit: 5a6c605730]
For hipLaunchKernelGGL(), hmod->kernargs is empty, thus we need
insert hmod->kernargs[name_str] which is empty.
Change-Id: I95f818d0525da84452e66c5778f0648a643843c7
[ROCm/clr commit: 05785772e3]
Fix the following issues:
1.Ignore hidden arguments of kernel functions.
2.Look up both origial function name and function name with .kd postfix
when argments are retrived from module.
3.Addition, fix compiling issue of LaunchKernel test app.
Change-Id: I9400943f2f02433cb4409b19c0cac3626c2bc454
[ROCm/clr commit: eb8c4d1ce2]
* Fix bug in LaunchKernel test
Instead of passing the address of the gpu buffer, pass the address
of the pointer that holds the address of the gpu buffer
* Fix hipLaunchKernel's kernarg buffer construction.
The hipLaunchKernel implementation should rely on ihipModuleLaunchKernel
to construct the kernarg buffer correctly based on kernel metadata.
* Fix a bug in get_functions where the Kernel_descriptor wasn't constructed with the correct kernarg layout information.
* Fix a bug in kernarg layout parsing dealing with kernel without any arg
* teach ihipModuleLaunchKernel to handle kernel without any arg
* Add a more interesting test
[ROCm/clr commit: 14e235378f]
* Detect when an explicit printf buffer flush is required
in a device/stream synchronization function.
* hip_module.cpp: add missing hc_am.hpp header
[ROCm/clr commit: 26b50e1e1b]
This PR is a follow-up on PR# #1698 and it makes two more APIs (hipLaunchCooperativeKernel/hipLaunchCooperativeKernelMultiDevice) inline so that they can work correctly with lazy binding.
[ROCm/clr commit: 857052be1e]
[Background] it was found that if lazy linking used for a library that calls hipExtLaunchMultiKernelMultiDevice API then this API can get the wrong program_state object for looking up device kernels leading to a "No device code available" error in this API.
To fix this issue, the API was refactored to be inline and get and pass the correct program_state to an internal hip API to request a multi-device kernel launch.
[ROCm/clr commit: 8eaea4d114]
* [hip] add support for implicit kernel argument for multi-grid sync
* modified code for calculating the prev_sum
* change the impCoopArg type to size_t
* add memory clean up
* launch init_gws and main kernels into two separate loops
[ROCm/clr commit: 93c688a0c9]
* guard against division by zero for no VGPR usage (e.g., in an empty kernel)
* fix bracket format
* clean up parenthesis
[ROCm/clr commit: d16963c9d5]
* [hip] add initial implementation for hipLaunchCooperativeKernel API
* [hip] use total number of work groups to initialize the GWS resource
* [hip] use only one argument for init_gws kernel
* [hip] use the device associated with the stream for checking the device properties
[ROCm/clr commit: 32ce882d6e]
* Add support for hipFunGetAttribute
* Support NVCC path
* Test using sample module_api_global
* Try fixing CI build failure due to hip_prof_gen scan
* Fix for CI build issue
* Resolve conflict
* Rebase and resolve conflicts with master
* Fix build error
* Fix NVCC path build error
[ROCm/clr commit: 8b6317d041]
* all thread local access now through single struct
* clean up old commented-out code, more use of GET_TLS()
* fewer calls to GET_TLS by passing tls as a funtion argument
* revert unnecessary change to printf
* fix failing tests due to TLS change
* fix merge conflicts in ihipOccupancyMaxActiveBlocksPerMultiprocessor
[ROCm/clr commit: f337ae1edb]
* Added support of hipOccupancyMaxActiveBlocksPerMultiprocessor & hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags APIs
* Taking into account of SGPR usage to determine the max active blocks in hipOccupancyMaxActiveBlocksPerMultiprocessor()
[ROCm/clr commit: 7b9801fe9a]
* Add hip init in hipExtLaunchMultiKernelMultiDevice
* Add more logstatus for multiple return paths
* Fix missing i in function name
[ROCm/clr commit: 8df47255c5]
* [hip] implement the hipExtLaunchMultiKernelMultiDevice API
* add a guard to check the HCC version for acquire_locked_hsa_queue() API which was introdued in HCC for ROCm 2.5
* modified code based on the requested changes
* changes to lock all streams before launching kernels for each device and unlock them after the dispatches
* check each stream to be valid before starting to lock all the streams
[ROCm/clr commit: d6ad690cb6]
* Implement the hipOccupancyMaxPotentialBlockSize function
* Replaced hipGetDeviceProperties() call by ihipGetDeviceProperties() in ihipOccupancyMaxPotentialBlockSize()
* Add test for hipOccupancyMaxPotentialBlockSize in Module API
* Added extern declaration for ihipGetDeviceProperties() to be accessed inside ihipOccupancyMaxPotentialBlockSize()
* fixed hipOccupancyMaxPotentialBlockSize test build issue
* Fix hipOccupancyMaxPotentialBlockSize dtest
* Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest
* Revert "Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest"
This reverts commit 0480ff56f1441fc515d2c26ce33783e303423938.
* Disable hipOccupancyMaxPotentialBlockSize dtest on NVCC
* move extern declaration of ihipGetDeviceProperties to hip_module.cpp
* Update the limiation of 32 wavefronts per CU and 800/512 SGPRs for VI/pre-VI chips to calculate the occupancy
[ROCm/clr commit: 28c34ead70]
Appended 48 empty bytes to the kernarg area at runtime. The implicit arguments are enabled primarily for the hostcall services
and it is completely abstracted from the user code. Enabled it for both hip-clang and hip-hcc.
[ROCm/clr commit: 214ec53da3]
Use the code object manager library to parse the code object metadata. Both
code object v2 and v3 formats are now supported for HCC generated binaries.
[ROCm/clr commit: de89102528]
* Add ersatz for NVRTC.
* Fix extraneous paren and use correct namespace.
* Use lowerCamelCase (yuck, yuck) consistently.
* Link against FS when building hiprtc lib.
* Correctly mark Manipulators. Fix dual compile.
* Add unit tests. Extend HIT to accept linker options.
* Make sure the HIPRTC library is installed.
* Better logging. Try to auto-detect the target.
* Stop specifying the target explicitly.
* Add missing flavour of `hipModuleLaunchKernel`.
* Program was already destroyed.
* Don't use `--genco`. Fix mangled name trimming.
* Fix HIPRTC breakage due to upstream noise.
* [dtests] Replace RUN -> TEST in hiprtc tests
Change-Id: Ie499e92dfe4e5c94634b1c2b76cf52d241bcfea3
* [hit] Set HIP_PATH to HIP_ROOT_DIR for all tests
Change-Id: Ib0ad1f99bc71c03e363e055dd508a7a4a210680a
[ROCm/clr commit: a538eb705a]
Tao Sang
Lad, Aditya
agodavar
Rahul Garg
Sarbojit2019
Maneesh Gupta
Siu Chi Chan
satyanveshd
Aryan Salmanpour
Alex Voicu
searlmc1
Nick Curtis
Jeff Daily
Rahul Garg
wkwchau
cdevadas
Aryan Salmanpour
cdevadas
Laurent Morichetti