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

P4 to Git Change 2003445 by yaxunl@yaxunl-lc10 on 2019/09/24 16:56:01

Browse Source 
SWDEV-145570 - Add new kernel launching API for hip-clang

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#28 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#27 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#41 edit
This commit is contained in:
foreman
2019-09-24 16:58:14 -04:00
parent 4f6bbfa473
commit 35c8a50c99
3 changed files with 57 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/api/hip/hip_platform.cpp
+51
View File
@@ -411,6 +411,35 @@ extern "C" hipError_t hipConfigureCall(
HIP_RETURN(hipSuccess);
}
extern "C" hipError_t __hipPushCallConfiguration(
dim3 gridDim,
dim3 blockDim,
size_t sharedMem,
hipStream_t stream)
{
HIP_INIT_API(gridDim, blockDim, sharedMem, stream);
PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream);
HIP_RETURN(hipSuccess);
}
extern "C" hipError_t __hipPopCallConfiguration(dim3 *gridDim,
dim3 *blockDim,
size_t *sharedMem,
hipStream_t *stream) {
HIP_INIT_API(gridDim, blockDim, sharedMem, stream);
ihipExec_t exec;
PlatformState::instance().popExec(exec);
*gridDim = exec.gridDim_;
*blockDim = exec.blockDim_;
*sharedMem = exec.sharedMem_;
*stream = exec.hStream_;
HIP_RETURN(hipSuccess);
}
extern "C" hipError_t hipSetupArgument(
const void *arg,
size_t size,
@@ -449,6 +478,28 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction)
exec.sharedMem_, exec.hStream_, nullptr, extra));
}
extern "C" hipError_t hipLaunchKernel(const void *hostFunction,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMemBytes,
hipStream_t stream)
{
HIP_INIT_API(hostFunction, gridDim, blockDim, args, sharedMemBytes,
stream);
int deviceId = ihipGetDevice();
hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId);
if (func == nullptr) {
HIP_RETURN(hipErrorUnknown);
}
HIP_RETURN(hipModuleLaunchKernel(func, gridDim.x, gridDim.y, gridDim.z,
blockDim.x, blockDim.y, blockDim.z,
sharedMemBytes, stream, args, nullptr));
}
hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
size_t size = 0;
if(!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), devPtr, &size)) {