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

Add new kernel launching API for hip-clang

Browse Source 
[ROCm/hip commit: 3c80389584]
This commit is contained in:
Yaxun (Sam) Liu
2019-09-24 12:07:32 -04:00
parent 5e72734b0d
commit 9ef2687f45
2 changed files with 64 additions and 31 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+59 -28
View File
@@ -1519,34 +1519,6 @@ hipError_t hipMemcpyToSymbol(void*, const void*, size_t, size_t, hipMemcpyKind,
} // Namespace hip_impl.
#endif
#ifdef __cplusplus
extern "C" {
#endif
/**
* @brief C compliant kernel launch API
*
* @param [in] function_address - kernel function pointer.
* @param [in] numBlocks - number of blocks
* @param [in] dimBlocks - dimension of a block
* @param [in] args - kernel arguments
* @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel. The
* Kernel can access this with HIP_DYNAMIC_SHARED.
* @param [in] stream - Stream where the kernel should be dispatched. May be 0, in which case th
* default stream is used with associated synchronization rules.
*
* @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice
*
*/
hipError_t hipLaunchKernel(const void* function_address,
dim3 numBlocks, dim3 dimBlocks, void** args,
size_t sharedMemBytes, hipStream_t stream);
#ifdef __cplusplus
}
#endif
#if defined(__cplusplus)
extern "C" {
#endif
@@ -3055,6 +3027,65 @@ hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset);
hipError_t hipLaunchByPtr(const void* func);
/**
* @brief Push configuration of a kernel launch.
*
* @param [in] gridDim grid dimension specified as multiple of blockDim.
* @param [in] blockDim block dimensions specified in work-items
* @param [in] sharedMem Amount of dynamic shared memory to allocate for this kernel. The
* kernel can access this with HIP_DYNAMIC_SHARED.
* @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case the
* default stream is used with associated synchronization rules.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
*
*/
hipError_t __hipPushCallConfiguration(dim3 gridDim,
dim3 blockDim,
size_t sharedMem __dparm(0),
hipStream_t stream __dparm(0));
/**
* @brief Pop configuration of a kernel launch.
*
* @param [out] gridDim grid dimension specified as multiple of blockDim.
* @param [out] blockDim block dimensions specified in work-items
* @param [out] sharedMem Amount of dynamic shared memory to allocate for this kernel. The
* kernel can access this with HIP_DYNAMIC_SHARED.
* @param [out] stream Stream where the kernel should be dispatched. May be 0, in which case the
* default stream is used with associated synchronization rules.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
*
*/
hipError_t __hipPopCallConfiguration(dim3 *gridDim,
dim3 *blockDim,
size_t *sharedMem,
hipStream_t *stream);
/**
* @brief C compliant kernel launch API
*
* @param [in] function_address - kernel stub function pointer.
* @param [in] numBlocks - number of blocks
* @param [in] dimBlocks - dimension of a block
* @param [in] args - kernel arguments
* @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel. The
* Kernel can access this with HIP_DYNAMIC_SHARED.
* @param [in] stream - Stream where the kernel should be dispatched. May be 0, in which case th
* default stream is used with associated synchronization rules.
*
* @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice
*
*/
hipError_t hipLaunchKernel(const void* function_address,
dim3 numBlocks,
dim3 dimBlocks,
void** args,
size_t sharedMemBytes __dparm(0),
hipStream_t stream __dparm(0));
/**
* @}