From fcb0a3d4e2a1863b6d1deeb1b30a0175a8152017 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary <51944368+cjatin@users.noreply.github.com> Date: Mon, 8 Jul 2019 16:30:38 +0530 Subject: [PATCH 1/5] Adding bounds check before hipMemset (#1190) * Adding bounds check in ihipMemset * Adding ihipMemPtrGetInfo to hipMemPtrGetInfo --- hipamd/src/hip_memory.cpp | 51 +++++++++++++++++++++++---------------- 1 file changed, 30 insertions(+), 21 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index b85ba61584..720d04deb8 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/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)); } From d0afaff28003b3b26f84feaf3e6c8b3d30ac0754 Mon Sep 17 00:00:00 2001 From: darkbuck Date: Mon, 8 Jul 2019 07:00:51 -0400 Subject: [PATCH 2/5] Fix even more typo. (#1204) --- hipamd/include/hip/hcc_detail/hip_vector_types.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index 2ce216e409..07908c71a4 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/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; } }; From bb72dd047ca32a038830cb8b174be8db857db7d1 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 10 Jul 2019 03:04:52 +0000 Subject: [PATCH 3/5] [ci] Enable tests on ROCm 2.6 (#1208) --- hipamd/Jenkinsfile | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/hipamd/Jenkinsfile b/hipamd/Jenkinsfile index 744f163d61..e583ff5019 100644 --- a/hipamd/Jenkinsfile +++ b/hipamd/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 From 034e5c0311e192fc6a7e10f5a9df13489be973a2 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Tue, 9 Jul 2019 23:05:04 -0400 Subject: [PATCH 4/5] Fix a compilation error when templated operator<< fails if 'U' type is std::ostream. (#1209) --- hipamd/include/hip/hcc_detail/hip_vector_types.h | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index 07908c71a4..953e4eb5f8 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -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<<( From b1d8736f4a46ec0d3c96f6eebae82d9b40c6aed1 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Tue, 9 Jul 2019 23:05:19 -0400 Subject: [PATCH 5/5] [hip][samples] fix a typo in module_api/defaultDriver sample (#1210) --- hipamd/samples/0_Intro/module_api/defaultDriver.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/samples/0_Intro/module_api/defaultDriver.cpp b/hipamd/samples/0_Intro/module_api/defaultDriver.cpp index 9fa0b0ee09..ea36aabcf4 100644 --- a/hipamd/samples/0_Intro/module_api/defaultDriver.cpp +++ b/hipamd/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; } }