diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index 744f163d61..e583ff5019 100644 --- a/projects/hip/Jenkinsfile +++ b/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 diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h index 2ce216e409..953e4eb5f8 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/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::value>::type, typename std::enable_if{}>* = nullptr> inline __host__ __device__ HIP_vector_type operator<<( diff --git a/projects/hip/samples/0_Intro/module_api/defaultDriver.cpp b/projects/hip/samples/0_Intro/module_api/defaultDriver.cpp index 9fa0b0ee09..ea36aabcf4 100644 --- a/projects/hip/samples/0_Intro/module_api/defaultDriver.cpp +++ b/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; } } diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index b85ba61584..720d04deb8 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/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 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)); }