Merge 'master' into 'amd-master'

Change-Id: Ic61bd45a4904ddb1645e90b9e8fb42ffeab4825a

packaging/hip_base.txt

@@ -28,7 +28,7 @@ set(CPACK_BINARY_DEB "ON")
 set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm")
 set(CPACK_DEBIAN_PACKAGE_DEPENDS "perl (>= 5.0)")
 set(CPACK_BINARY_RPM "ON")
-set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64")
+set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}")
 set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst")
 set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm")
 set(CPACK_RPM_PACKAGE_AUTOREQPROV " no")

packaging/hip_doc.txt

@@ -33,7 +33,7 @@ set(CPACK_GENERATOR "TGZ;DEB;RPM")
 set(CPACK_BINARY_DEB "ON")
 set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION})")
 set(CPACK_BINARY_RPM "ON")
-set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64")
+set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}")
 set(CPACK_RPM_PACKAGE_AUTOREQPROV " no")
 set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}")
 set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt")

packaging/hip_hcc.txt

@@ -37,7 +37,7 @@ else()
     set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@)")
 endif()
 set(CPACK_BINARY_RPM "ON")
-set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64")
+set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}")
 set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst")
 set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm")
 set(CPACK_RPM_PACKAGE_AUTOREQPROV " no")

packaging/hip_nvcc.txt

@@ -20,7 +20,7 @@ set(CPACK_BINARY_DEB "ON")
 #set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm")
 set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), cuda (>= 7.5)")
 set(CPACK_BINARY_RPM "ON")
-set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64")
+set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}")
 #set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst")
 #set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm")
 set(CPACK_RPM_PACKAGE_AUTOREQPROV " no")

packaging/hip_samples.txt

@@ -21,7 +21,7 @@ set(CPACK_GENERATOR "TGZ;DEB;RPM")
 set(CPACK_BINARY_DEB "ON")
 set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION})")
 set(CPACK_BINARY_RPM "ON")
-set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64")
+set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}")
 set(CPACK_RPM_PACKAGE_AUTOREQPROV " no")
 set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}")
 set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt")

src/hip_device.cpp

@@ -60,12 +60,12 @@ hipError_t ihipGetDeviceCount(int *count)
         *count = g_deviceCnt;
 
         if (*count > 0) {
-            e = ihipLogStatus(hipSuccess);
+            e = hipSuccess;
         } else {
-            e = ihipLogStatus(hipErrorNoDevice);
+            e = hipErrorNoDevice;
         }
     } else {
-        e = ihipLogStatus(hipErrorInvalidValue);
+        e = hipErrorInvalidValue;
     }
     return e;
 }
@@ -73,7 +73,7 @@ hipError_t ihipGetDeviceCount(int *count)
 hipError_t hipGetDeviceCount(int *count)
 {
     HIP_INIT_API(count);
-    return ihipGetDeviceCount(count);
+    return ihipLogStatus(ihipGetDeviceCount(count));
 }
 
 hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig)
@@ -205,7 +205,7 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device
     hipError_t e = hipSuccess;
 
     if(pi == nullptr) {
-        return ihipLogStatus(hipErrorInvalidValue);
+        return hipErrorInvalidValue;
     }
 
     auto * hipDevice = ihipGetDevice(device);

src/hip_hcc.cpp

@@ -156,7 +156,7 @@ thread_local TidInfo tls_tidInfo;
 //=================================================================================================
 // Top-level "free" functions:
 //=================================================================================================
-void recordApiTrace(std::string *fullStr, const std::string &apiStr)
+uint64_t recordApiTrace(std::string *fullStr, const std::string &apiStr)
 {
     auto apiSeqNum = tls_tidInfo.apiSeqNum();
     auto tid = tls_tidInfo.tid();
@@ -178,10 +178,14 @@ void recordApiTrace(std::string *fullStr, const std::string &apiStr)
     *fullStr += " ";
     *fullStr += apiStr;
 
+    uint64_t apiStartTick = getTicks();
+
 
     if (COMPILE_HIP_DB && HIP_TRACE_API) {
-        fprintf (stderr, "%s<<hip-api tid:%s%s\n" , API_COLOR, fullStr->c_str(), API_COLOR_END);
+        fprintf (stderr, "%s<<hip-api tid:%s @%lu%s\n" , API_COLOR, fullStr->c_str(), apiStartTick, API_COLOR_END);
     }
+
+    return apiStartTick;
 }
 
 

src/hip_hcc_internal.h

@@ -244,23 +244,28 @@ static const DbName dbName [] =
 #endif
 
 
-
+static inline uint64_t getTicks()
+{
+    return hc::get_system_ticks();
+}
 
 //---
-extern void recordApiTrace(std::string *fullStr, const std::string &apiStr);
+extern uint64_t recordApiTrace(std::string *fullStr, const std::string &apiStr);
 
 #if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1)
 #define API_TRACE(forceTrace, ...)\
+uint64_t hipApiStartTick;\
 {\
     tls_tidInfo.incApiSeqNum();\
     if (forceTrace || (HIP_PROFILE_API || (COMPILE_HIP_DB && (HIP_TRACE_API & (1<<TRACE_ALL))))) {\
         std::string apiStr = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')';\
         std::string fullStr;\
-        recordApiTrace(&fullStr, apiStr);\
+        hipApiStartTick = recordApiTrace(&fullStr, apiStr);\
         if      (HIP_PROFILE_API == 0x1) {MARKER_BEGIN(__func__, "HIP") }\
         else if (HIP_PROFILE_API == 0x2) {MARKER_BEGIN(fullStr.c_str(), "HIP"); }\
     }\
 }
+
 #else
 // Swallow API_TRACE
 #define API_TRACE(IS_CMD, ...)\
@@ -302,7 +307,10 @@ extern void recordApiTrace(std::string *fullStr, const std::string &apiStr);
         tls_lastHipError = localHipStatus;\
         \
         if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API & (1<<TRACE_ALL)) {\
-            fprintf(stderr, "  %ship-api tid:%d.%lu %-30s ret=%2d (%s)>>%s\n", (localHipStatus == 0) ? API_COLOR:KRED, tls_tidInfo.tid(),tls_tidInfo.apiSeqNum(),  __func__, localHipStatus, ihipErrorString(localHipStatus), API_COLOR_END);\
+            auto ticks = getTicks() - hipApiStartTick;\
+            fprintf(stderr, "  %ship-api tid:%d.%lu %-30s ret=%2d (%s)>> +%lu ns%s\n", \
+                    (localHipStatus == 0) ? API_COLOR:KRED, tls_tidInfo.tid(),tls_tidInfo.apiSeqNum(),  \
+                    __func__, localHipStatus, ihipErrorString(localHipStatus), ticks, API_COLOR_END);\
         }\
         if (HIP_PROFILE_API) { MARKER_END(); }\
         localHipStatus;\

src/hip_module.cpp

@@ -374,6 +374,8 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname)
 
 hipError_t hipModuleUnload(hipModule_t hmod)
 {
+    HIP_INIT_API(hmod);
+
     // TODO - improve this synchronization so it is thread-safe.
     // Currently we want for all inflight activity to complete, but don't prevent another
     // thread from launching new kernels before we finish this operation.
@@ -408,7 +410,7 @@ hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char
     hipError_t ret = hipSuccess;
 
     if (name == nullptr){
-        return ihipLogStatus(hipErrorInvalidValue);
+        return (hipErrorInvalidValue);
     }
 
     if (ctx == nullptr){
@@ -431,7 +433,7 @@ hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char
         hsa_executable_symbol_t 	symbol;
         status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &symbol);
         if(status != HSA_STATUS_SUCCESS){
-            return ihipLogStatus(hipErrorNotFound);
+            return hipErrorNotFound;
         }
 
         status = hsa_executable_symbol_get_info(symbol,
@@ -503,11 +505,11 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f,
             if(config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END){
                 kernArgSize = *(size_t*)(config[3]);
             } else {
-                return ihipLogStatus(hipErrorNotInitialized);
+                return hipErrorNotInitialized;
             }
 
         }else{
-            return ihipLogStatus(hipErrorInvalidValue);
+            return hipErrorInvalidValue;
         }
 
 
@@ -624,7 +626,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes,
     }
     else{
         hipFunction_t func;
-        ihipModuleGetSymbol(&func, hmod, name);
+        ret = ihipModuleGetSymbol(&func, hmod, name);
         *bytes = PrintSymbolSizes(hmod->ptr, name) + sizeof(amd_kernel_code_t);
         *dptr = reinterpret_cast<void*>(func->_object);
         return ihipLogStatus(ret);

src/hip_stream.cpp

@@ -140,7 +140,6 @@ hipError_t hipStreamQuery(hipStream_t stream)
 //---
 hipError_t hipStreamSynchronize(hipStream_t stream)
 {
-    HIP_INIT_API(stream);
     HIP_INIT_SPECIAL_API(TRACE_SYNC, stream);
 
     hipError_t e = hipSuccess;

tests/src/runtimeApi/memory/p2p_copy_coherency.cpp

@@ -33,11 +33,13 @@ THE SOFTWARE.
 
 #ifdef __HIP_PLATFORM_HCC__
 #include <hc_am.hpp>
-#define USE_HCC_MEMTRACKER 0
 #endif
 
+#define USE_HCC_MEMTRACKER 0  /* Debug flag to show the memtracker periodically */
 
-int elementSizes[] = {16, 1024,524288};
+
+
+int elementSizes[] = {1, 16, 1024, 524288, 16*1000*1000};
 int nSizes  = sizeof(elementSizes) / sizeof(int);
 
 int enablePeers(int dev0, int dev1)
@@ -57,8 +59,9 @@ int enablePeers(int dev0, int dev1)
     return 0;
 };
 
+// Set value of array to specified 32-bit integer:
 __global__ void
-memsetIntKernel(/*hipLaunchParm lp,*/ int * ptr, const int val, size_t numElements)
+memsetIntKernel(int * ptr, const int val, size_t numElements)
 {
     int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
     int stride = hipBlockDim_x * hipGridDim_x ;
@@ -68,7 +71,7 @@ memsetIntKernel(/*hipLaunchParm lp,*/ int * ptr, const int val, size_t numElemen
 };
 
 __global__ void
-memcpyIntKernel(/*hipLaunchParm lp, */const int * src, int* dst, size_t numElements)
+memcpyIntKernel(const int * src, int* dst, size_t numElements)
 {
     int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
     int stride = hipBlockDim_x * hipGridDim_x ;
@@ -77,6 +80,9 @@ memcpyIntKernel(/*hipLaunchParm lp, */const int * src, int* dst, size_t numEleme
     }
 };
 
+
+// CHeck arrays in reverse order, to more easily detect cases where 
+// the copy is "partially" done.
 void checkReverse(const int *ptr, int numElements, int expected) {
     for (int i=numElements-1; i>=0; i--) {
         if (ptr[i] != expected) {
@@ -88,7 +94,8 @@ void checkReverse(const int *ptr, int numElements, int expected) {
     printf ("test:   OK\n");
 }
 
-void runTest(bool stepAIsCopy, bool hostSync, hipStream_t gpu0Stream, hipStream_t gpu1Stream, int numElements,
+
+void runTestImpl(bool stepAIsCopy, bool hostSync, hipStream_t gpu0Stream, hipStream_t gpu1Stream, int numElements,
              int * dataGpu0_0, int * dataGpu0_1, int *dataGpu1, int *dataHost, int expected)
 {
     hipEvent_t e;
@@ -96,7 +103,7 @@ void runTest(bool stepAIsCopy, bool hostSync, hipStream_t gpu0Stream, hipStream_
         HIPCHECK(hipEventCreateWithFlags(&e,0));
     }
     const size_t sizeElements = numElements * sizeof(int);
-    printf ("test: runTest with %zu bytes %s with hostSync %s\n", sizeElements, stepAIsCopy ? "copy" : "kernel", hostSync ? "enabled" : "disabled");
+    printf ("test: runTestImpl with %zu bytes %s with hostSync %s\n", sizeElements, stepAIsCopy ? "copy" : "kernel", hostSync ? "enabled" : "disabled");
 
     hipStream_t stepAStream = gpu0Stream;
 
@@ -127,9 +134,12 @@ void runTest(bool stepAIsCopy, bool hostSync, hipStream_t gpu0Stream, hipStream_
     HIPCHECK(hipStreamSynchronize(gpu0Stream));
 
     checkReverse(dataHost, numElements, expected);
+    if(!hostSync) {
+        HIPCHECK(hipEventDestroy(e));
+    }
 }
 
-void testMultiGpu(int dev0, int dev1, int numElements, bool hostSync, bool useMemcpy)
+void testMultiGpu(int dev0, int dev1, int numElements, bool hostSync)
 {
     const size_t sizeElements = numElements * sizeof(int);
 
@@ -163,12 +173,15 @@ void testMultiGpu(int dev0, int dev1, int numElements, bool hostSync, bool useMe
 #endif
 
     printf ("  test: init complete\n");
-    runTest(useMemcpy , hostSync, gpu0Stream, gpu1Stream, numElements, dataGpu0_0,dataGpu0_1, dataGpu1, dataHost, expected);
+    runTestImpl(true, hostSync, gpu0Stream, gpu1Stream, numElements, dataGpu0_0,dataGpu0_1, dataGpu1, dataHost, expected);
 
     HIPCHECK(hipFree(dataGpu0_0));
     HIPCHECK(hipFree(dataGpu0_1));
     HIPCHECK(hipFree(dataGpu1));
     HIPCHECK(hipHostFree(dataHost));
+
+    HIPCHECK(hipStreamDestroy(gpu0Stream));
+    HIPCHECK(hipStreamDestroy(gpu1Stream));
 };
 
 int main(int argc, char *argv[])
@@ -192,11 +205,9 @@ int main(int argc, char *argv[])
         return -1;
     };
 
-    for(int index = 1;index < nSizes;index++) {
-        testMultiGpu(dev0, dev1, elementSizes[index] , false /* GPU Synchronization*/, true);
-        testMultiGpu(dev0, dev1, elementSizes[index] , true /*Host Synchronization*/, true);
-        testMultiGpu(dev0, dev1, elementSizes[index] , true /*Host Synchronization*/, false);
-        testMultiGpu(dev0, dev1, elementSizes[index] , false /*Host Synchronization*/, false);
+    for(int index = 0;index < nSizes;index++) {
+        testMultiGpu(dev0, dev1, elementSizes[index] , false /*GPU Synchronization*/);
+        testMultiGpu(dev0, dev1, elementSizes[index] , true  /*Host Synchronization*/);
     }
 
 

tests/src/runtimeApi/synchronization/copy_coherency.cpp

@@ -0,0 +1,391 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+//ROCM_TARGET=gfx900 hipcc --genco memcpyInt.device.cpp -o memcpyInt.hsaco
+//hipcc copy_coherency.cpp  -I ~/X/HIP/tests/src/ ~/X/HIP/tests/src/test_common.cpp
+
+
+// TODO - add code object support here.
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
+ * RUN: %t 
+ * HIT_END
+ */
+
+
+// Test cache management (fences) and synchronization between kernel and copy commands.
+// Exhaustively tests 3 command types (copy, kernel, module kernel), 
+// many sync types (see SyncType), followed by another command, across a sweep 
+// of data sizes designed to stress various levels of the memory hierarchy.
+
+#include "hip/hip_runtime.h"
+#include "test_common.h"
+
+// TODO - turn this back on when test infra can copy the module files to use as test inputs.
+#define SKIP_MODULE_KERNEL 1
+
+
+class MemcpyFunction
+{
+public:
+    MemcpyFunction(const char *fileName, const char *functionName) { load(fileName, functionName); };
+    void load(const char *fileName, const char *functionName);
+    void launch(int * dst, const int * src, size_t numElements, hipStream_t s);
+
+private:
+    hipFunction_t _function;
+    hipModule_t   _module;
+};
+
+
+void MemcpyFunction::load(const char *fileName, const char *functionName)
+{
+#if SKIP_MODULE_KERNEL!=1
+    HIPCHECK(hipModuleLoad(&_module, fileName));
+    HIPCHECK(hipModuleGetFunction(&_function, _module, functionName));
+#endif
+};
+
+
+void MemcpyFunction::launch(int * dst, const int * src, size_t numElements, hipStream_t s)
+{
+    struct {
+        uint32_t _hidden[6];
+        int*        _dst;
+        const int*  _src;
+        size_t      _numElements;
+    } args;
+
+    args._dst = dst;
+    args._src = src;
+    args._numElements = numElements;
+
+    size_t size = sizeof(args);
+    void *config[] = {
+      HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
+      HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
+      HIP_LAUNCH_PARAM_END
+    };
+
+    unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
+    HIPCHECK(hipModuleLaunchKernel(_function, blocks, 1, 1, threadsPerBlock, 1, 1,   0/*dynamicShared*/, s, NULL, (void**)&config));
+};
+
+bool g_warnOnFail = true;
+//int g_elementSizes[] = {1, 16, 1024, 524288, 16*1000*1000}; // TODO
+int g_elementSizes[] = {128*1000, 256*1000, 16*1000*1000};
+
+MemcpyFunction g_moduleMemcpy("memcpyInt.hsaco", "memcpyIntKernel");
+
+
+
+// Set value of array to specified 32-bit integer:
+__global__ void
+memsetIntKernel(int * ptr, const int val, size_t numElements)
+{
+    int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+    int stride = hipBlockDim_x * hipGridDim_x ;
+    for (size_t i= gid; i< numElements; i+=stride){
+       ptr[i] = val;
+    }
+};
+
+__global__ void
+memcpyIntKernel(int *dst, const int * src, size_t numElements)
+{
+    int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+    int stride = hipBlockDim_x * hipGridDim_x ;
+    for (size_t i= gid; i< numElements; i+=stride){
+       dst[i] = src[i];
+    }
+};
+
+
+// CHeck arrays in reverse order, to more easily detect cases where 
+// the copy is "partially" done.
+void checkReverse(const int *ptr, int numElements, int expected) {
+    int mismatchCnt = 0;
+    for (int i=numElements-1; i>=0; i--) {
+        if (ptr[i] != expected) {
+            fprintf (stderr, "%s**error: i=%d, ptr[i] == (%x) , does not equal expected (%x)\n%s", KRED, i, ptr[i], expected, KNRM);
+            if (!g_warnOnFail) {
+                assert (ptr[i] == expected);
+            }
+            if (++mismatchCnt >= 10) {
+                break;
+            }
+        }
+    }
+
+    fprintf (stderr, "test:   OK\n");
+}
+
+#define ENUM_CASE_STR(x) case x: return #x
+
+enum CmdType {
+    COPY,
+    KERNEL,
+    MODULE_KERNEL,
+    MAX_CmdType
+};
+
+
+const char * CmdTypeStr(CmdType c)
+{
+    switch(c) {
+        ENUM_CASE_STR(COPY);
+        ENUM_CASE_STR(KERNEL);
+        ENUM_CASE_STR(MODULE_KERNEL);
+        default: return "UNKNOWN";
+    };
+} 
+
+
+enum SyncType {
+    NONE,
+    EVENT_QUERY,
+    EVENT_SYNC,
+    STREAM_WAIT_EVENT,
+    STREAM_QUERY,
+    STREAM_SYNC,
+    DEVICE_SYNC,
+    MAX_SyncType
+};
+
+
+const char * SyncTypeStr(SyncType s)
+{
+    switch(s) {
+        ENUM_CASE_STR(NONE);
+        ENUM_CASE_STR(EVENT_QUERY);
+        ENUM_CASE_STR(EVENT_SYNC);
+        ENUM_CASE_STR(STREAM_WAIT_EVENT);
+        ENUM_CASE_STR(STREAM_QUERY);
+        ENUM_CASE_STR(STREAM_SYNC);
+        ENUM_CASE_STR(DEVICE_SYNC);
+        default: return "UNKNOWN";
+    };
+};
+
+
+void runCmd(CmdType cmd, int *dst, const int *src, hipStream_t s, size_t numElements)
+{
+    switch (cmd) {
+        case COPY:
+            HIPCHECK(hipMemcpyAsync(dst, src, numElements*sizeof(int), hipMemcpyDeviceToDevice, s));
+            break;
+        case KERNEL:
+            {
+                unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
+                hipLaunchKernelGGL(memcpyIntKernel, dim3(blocks), dim3(threadsPerBlock), 0, s,
+                                   dst, src, numElements);
+            }
+            break;
+        case MODULE_KERNEL:
+            g_moduleMemcpy.launch(dst, src, numElements, s);
+            break;
+        default:
+            failed("unknown cmd=%d type", cmd);
+    };
+}
+
+void resetInputs( int * Ad, int * Bd, int *Cd, int *Ch, size_t numElements, int expected)
+{
+    unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
+    hipLaunchKernelGGL(memsetIntKernel, dim3(blocks), dim3(threadsPerBlock), 0, hipStream_t(0),
+                       Ad, expected, numElements);
+    hipLaunchKernelGGL(memsetIntKernel, dim3(blocks), dim3(threadsPerBlock), 0, hipStream_t(0),
+                       Bd, 0xDEADBEEF, numElements); // poison with bad value to ensure is overwritten correctly
+    hipLaunchKernelGGL(memsetIntKernel, dim3(blocks), dim3(threadsPerBlock), 0, hipStream_t(0),
+                       Bd, 0xF000BA55, numElements); // poison with bad value to ensure is overwritten correctly
+    memset(Ch, 13, numElements*sizeof(int)); // poison with bad value to ensure is overwritten correctly
+    HIPCHECK(hipDeviceSynchronize());
+}
+
+// Intended to test proper synchronization and cache flushing between CMDA and CMDB.
+// CMD are of type CmdType. All command copy memory, using either hipMemcpyAsync or kernel implementations.
+// CmdA copies from Ad to Bd, 
+// Some form of synchronization is applied.
+// Then cmdB copies from Bd to Cd.
+//
+// Cd is then copied to host Ch using a memory copy.
+//
+// Correct result at the end is that Ch contains the contents originally in Ad (integer 0x42)
+void runTestImpl(CmdType cmdAType, SyncType syncType, CmdType cmdBType, 
+            hipStream_t stream1, hipStream_t stream2, int numElements,
+             int * Ad, int * Bd, int *Cd, int *Ch, 
+             int expected)
+{
+    hipEvent_t e;
+    HIPCHECK(hipEventCreateWithFlags(&e,0));
+
+    resetInputs(Ad, Bd, Cd, Ch, numElements, expected);
+
+    const size_t sizeElements = numElements * sizeof(int);
+    fprintf (stderr, "test: runTest with %zu bytes (%6.2f MB) cmdA=%s; sync=%s; cmdB=%s\n", 
+            sizeElements, (double) (sizeElements/1024.0), CmdTypeStr(cmdAType), SyncTypeStr(syncType), CmdTypeStr(cmdBType));
+
+    if (SKIP_MODULE_KERNEL && ((cmdAType == MODULE_KERNEL) || (cmdBType == MODULE_KERNEL))) {
+        fprintf (stderr, "warn: skipping since test infra does not yet support modules\n");
+        return;
+    }
+            
+
+    // Step A:
+    runCmd(cmdAType, Bd, Ad, stream1, numElements);
+
+
+    // Sync in-between?
+    switch (syncType) {
+        case NONE:
+            break;
+        case EVENT_QUERY:
+            {
+                hipError_t st = hipErrorNotReady;
+                HIPCHECK(hipEventRecord(e, stream1));
+                do {
+                    st = hipEventQuery(e);
+                } while (st == hipErrorNotReady);
+                HIPCHECK(st);
+            }
+            break;
+        case EVENT_SYNC:
+            HIPCHECK(hipEventRecord(e, stream1));
+            HIPCHECK(hipEventSynchronize(e));
+            break;
+        case STREAM_WAIT_EVENT:
+            HIPCHECK(hipEventRecord(e, stream1));
+            HIPCHECK(hipStreamWaitEvent(stream2, e, 0));
+            break;
+        case STREAM_QUERY:
+            {
+                hipError_t st = hipErrorNotReady;
+                do {
+                    st = hipStreamQuery(stream1);
+                } while (st == hipErrorNotReady);
+                HIPCHECK(st);
+            }
+            break;
+        case STREAM_SYNC:
+            HIPCHECK(hipStreamSynchronize(stream1));
+            break;
+        case DEVICE_SYNC:
+            HIPCHECK(hipDeviceSynchronize());
+            break;
+        default:
+            fprintf(stderr, "warning: unknown sync type=%s", SyncTypeStr(syncType));
+            return; // FIXME, this doesn't clean up
+            //failed("unknown sync type=%s", SyncTypeStr(syncType));
+    };
+
+
+    runCmd(cmdBType, Cd, Bd, stream2, numElements);
+
+
+    // Copy back to host, use async copy to avoid any extra synchronization that might mask issues.
+    HIPCHECK(hipMemcpyAsync(Ch, Cd, sizeElements, hipMemcpyDeviceToHost, stream2));
+    HIPCHECK(hipStreamSynchronize(stream2));
+
+    checkReverse(Ch, numElements, expected);
+
+    HIPCHECK(hipEventDestroy(e));
+};
+
+
+void testWrapper(size_t numElements)
+{
+    const size_t sizeElements = numElements * sizeof(int);
+    const int expected = 0x42;
+    int * Ad, * Bd, *Cd, *Ch;
+
+    HIPCHECK(hipMalloc(&Ad, sizeElements));
+    HIPCHECK(hipMalloc(&Bd, sizeElements));
+    HIPCHECK(hipMalloc(&Cd, sizeElements));
+    HIPCHECK(hipHostMalloc(&Ch, sizeElements));  // Ch is the end array
+
+
+
+    hipStream_t stream1, stream2;
+
+    HIPCHECK(hipStreamCreate(&stream1));
+    HIPCHECK(hipStreamCreate(&stream2));
+
+
+    HIPCHECK(hipDeviceSynchronize());
+    fprintf (stderr, "test: init complete, start running tests\n");
+
+
+    runTestImpl(COPY, EVENT_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+
+    for (int cmdA=0; cmdA<MAX_CmdType; cmdA++) {
+        for (int cmdB=0; cmdB<MAX_CmdType; cmdB++) {
+            for (int syncMode=0; syncMode<MAX_SyncType; syncMode++) {
+                switch(syncMode) {
+                    //case NONE::
+                    case EVENT_QUERY:
+                    case EVENT_SYNC:
+                    case STREAM_WAIT_EVENT:
+                    //case STREAM_QUERY:
+                    case STREAM_SYNC:
+                    case DEVICE_SYNC:
+                        runTestImpl(CmdType(cmdA), SyncType(syncMode), CmdType(cmdB), stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+                        break;
+                    default:
+                        break;
+
+                }
+            }
+        }
+    }
+
+#if 0
+    runTestImpl(COPY, STREAM_SYNC, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+    runTestImpl(COPY, STREAM_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+    runTestImpl(COPY, STREAM_WAIT_EVENT, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+
+    runTestImpl(COPY, STREAM_WAIT_EVENT, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+#endif
+
+    HIPCHECK(hipFree(Ad));
+    HIPCHECK(hipFree(Bd));
+    HIPCHECK(hipFree(Cd));
+    HIPCHECK(hipHostFree(Ch));
+
+    HIPCHECK(hipStreamDestroy(stream1));
+    HIPCHECK(hipStreamDestroy(stream2));
+}
+
+
+int main(int argc, char *argv[])
+{
+
+    for(int index = 0;index < sizeof(g_elementSizes) / sizeof(int); index++) {
+        size_t numElements = g_elementSizes[index];
+        testWrapper(numElements);
+    }
+
+    passed();
+}
+
+
+// TODO
+// - test environment variables

tests/src/runtimeApi/synchronization/memcpyInt.device.cpp

@@ -0,0 +1,13 @@
+#include <hip/hip_runtime.h>
+
+
+    
+extern "C" __global__ void
+memcpyIntKernel(hipLaunchParm lp, int *dst, const int * src, size_t numElements)
+{
+    int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+    int stride = hipBlockDim_x * hipGridDim_x ;
+    for (size_t i= gid; i< numElements; i+=stride){
+       dst[i] = src[i];
+    }
+};

tests/src/test_common.h

@@ -23,7 +23,6 @@ THE SOFTWARE.
 #include <stddef.h>
 
 #include "hip/hip_runtime.h"
-#include "hip/hip_texture_types.h"
 #include "hip/hip_runtime_api.h"
 
 #define HC __attribute__((hc))