From ffcfc9536078569ada49e51a79f85ffcd572b86e Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Sat, 18 Jun 2016 11:28:20 -0500 Subject: [PATCH] able to pass non-dim launch parm to kernel launch Change-Id: I0411849a27efcba597a1a9aa08be179635e04988 --- hipamd/include/hcc_detail/hip_runtime.h | 13 ++---- hipamd/src/hip_hcc.cpp | 62 ++++++++++++++++++++++++- hipamd/tests/src/CMakeLists.txt | 3 +- hipamd/tests/src/hipLaunchParm.cpp | 36 ++++++++++++++ 4 files changed, 104 insertions(+), 10 deletions(-) create mode 100644 hipamd/tests/src/hipLaunchParm.cpp diff --git a/hipamd/include/hcc_detail/hip_runtime.h b/hipamd/include/hcc_detail/hip_runtime.h index b0b6ecfe6c..1b536853d6 100644 --- a/hipamd/include/hcc_detail/hip_runtime.h +++ b/hipamd/include/hcc_detail/hip_runtime.h @@ -529,7 +529,10 @@ __device__ float __dsqrt_rz(double x); #define HIP_KERNEL_NAME(...) __VA_ARGS__ #ifdef __HCC_CPP__ -hipStream_t ihipPreLaunchKernel(hipStream_t stream, grid_launch_parm *lp); +hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp); +hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp); +hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp); +hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp); void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp); // TODO - move to common header file. @@ -540,14 +543,8 @@ void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp); #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \ do {\ grid_launch_parm lp;\ - lp.gridDim.x = _numBlocks3D.x; \ - lp.gridDim.y = _numBlocks3D.y; \ - lp.gridDim.z = _numBlocks3D.z; \ - lp.groupDim.x = _blockDim3D.x; \ - lp.groupDim.y = _blockDim3D.y; \ - lp.groupDim.z = _blockDim3D.z; \ lp.groupMemBytes = _groupMemBytes; \ - hipStream_t trueStream = (ihipPreLaunchKernel(_stream, &lp)); \ + hipStream_t trueStream = (ihipPreLaunchKernel(_stream, _numBlocks3D, _blockDim3D, &lp)); \ if (HIP_TRACE_API) {\ fprintf(stderr, KGRN "<gridDim.x = grid.x; + lp->gridDim.y = grid.y; + lp->gridDim.z = grid.z; + lp->groupDim.x = block.x; + lp->groupDim.y = block.y; + lp->groupDim.z = block.z; + stream->lockopen_preKernelCommand(); +// *av = &stream->_av; + lp->av = &stream->_av; + lp->cf = new hc::completion_future; +// lp->av = static_cast(av); +// lp->cf = static_cast(malloc(sizeof(hc::completion_future))); + return (stream); +} +hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp) +{ + std::call_once(hip_initialized, ihipInit); + stream = ihipSyncAndResolveStream(stream); + lp->gridDim.x = grid; + lp->gridDim.y = 0; + lp->gridDim.z = 0; + lp->groupDim.x = block.x; + lp->groupDim.y = block.y; + lp->groupDim.z = block.z; + stream->lockopen_preKernelCommand(); +// *av = &stream->_av; + lp->av = &stream->_av; + lp->cf = new hc::completion_future; +// lp->av = static_cast(av); +// lp->cf = static_cast(malloc(sizeof(hc::completion_future))); + return (stream); +} +hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp) +{ + std::call_once(hip_initialized, ihipInit); + stream = ihipSyncAndResolveStream(stream); + lp->gridDim.x = grid.x; + lp->gridDim.y = grid.y; + lp->gridDim.z = grid.z; + lp->groupDim.x = block; + lp->groupDim.y = 0; + lp->groupDim.z = 0; + stream->lockopen_preKernelCommand(); +// *av = &stream->_av; + lp->av = &stream->_av; + lp->cf = new hc::completion_future; +// lp->av = static_cast(av); +// lp->cf = static_cast(malloc(sizeof(hc::completion_future))); + return (stream); +} +hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp) +{ + std::call_once(hip_initialized, ihipInit); + stream = ihipSyncAndResolveStream(stream); + lp->gridDim.x = grid; + lp->gridDim.y = 0; + lp->gridDim.z = 0; + lp->groupDim.x = block; + lp->groupDim.y = 0; + lp->groupDim.z = 0; stream->lockopen_preKernelCommand(); // *av = &stream->_av; lp->av = &stream->_av; diff --git a/hipamd/tests/src/CMakeLists.txt b/hipamd/tests/src/CMakeLists.txt index 2734abb911..eed4356405 100644 --- a/hipamd/tests/src/CMakeLists.txt +++ b/hipamd/tests/src/CMakeLists.txt @@ -178,10 +178,11 @@ build_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp) build_hip_executable (hipPeerToPeer_simple hipPeerToPeer_simple.cpp) build_hip_executable (hipTestMemcpyPin hipTestMemcpyPin.cpp) #build_hip_executable (hipDynamicShared hipDynamicShared.cpp) - +build_hip_executable (hipLaunchParm hipLaunchParm.cpp) make_test(hipEventRecord --iterations 10) make_test(hipEnvVarDriver " " ) +make_test(hipLaunchParm " ") #TODO -reenable #make_test(hipPointerAttrib " " ) diff --git a/hipamd/tests/src/hipLaunchParm.cpp b/hipamd/tests/src/hipLaunchParm.cpp new file mode 100644 index 0000000000..c6d28fcd3a --- /dev/null +++ b/hipamd/tests/src/hipLaunchParm.cpp @@ -0,0 +1,36 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include"hip_runtime.h" +#include"test_common.h" +#include"hip_runtime_api.h" +#include + +__global__ void vAdd(hipLaunchParm lp, float *a){} + +int main() +{ + float *Ad; + hipMalloc((void**)&Ad, 1024); + hipLaunchKernel(vAdd, 1024, 1, 0, 0, Ad); + hipLaunchKernel(vAdd, 1024, dim3(1), 0, 0, Ad); + hipLaunchKernel(vAdd, dim3(1024), 1, 0, 0, Ad); + hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad); + passed(); +}