From 89b576da656f3d603249cfc91d0990cc967d062d Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 11 Oct 2016 12:09:58 -0500 Subject: [PATCH] Added feature for memcpy to Symbol 1. Currently works only for __attribute__((addrspace(1)) 2. Need to pass in string for name of the variable 3. Added test to check functionality Change-Id: I4c3cc1bf151cb5423e4aef59fcc4ad5693b31641 --- hipamd/include/hip/hcc_detail/hip_hcc.h | 2 +- hipamd/src/hip_memory.cpp | 25 ++++----- .../src/deviceLib/hipTestDeviceSymbol.cpp | 54 +++++++++++++++++++ 3 files changed, 64 insertions(+), 17 deletions(-) create mode 100644 hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp diff --git a/hipamd/include/hip/hcc_detail/hip_hcc.h b/hipamd/include/hip/hcc_detail/hip_hcc.h index f3967247a6..1d067432ea 100644 --- a/hipamd/include/hip/hcc_detail/hip_hcc.h +++ b/hipamd/include/hip/hcc_detail/hip_hcc.h @@ -29,7 +29,7 @@ THE SOFTWARE. #error("This version of HIP requires a newer version of HCC."); #endif -// #define USE_MEMCPYTOSYMBOL +#define USE_MEMCPYTOSYMBOL // diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 5443c43344..278502973a 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -102,13 +102,13 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) HIP_INIT_API(ptr, sizeBytes); hipError_t hip_status = hipSuccess; - // return NULL pointer when malloc size is 0 + // return NULL pointer when malloc size is 0 if (sizeBytes == 0) { *ptr = NULL; return ihipLogStatus(hipSuccess); } - + auto ctx = ihipGetTlsDefaultCtx(); if (ctx) { @@ -185,7 +185,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) }; // width in bytes -hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) +hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) { HIP_INIT_API(ptr, pitch, width, height); @@ -228,7 +228,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height return ihipLogStatus(hip_status); } -hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) +hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { hipChannelFormatDesc cd; cd.x = x; cd.y = y; cd.z = z; cd.w = w; @@ -237,7 +237,7 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel } hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, - size_t width, size_t height, unsigned int flags) + size_t width, size_t height, unsigned int flags) { HIP_INIT_API(array, desc, width, height, flags); @@ -396,7 +396,9 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou //int depSignalCnt = ctx._default_stream->preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D); assert(0); // Need to properly synchronize the copy - do something with depSignal if != NULL. - ctx->_acc.memcpy_symbol(symbolName, (void*) src,count, offset); + hc::accelerator acc = ctx->getDevice()->_acc; + + acc.memcpy_symbol(symbolName, (void*) src,count, offset); #endif return ihipLogStatus(hipSuccess); } @@ -715,7 +717,7 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, // TODO - make member function of stream? template hc::completion_future -ihipMemsetKernel(hipStream_t stream, +ihipMemsetKernel(hipStream_t stream, LockedAccessor_StreamCrit_t &crit, T * ptr, T val, size_t sizeBytes) { @@ -969,16 +971,7 @@ hipError_t hipFreeArray(hipArray* array) return ihipLogStatus(hipStatus); } -// Stubs of threadfence operations -__device__ void __threadfence_block(void){ - // no-op -} - -__device__ void __threadfence(void){ - // no-op -} __device__ void __threadfence_system(void){ // no-op } - diff --git a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp new file mode 100644 index 0000000000..359bc5d6db --- /dev/null +++ b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -0,0 +1,54 @@ +/* +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 +#include +#include + +#define NUM 1024 +#define SIZE 1024*4 + +__attribute__((address_space(1))) int global[NUM]; + + +__global__ void Assign(hipLaunchParm lp, int* Out) +{ + int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + Out[tid] = global[tid]; +} + +int main() +{ + int *A, *B, *Ad; + A = new int[NUM]; + B = new int[NUM]; + for(unsigned i=0;i