From f27e6709ebc7e9242dea3376ec5e2c90501ea1cf Mon Sep 17 00:00:00 2001 From: foreman Date: Fri, 2 Mar 2018 17:55:48 -0500 Subject: [PATCH] P4 to Git Change 1522211 by lmoriche@lmoriche_opencl_dev2 on 2018/03/02 17:41:47 SWDEV-145570 - [HIP] - Hip Rearchitecture - Rename cuda* launch functions -> hip* - Add more function prototypes to compile the HIP tests Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_device.cpp#3 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#4 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#1 add --- hipamd/api/hip/hip_context.cpp | 13 ++++++ hipamd/api/hip/hip_device.cpp | 19 +++++++++ hipamd/api/hip/hip_hcc.def.in | 14 +++--- hipamd/api/hip/hip_hcc.map.in | 14 +++--- hipamd/api/hip/hip_memory.cpp | 56 +++++++++++++++++++++--- hipamd/api/hip/hip_module.cpp | 19 +++++++++ hipamd/api/hip/hip_platform.cpp | 14 +++--- hipamd/api/hip/hip_stream.cpp | 76 +++++++++++++++++++++++++++++++++ 8 files changed, 199 insertions(+), 26 deletions(-) create mode 100644 hipamd/api/hip/hip_stream.cpp diff --git a/hipamd/api/hip/hip_context.cpp b/hipamd/api/hip/hip_context.cpp index 0e6ff2116a..78e65e99c5 100644 --- a/hipamd/api/hip/hip_context.cpp +++ b/hipamd/api/hip/hip_context.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. #include "hip_internal.hpp" #include "platform/runtime.hpp" +#include "utils/versions.hpp" amd::Context* g_context = nullptr; @@ -56,3 +57,15 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) return hipSuccess; } +hipError_t hipRuntimeGetVersion(int *runtimeVersion) +{ + HIP_INIT_API(runtimeVersion); + + if (!runtimeVersion) { + return hipErrorInvalidValue; + } + + *runtimeVersion = AMD_PLATFORM_BUILD_NUMBER; + + return hipSuccess; +} diff --git a/hipamd/api/hip/hip_device.cpp b/hipamd/api/hip/hip_device.cpp index 7296eabb17..b5da0c34b5 100644 --- a/hipamd/api/hip/hip_device.cpp +++ b/hipamd/api/hip/hip_device.cpp @@ -358,6 +358,25 @@ hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device) { return hipSuccess; } +hipError_t hipSetDevice(int deviceId) +{ + HIP_INIT_API(deviceId); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + +hipError_t hipDeviceReset(void) +{ + HIP_INIT_API(); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + + hipError_t hipDeviceSynchronize(void) { // FIXME: should wait on all streams diff --git a/hipamd/api/hip/hip_hcc.def.in b/hipamd/api/hip/hip_hcc.def.in index 6b4793ea94..10113dc2bd 100644 --- a/hipamd/api/hip/hip_hcc.def.in +++ b/hipamd/api/hip/hip_hcc.def.in @@ -120,10 +120,10 @@ hipStreamGetFlags hipStreamQuery hipStreamSynchronize hipStreamWaitEvent -__cudaRegisterFatBinary -__cudaRegisterFunction -__cudaRegisterVariable -__cudaUnregisterFatBinary -cudaConfigureCall -cudaSetupArgument -cudaLaunch +__hipRegisterFatBinary +__hipRegisterFunction +__hipRegisterVariable +__hipUnregisterFatBinary +hipConfigureCall +hipSetupArgument +hipLaunchByPtr diff --git a/hipamd/api/hip/hip_hcc.map.in b/hipamd/api/hip/hip_hcc.map.in index e4025606bc..a4153ee56f 100644 --- a/hipamd/api/hip/hip_hcc.map.in +++ b/hipamd/api/hip/hip_hcc.map.in @@ -121,13 +121,13 @@ global: hipStreamQuery; hipStreamSynchronize; hipStreamWaitEvent; - __cudaRegisterFatBinary; - __cudaRegisterFunction; - __cudaRegisterVariable; - __cudaUnregisterFatBinary; - cudaConfigureCall; - cudaSetupArgument; - cudaLaunch; + __hipRegisterFatBinary; + __hipRegisterFunction; + __hipRegisterVariable; + __hipUnregisterFatBinary; + hipConfigureCall; + hipSetupArgument; + hipLaunchByPtr; local: *; }; diff --git a/hipamd/api/hip/hip_memory.cpp b/hipamd/api/hip/hip_memory.cpp index 2dba003ba6..0911f61e4c 100644 --- a/hipamd/api/hip/hip_memory.cpp +++ b/hipamd/api/hip/hip_memory.cpp @@ -54,6 +54,15 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) return hipSuccess; } +hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) +{ + HIP_INIT_API(ptr, sizeBytes, flags); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + hipError_t hipFree(void* ptr) { if (!is_valid(reinterpret_cast(ptr))) { @@ -63,6 +72,20 @@ hipError_t hipFree(void* ptr) return hipSuccess; } +hipError_t hipMemcpyAsync(void* dst, + const void* src, + size_t sizeBytes, + hipMemcpyKind kind, + hipStream_t stream) +{ + HIP_INIT_API(dst, src, sizeBytes, kind, stream); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + + hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { HIP_INIT_API(dst, src, sizeBytes, kind); @@ -76,20 +99,17 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind return hipErrorOutOfMemory; } - amd::Buffer* srcBuffer = as_amd(reinterpret_cast(const_cast(src)))->asBuffer(); - amd::Buffer* dstBuffer = as_amd(reinterpret_cast(dst))->asBuffer(); - amd::Command* command; amd::Command::EventWaitList waitList; switch (kind) { case hipMemcpyDeviceToHost: command = new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER, waitList, - *srcBuffer, 0, sizeBytes, dst); + *as_amd(reinterpret_cast(const_cast(src)))->asBuffer(), 0, sizeBytes, dst); break; case hipMemcpyHostToDevice: command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER, waitList, - *dstBuffer, 0, sizeBytes, src); + *as_amd(reinterpret_cast(dst))->asBuffer(), 0, sizeBytes, src); break; default: assert(!"Shouldn't reach here"); @@ -117,3 +137,29 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind return hipSuccess; } +hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream ) +{ + HIP_INIT_API(dst, value, sizeBytes, stream); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + +hipError_t hipMemset(void* dst, int value, size_t sizeBytes) +{ + HIP_INIT_API(dst, value, sizeBytes); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + +hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) +{ + HIP_INIT_API(ptr, size); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} diff --git a/hipamd/api/hip/hip_module.cpp b/hipamd/api/hip/hip_module.cpp index dde0c4e790..fd7729c6e5 100644 --- a/hipamd/api/hip/hip_module.cpp +++ b/hipamd/api/hip/hip_module.cpp @@ -47,6 +47,25 @@ static uint64_t ElfSize(const void *emi) return total_size; } +hipError_t hipModuleLoad(hipModule_t *module, const char *fname) +{ + HIP_INIT_API(module, fname); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + + +hipError_t hipModuleUnload(hipModule_t hmod) +{ + HIP_INIT_API(hmod); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + hipError_t hipModuleLoadData(hipModule_t *module, const void *image) { HIP_INIT_API(module, image); diff --git a/hipamd/api/hip/hip_platform.cpp b/hipamd/api/hip/hip_platform.cpp index aed3342483..0cc6a3b1c2 100644 --- a/hipamd/api/hip/hip_platform.cpp +++ b/hipamd/api/hip/hip_platform.cpp @@ -54,7 +54,7 @@ struct __CudaPartHeader{ unsigned int subarch; }; -extern "C" hipModule_t __cudaRegisterFatBinary(void* bundle) +extern "C" hipModule_t __hipRegisterFatBinary(void* bundle) { if (!amd::Runtime::initialized()) { // FIXME: fix initialization hipInit(0); @@ -96,7 +96,7 @@ extern "C" hipModule_t __cudaRegisterFatBinary(void* bundle) std::map g_functions; -extern "C" void __cudaRegisterFunction( +extern "C" void __hipRegisterFunction( hipModule_t module, const void* hostFunction, char* deviceFunction, @@ -120,7 +120,7 @@ extern "C" void __cudaRegisterFunction( g_functions.insert(std::make_pair(hostFunction, reinterpret_cast(as_cl(kernel)))); } -extern "C" void __cudaRegisterVar( +extern "C" void __hipRegisterVar( hipModule_t module, char* hostVar, char* deviceVar, @@ -132,7 +132,7 @@ extern "C" void __cudaRegisterVar( { } -extern "C" void __cudaUnregisterFatBinary( +extern "C" void __hipUnregisterFatBinary( hipModule_t module ) { @@ -143,7 +143,7 @@ dim3 g_blockDim; // FIXME: place in execution stack size_t g_sharedMem; // FIXME: place in execution stack hipStream_t g_stream; // FIXME: place in execution stack -extern "C" hipError_t cudaConfigureCall( +extern "C" hipError_t hipConfigureCall( dim3 gridDim, dim3 blockDim, size_t sharedMem, @@ -161,7 +161,7 @@ extern "C" hipError_t cudaConfigureCall( char* g_arguments[1024]; // FIXME: needs to grow -extern "C" hipError_t cudaSetupArgument( +extern "C" hipError_t hipSetupArgument( const void *arg, size_t size, size_t offset) @@ -172,7 +172,7 @@ extern "C" hipError_t cudaSetupArgument( return hipSuccess; } -extern "C" hipError_t cudaLaunch(const void *hostFunction) +extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) { std::map::iterator it; if ((it = g_functions.find(hostFunction)) == g_functions.end()) diff --git a/hipamd/api/hip/hip_stream.cpp b/hipamd/api/hip/hip_stream.cpp new file mode 100644 index 0000000000..efecb5174d --- /dev/null +++ b/hipamd/api/hip/hip_stream.cpp @@ -0,0 +1,76 @@ +/* +Copyright (c) 2015 - 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 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 + +#include "hip_internal.hpp" + +hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) +{ + HIP_INIT_API(stream, flags); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + + +hipError_t hipStreamCreate(hipStream_t *stream) +{ + HIP_INIT_API(stream); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + + +hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) +{ + HIP_INIT_API(stream, flags); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + + +hipError_t hipStreamSynchronize(hipStream_t stream) +{ + HIP_INIT_API(stream); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + + +hipError_t hipStreamDestroy(hipStream_t stream) +{ + HIP_INIT_API(stream); + + assert(0 && "Unimplemented"); + + return hipErrorUnknown; +} + +