Merge 'master' into 'amd-master'

Change-Id: I9e1db1ea296de411edef79d9ffd0c050247feb02


[ROCm/hip commit: 682e711379]

projects/hip/Jenkinsfile

@@ -296,13 +296,13 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_
 String build_config = 'Release'
 String job_name = env.JOB_NAME.toLowerCase( )
 
-// The following launches 3 builds in parallel: rocm-head, rocm-2.4.x and cuda-10.x
-parallel rocm_2_4:
+// The following launches 3 builds in parallel: rocm-head, rocm-2.6.x and cuda-10.x
+parallel rocm_2_6:
 {
   node('hip-rocm')
   {
-    String hcc_ver = 'rocm-2.4.x'
-    String from_image = 'ci_test_nodes/rocm-2.4.x/ubuntu-16.04:latest'
+    String hcc_ver = 'rocm-2.6.x'
+    String from_image = 'ci_test_nodes/rocm-2.6.x/ubuntu-16.04:latest'
     String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video'
 
     // Checkout source code, dependencies and version files

projects/hip/include/hip/hcc_detail/hip_vector_types.h

@@ -218,7 +218,7 @@ THE SOFTWARE.
                 r[0] = d[0] == x.d[0];
                 r[1] = d[1] == x.d[1];
                 r[2] = d[2] == x.d[2];
-                return -r;
+                return r;
             }
         };
 
@@ -738,6 +738,7 @@ THE SOFTWARE.
         typename T,
         unsigned int n,
         typename U,
+        typename std::enable_if<std::is_arithmetic<U>::value>::type,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
     inline __host__ __device__
     HIP_vector_type<T, n> operator<<(

projects/hip/samples/0_Intro/module_api/defaultDriver.cpp

@@ -68,7 +68,7 @@ int main() {
     for (uint32_t i = 0; i < LEN; i++) {
         if (A[i] != B[i]) {
             mismatchCount++;
-            std::cout << "error: mismatch " << A[i] << " != " << C[i] << std::endl;
+            std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
         }
     }
 

projects/hip/src/hip_memory.cpp

@@ -1506,6 +1506,29 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height,
 }
 }  // namespace
 
+//Get the allocated size
+hipError_t ihipMemPtrGetInfo(void* ptr, size_t* size) {
+    hipError_t e = hipSuccess;
+    if (ptr != nullptr && size != nullptr) {
+        *size = 0;
+        hc::accelerator acc;
+#if (__hcc_workweek__ >= 17332)
+        hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0);
+#else
+        hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
+#endif
+        am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr);
+        if (status == AM_SUCCESS) {
+            *size = amPointerInfo._sizeBytes;
+        } else {
+            e = hipErrorInvalidValue;
+        }
+    } else {
+        e = hipErrorInvalidValue;
+    }
+    return e;
+}
+
 template <typename T>
 void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t count) {
     static constexpr uint32_t block_dim = 256;
@@ -1532,13 +1555,17 @@ typedef enum ihipMemsetDataType {
     ihipMemsetDataTypeInt    = 2
 }ihipMemsetDataType;
 
-hipError_t ihipMemset(void* dst, int  value, size_t count, hipStream_t stream, enum ihipMemsetDataType copyDataType  )
+hipError_t ihipMemset(void* dst, int  value, size_t count, hipStream_t stream, enum ihipMemsetDataType copyDataType)
 {
     hipError_t e = hipSuccess;
 
     if (count == 0) return e;
 
-    if (stream && (dst != NULL)) {
+    size_t allocSize = 0;
+    bool isInbound = (ihipMemPtrGetInfo(dst, &allocSize) == hipSuccess);
+    isInbound &= (allocSize >= count);
+
+    if (stream && (dst != NULL) && isInbound) {
         if(copyDataType == ihipMemsetDataTypeChar){
             if ((count & 0x3) == 0) {
                 // use a faster dword-per-workitem copy:
@@ -1898,25 +1925,7 @@ hipError_t hipMemGetInfo(size_t* free, size_t* total) {
 hipError_t hipMemPtrGetInfo(void* ptr, size_t* size) {
     HIP_INIT_API(hipMemPtrGetInfo, ptr, size);
 
-    hipError_t e = hipSuccess;
-
-    if (ptr != nullptr && size != nullptr) {
-        hc::accelerator acc;
-#if (__hcc_workweek__ >= 17332)
-        hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0);
-#else
-        hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
-#endif
-        am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr);
-        if (status == AM_SUCCESS) {
-            *size = amPointerInfo._sizeBytes;
-        } else {
-            e = hipErrorInvalidValue;
-        }
-    } else {
-        e = hipErrorInvalidValue;
-    }
-    return ihipLogStatus(e);
+    return ihipLogStatus(ihipMemPtrGetInfo(ptr, size));
 }