From 8a53697b6c11d6de4fcd482270950d50d6ef88c2 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 11 Oct 2019 02:05:06 +0530 Subject: [PATCH] Add hipExtLaunchKernelGGL --- .../hip/hcc_detail/functional_grid_launch.hpp | 33 ++++++-- tests/src/kernel/hipExtLaunchKernelGGL.cpp | 83 +++++++++++++++++++ 2 files changed, 111 insertions(+), 5 deletions(-) create mode 100644 tests/src/kernel/hipExtLaunchKernelGGL.cpp diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index c493eec933..b06efc340a 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -26,6 +26,7 @@ THE SOFTWARE. #include "helpers.hpp" #include "program_state.hpp" #include "hip_runtime_api.h" +#include "hip/hip_hcc.h" #include #include @@ -115,14 +116,15 @@ void hipLaunchKernelGGLImpl( const dim3& dimBlocks, std::uint32_t sharedMemBytes, hipStream_t stream, - void** kernarg) { + void** kernarg, hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags){ const auto& kd = hip_impl::get_program_state().kernel_descriptor(function_address, target_agent(stream)); - hipModuleLaunchKernel(kd, numBlocks.x, numBlocks.y, numBlocks.z, - dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, - stream, nullptr, kernarg); + hipExtModuleLaunchKernel(kd, (numBlocks.x * dimBlocks.x) , + (numBlocks.y * dimBlocks.y), (numBlocks.z * dimBlocks.z), + dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, + stream, nullptr, kernarg, startEvent, stopEvent, flags); } } // Namespace hip_impl. @@ -165,6 +167,27 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, auto kernarg = hip_impl::make_kernarg(kernel, std::tuple{std::move(args)...}); std::size_t kernarg_size = kernarg.size(); + void* config[]{ + HIP_LAUNCH_PARAM_BUFFER_POINTER, + kernarg.data(), + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &kernarg_size, + HIP_LAUNCH_PARAM_END}; + + hip_impl::hipLaunchKernelGGLImpl(reinterpret_cast(kernel), + numBlocks, dimBlocks, sharedMemBytes, + stream, &config[0], nullptr, nullptr, 0); +} + +template +inline +void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, + hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags, Args... args) { + hip_impl::hip_init(); + auto kernarg = hip_impl::make_kernarg(kernel, std::tuple{std::move(args)...}); + std::size_t kernarg_size = kernarg.size(); + void* config[]{ HIP_LAUNCH_PARAM_BUFFER_POINTER, kernarg.data(), @@ -174,7 +197,7 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, hip_impl::hipLaunchKernelGGLImpl(reinterpret_cast(kernel), numBlocks, dimBlocks, sharedMemBytes, - stream, &config[0]); + stream, &config[0], startEvent , stopEvent, flags); } #pragma GCC visibility pop diff --git a/tests/src/kernel/hipExtLaunchKernelGGL.cpp b/tests/src/kernel/hipExtLaunchKernelGGL.cpp new file mode 100644 index 0000000000..52ac671818 --- /dev/null +++ b/tests/src/kernel/hipExtLaunchKernelGGL.cpp @@ -0,0 +1,83 @@ +/* +Copyright (c) 2015-2017 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. +*/ +// Test the Grid_Launch syntax. + +/* HIT_START + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "test_common.h" + + +// __device__ maps to __attribute__((hc)) +__device__ int foo(int i) { return i + 1; } + +//--- +// Syntax we would like to support with GRID_LAUNCH enabled: +template +__global__ void vectorADD2(T* A_d, T* B_d, T* C_d, size_t N) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < N; i += stride) { + double foo = __hiloint2double(A_d[i], B_d[i]); + C_d[i] = __double2loint(foo) + __double2hiint(foo); // A_d[i] + B_d[i] ; + } +} + +int test_gl2(size_t N) { + size_t Nbytes = N * sizeof(int); + + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + + // Full vadd in one large chunk, to get things started: + HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipExtLaunchKernelGGL(vectorADD2, dim3(blocks), dim3(threadsPerBlock),0,0,nullptr,nullptr,0, A_d, B_d, C_d, N); + + HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIPCHECK(hipDeviceSynchronize()); + + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + return 0; +} + +int main(int argc, char* argv[]) { + HipTest::parseStandardArguments(argc, argv, true); + + test_gl2(N); + + passed(); +}