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

@@ -29,7 +29,7 @@ THE SOFTWARE.
 #error("This version of HIP requires a newer version of HCC.");
 #endif
 
-// #define USE_MEMCPYTOSYMBOL
+#define USE_MEMCPYTOSYMBOL
 //
 
 

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 <typename T>
 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
 }
-

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<hip/hip_runtime.h>
+#include<hip/hip_runtime_api.h>
+#include<iostream>
+
+#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<NUM;i++)
+    {
+        A[i] = -1*i;
+        B[i] = 0;
+    }
+
+    hipMalloc((void**)&Ad, SIZE);
+    hipMemcpyToSymbol("global", A, SIZE, 0, hipMemcpyHostToDevice);
+    hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad);
+    hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
+    for(unsigned i=0;i<NUM;i++) {
+        assert(A[i] == B[i]);
+    }
+}