AI/rocm-systems
2
0 Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[hip] implement the hipExtLaunchMultiKernelMultiDevice API (#1165)

Browse Source 
* [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/hip commit: 96dc74897d]
This commit is contained in:
Aryan Salmanpour authored and Maneesh Gupta committed 2019-06-20 05:59:05 +05:30
1 parent 40bd111519
commit 45fa752888
5 files changed
+113 -29
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

projects/hip/include/hip/hcc_detail/hip_runtime.h
+5 -5
View File
@@ -270,14 +270,14 @@ static inline __device__ void printf(const char* format, All... all) {}
#if defined __HCC_CPP__
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block,
grid_launch_parm* lp, const char* kernelNameStr);
grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0);
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block,
grid_launch_parm* lp, const char* kernelNameStr);
grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0);
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block,
grid_launch_parm* lp, const char* kernelNameStr);
grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0);
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block,
grid_launch_parm* lp, const char* kernelNameStr);
extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp);
grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0);
extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp, bool unlockPostponed = 0);
#if GENERIC_GRID_LAUNCH == 0
//#warning "Original hipLaunchKernel defined"
projects/hip/include/hip/hcc_detail/program_state.hpp
+1
View File
@@ -77,6 +77,7 @@ class kernargs_size_align {
public:
std::size_t size(std::size_t n) const;
std::size_t alignment(std::size_t n) const;
const void* getHandle() const {return handle;};
private:
const void* handle;
friend kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t);