diff --git a/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp b/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp index e4a5d79f7d..c493eec933 100644 --- a/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp @@ -177,13 +177,4 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, stream, &config[0]); } -template -[[deprecated("hipLaunchKernel is deprecated and will be removed in the next " - "version of HIP; please upgrade to hipLaunchKernelGGL.")]] -inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks, - std::uint32_t groupMemBytes, hipStream_t stream, Args... args) { - hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{}, - std::move(args)...); -} - #pragma GCC visibility pop diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 97ff4bd1d7..c93aa67994 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1520,6 +1520,34 @@ 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 diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime.h b/hipamd/include/hip/nvcc_detail/hip_runtime.h index 3a18f11614..d159e116c0 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime.h @@ -31,11 +31,6 @@ THE SOFTWARE. typedef int hipLaunchParm; -#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ - do { \ - kernelName<<>>(0, ##__VA_ARGS__); \ - } while (0) - #define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ do { \ kernelName<<>>(__VA_ARGS__); \ diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 532e1d827e..6150969b25 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -1241,6 +1241,13 @@ inline static hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* im cuModuleLoadDataEx(module, image, numOptions, options, optionValues)); } +inline static hipError_t hipLaunchKernel(const void* function_address, dim3 numBlocks, + dim3 dimBlocks, void** args, size_t sharedMemBytes, + hipStream_t stream) +{ + return hipCUDAErrorTohipError(cudaLaunchKernel(function_address,numBlocks,dimBlocks,args,sharedMemBytes,stream)); +} + inline static hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 50b3cc9b04..2548660a72 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -1238,3 +1238,31 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor( tls, numBlocks, f, blockSize, dynSharedMemPerBlk)); } + +hipError_t hipLaunchKernel( + const void* func_addr, dim3 numBlocks, dim3 dimBlocks, void** args, + size_t sharedMemBytes, hipStream_t stream) +{ + HIP_INIT_API(hipLaunchKernel,func_addr,numBlocks,dimBlocks,args,sharedMemBytes,stream); + + hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)func_addr, + hip_impl::target_agent(stream)); + + if(kd == nullptr || kd->_header == nullptr) + return ihipLogStatus(hipErrorInvalidValue); + + size_t szKernArg = kd->_header->kernarg_segment_byte_size; + + if(args == NULL && szKernArg != 0) + return ihipLogStatus(hipErrorInvalidValue); + + void* config[]{ + HIP_LAUNCH_PARAM_BUFFER_POINTER, + args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &szKernArg, + HIP_LAUNCH_PARAM_END}; + + return ihipLogStatus(ihipModuleLaunchKernel(tls, kd, numBlocks.x * dimBlocks.x, numBlocks.y * dimBlocks.y, numBlocks.z * dimBlocks.z, + dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, stream, nullptr, (void**)&config, nullptr, nullptr, 0)); +} diff --git a/hipamd/tests/src/gcc/LaunchKernel.c b/hipamd/tests/src/gcc/LaunchKernel.c new file mode 100644 index 0000000000..4ebe59d7c5 --- /dev/null +++ b/hipamd/tests/src/gcc/LaunchKernel.c @@ -0,0 +1,165 @@ + +/* Copyright (c) 2019-Present 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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + + +/* HIT_START + * BUILD_CMD: gpu.o %hc -I%hip-path/include -g -c %S/gpu.cpp -o %T/gpu.o EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: launchkernel.o %cc -D__HIP_PLATFORM_HCC__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: LaunchKernel %hc %T/launchkernel.o %T/gpu.o -g -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t DEPENDS gpu.o launchkernel.o EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t EXCLUDE_HIP_PLATFORM nvcc + * HIT_END + */ + + +#include "hip/hip_runtime.h" +#include +#include "LaunchKernel.h" + +#define HIPCHECK(error) \ + { \ + hipError_t localError = error; \ + if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + printf("%serror: '%s'(%d) from %s at %s:%d%s\n", "\x1B[31m", hipGetErrorString(localError), \ + localError, #error, __FILE__, __LINE__, "\x1B[0m"); \ + return false; \ + } \ + } + + +bool LaunchKernelArg() +{ + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + HIPCHECK(hipLaunchKernel(kernel, blocks, threads,NULL, 0, 0)); + + return true; +} + +bool LaunchKernelArg1() +{ + int A = 0; + int *A_d = NULL; + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + void* Args[]={A_d}; + HIPCHECK(hipLaunchKernel(kernel1, blocks, threads, Args,0,0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + + if(A != 333) + return false; + + return true; +} + +bool LaunchKernelArg2() +{ + int A = 0; + int B = 123; + int *A_d = NULL; + int *B_d = NULL; + + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + HIPCHECK(hipMalloc((void**)&B_d, sizeof(int))); + + // Copy data from host memory to device memory + HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice)); + + void* Args[]={A_d,B_d}; + HIPCHECK(hipLaunchKernel(kernel2, blocks, threads, Args,0,0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(B_d)); + + if(A != 123) + return false; + + return true; +} + +bool LaunchKernelArg3() +{ + int A = 321; + int B = 123; + int C = 0; + int *A_d = NULL; + int *B_d = NULL; + int *C_d = NULL; + + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + HIPCHECK(hipMalloc((void**)&B_d, sizeof(int))); + + HIPCHECK(hipMalloc((void**)&C_d, sizeof(int))); + + // Copy data from host memory to device memory + HIPCHECK(hipMemcpy(A_d,&A, sizeof(int), hipMemcpyHostToDevice)); + + HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice)); + + void* Args[]={A_d,B_d,C_d}; + HIPCHECK(hipLaunchKernel(kernel3, blocks, threads, Args,0,0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&C, C_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(B_d)); + HIPCHECK(hipFree(C_d)); + + if(C != 444) + return false; + + return true; +} + + +int main() +{ + + if( LaunchKernelArg() && + LaunchKernelArg1() && + LaunchKernelArg2() && + LaunchKernelArg3()) + { + printf("PASSED!\n"); + } + else + printf("FAILED\n"); +} diff --git a/hipamd/tests/src/gcc/LaunchKernel.h b/hipamd/tests/src/gcc/LaunchKernel.h new file mode 100644 index 0000000000..b7424d8b46 --- /dev/null +++ b/hipamd/tests/src/gcc/LaunchKernel.h @@ -0,0 +1,23 @@ + +/* Copyright (c) 2019-Present 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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +extern __global__ void kernel(); +extern __global__ void kernel1(int*); +extern __global__ void kernel2(int*,int*); +extern __global__ void kernel3(int*,int*,int*); diff --git a/hipamd/tests/src/gcc/gpu.cpp b/hipamd/tests/src/gcc/gpu.cpp new file mode 100644 index 0000000000..00223478fa --- /dev/null +++ b/hipamd/tests/src/gcc/gpu.cpp @@ -0,0 +1,46 @@ +/* + * Copyright (c) 2019-Present 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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * */ + + +#include + +extern "C" +{ + +__global__ void kernel() +{ + int a = 20; +} + +__global__ void kernel1(int *a) +{ + *a = 333; +} + +__global__ void kernel2(int *a, int*b) +{ + *a = *b; +} + +__global__ void kernel3(int *a, int*b, int* c) +{ + *c = *a+*b; +} + +}//extern "C"