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


[ROCm/hip commit: d8a344113f]

projects/hip/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;
+}

projects/hip/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

projects/hip/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

projects/hip/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:
     *;
 };

projects/hip/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<cl_mem>(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<cl_mem>(const_cast<void*>(src)))->asBuffer();
-  amd::Buffer* dstBuffer = as_amd(reinterpret_cast<cl_mem>(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<cl_mem>(const_cast<void*>(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<cl_mem>(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;
+}

projects/hip/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);

projects/hip/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<const void*, hipFunction_t> 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<hipFunction_t>(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<const void*, hipFunction_t>::iterator it;
   if ((it = g_functions.find(hostFunction)) == g_functions.end())

projects/hip/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 <hip/hip_runtime.h>
+
+#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;
+}
+
+