From c7f8ffe41eaff8143cf1de87ea1099804c3c65fa 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 [ROCm/hip commit: 5ed16432f8d43d99c423d0996ac5479952b5b4ea] --- projects/hip/src/hip_memory.cpp | 51 +++++++++++++++++++-------------- 1 file changed, 30 insertions(+), 21 deletions(-) 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)); } From 8368ca472ddfd2c887dfb88b4af1c70b3c7b1e2e 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) [ROCm/hip commit: c8b76a6a63c4519e90dd833a7321a79b77d2c746] --- projects/hip/include/hip/hcc_detail/hip_vector_types.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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..07908c71a4 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; } }; From 6b29461f65401c4faab49e29cba2e474ece84d17 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) [ROCm/hip commit: b5150c8bc7960a715a429e6b91b9130cec576f84] --- projects/hip/Jenkinsfile | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) 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 From d46430b450d55a9e4ac8e52aedca9495634a6a08 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) [ROCm/hip commit: e87b5a1e8aa9c850310721859f0f8fc95f0ebd3a] --- projects/hip/include/hip/hcc_detail/hip_vector_types.h | 1 + 1 file changed, 1 insertion(+) 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 07908c71a4..953e4eb5f8 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/hip/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 706d95829782cde965250bd545c1ed32eb1f30ca 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) [ROCm/hip commit: f21601733e2bb07c92b91b7981638cae80403a24] --- projects/hip/samples/0_Intro/module_api/defaultDriver.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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; } }