From 1aba3c4375ffe483eee98a0efe8236d65be6ff93 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 14 Mar 2017 22:11:34 +0530 Subject: [PATCH 1/4] Added hipMemsetD8 Change-Id: I6a230a036c9c46c72a77d5f93c16ce8a00c3f837 --- include/hip/hcc_detail/hip_runtime_api.h | 21 +++++--- src/hip_memory.cpp | 50 ++++++++++++++++++ tests/src/context/hipMemsetD8.cpp | 67 ++++++++++++++++++++++++ 3 files changed, 130 insertions(+), 8 deletions(-) create mode 100644 tests/src/context/hipMemsetD8.cpp diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 7f85aad28d..0d3ecc6613 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -1208,19 +1208,24 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp #endif /** - * @brief Copy data from src to dst asynchronously. + * @brief Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant byte value value. * - * It supports memory from host to device, - * device to host, device to device and host to host. - * - * @param[out] dst Data being copy to - * @param[in] src Data being copy from + * @param[out] dst Data being filled + * @param[in] constant value to be set * @param[in] sizeBytes Data size in bytes - * @param[in] accelerator_view Accelerator view which the copy is being enqueued - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized */ hipError_t hipMemset(void* dst, int value, size_t sizeBytes ); +/** + * @brief Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant byte value value. + * + * @param[out] dst Data ptr to be filled + * @param[in] constant value to be set + * @param[in] sizeBytes Data size in bytes + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized + */ +hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes ); /** * @brief Fills the first sizeBytes bytes of the memory area pointed to by dev with the constant byte value value. diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index c6b9406778..a92d11b847 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -1004,6 +1004,56 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes ) return ihipLogStatus(e); } +hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes ) +{ + HIP_INIT_CMD_API(dst, value, sizeBytes); + + hipError_t e = hipSuccess; + + hipStream_t stream = hipStreamNull; + // TODO - call an ihip memset so HIP_TRACE is correct. + stream = ihipSyncAndResolveStream(stream); + + if (stream) { + auto crit = stream->lockopen_preKernelCommand(); + + stream->ensureHaveQueue(crit); + hc::completion_future cf ; + + if ((sizeBytes & 0x3) == 0) { + // use a faster dword-per-workitem copy: + try { + uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value) ; + ihipMemsetKernel (stream, crit, static_cast (dst), value32, sizeBytes/sizeof(uint32_t), &cf); + } + catch (std::exception &ex) { + e = hipErrorInvalidValue; + } + } else { + // use a slow byte-per-workitem copy: + try { + ihipMemsetKernel (stream, crit, static_cast (dst), value, sizeBytes, &cf); + } + catch (std::exception &ex) { + e = hipErrorInvalidValue; + } + } + cf.wait(); + + stream->lockclose_postKernelCommand("hipMemsetD8", &crit->_av); + + + if (HIP_LAUNCH_BLOCKING) { + tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset in %s.\n", __func__, ToString(stream).c_str()); + cf.wait(); + tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed in %s.\n", __func__, ToString(stream).c_str()); + } + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} hipError_t hipMemGetInfo (size_t *free, size_t *total) { diff --git a/tests/src/context/hipMemsetD8.cpp b/tests/src/context/hipMemsetD8.cpp new file mode 100644 index 0000000000..1cd43696aa --- /dev/null +++ b/tests/src/context/hipMemsetD8.cpp @@ -0,0 +1,67 @@ +/* +Copyright (c) 2015-2016 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. +*/ +// Simple test for hipMemsetD8. +// Also serves as a template for other tests. + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * //Small copy + * RUN: %t -N 10 --memsetval 0x42 + * // Oddball size + * RUN: %t -N 10013 --memsetval 0x5a + * // Big copy + * RUN: %t -N 256M --memsetval 0xa6 + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "test_common.h" + +int main(int argc, char *argv[]) +{ + HipTest::parseStandardArguments(argc, argv, true); + size_t Nbytes = N*sizeof(char); + char *A_h; + hipDeviceptr_t A_d; + A_h = new char[Nbytes]; + + HIPCHECK ( hipMalloc(&A_d, Nbytes) ); + A_h = (char*)malloc(Nbytes); + + printf ("Size=%zu memsetval=%2x \n", Nbytes, memsetval); + HIPCHECK ( hipMemsetD8(A_d, memsetval, Nbytes) ); + + HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + + for (int i=0; i Date: Tue, 14 Mar 2017 23:49:21 +0530 Subject: [PATCH 2/4] hipMemsetD8 support for HIP/NVCC path Change-Id: I48eee8266afd7b45a12d5ce2c4849b687a006c0f --- include/hip/nvcc_detail/hip_runtime_api.h | 5 +++++ tests/src/context/hipMemsetD8.cpp | 6 +++--- 2 files changed, 8 insertions(+), 3 deletions(-) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index dc51290167..8c3e0da639 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -397,6 +397,11 @@ inline static hipError_t hipMemsetAsync(void* devPtr,int value, size_t count, hi return hipCUDAErrorTohipError(cudaMemsetAsync(devPtr, value, count, stream)); } +inline static hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes ) +{ + return hipCUResultTohipError(cuMemsetD8(dest, value, sizeBytes)); +} + inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int device) { cudaDeviceProp cdprop; diff --git a/tests/src/context/hipMemsetD8.cpp b/tests/src/context/hipMemsetD8.cpp index 1cd43696aa..3730fcb70b 100644 --- a/tests/src/context/hipMemsetD8.cpp +++ b/tests/src/context/hipMemsetD8.cpp @@ -45,13 +45,13 @@ int main(int argc, char *argv[]) hipDeviceptr_t A_d; A_h = new char[Nbytes]; - HIPCHECK ( hipMalloc(&A_d, Nbytes) ); + HIPCHECK ( hipMalloc((void **) &A_d, Nbytes) ); A_h = (char*)malloc(Nbytes); printf ("Size=%zu memsetval=%2x \n", Nbytes, memsetval); HIPCHECK ( hipMemsetD8(A_d, memsetval, Nbytes) ); - HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + HIPCHECK ( hipMemcpy(A_h, (void *) A_d, Nbytes, hipMemcpyDeviceToHost)); for (int i=0; i Date: Wed, 15 Mar 2017 12:03:05 +0530 Subject: [PATCH 3/4] hipcc: Fix warning when HCC_AMDGPU_TARGET is not defined Change-Id: I5cc6b0e9fb23ec78152d8bcfe9e7511e2fe91055 --- bin/hipcc | 41 ++++++++++++++++++++++------------------- 1 file changed, 22 insertions(+), 19 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index d2822fd0da..381c774c94 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -326,27 +326,30 @@ foreach $arg (@ARGV) } $toolArgs .= " $arg" unless $swallowArg; } -foreach my $target (split(/,/, $ENV{HCC_AMDGPU_TARGET})) +if(defined $ENV{HCC_AMDGPU_TARGET}) { - if($target eq 'gfx701') + foreach my $target (split(/,/, $ENV{HCC_AMDGPU_TARGET})) { - $target_gfx701 = 1; - } - if($target eq 'gfx801') - { - $target_gfx801 = 1; - } - if($target eq 'gfx802') - { - $target_gfx802 = 1; - } - if($target eq 'gfx803') - { - $target_gfx803 = 1; - } - if($target eq 'gfx900') - { - $target_gfx900 = 1; + if($target eq 'gfx701') + { + $target_gfx701 = 1; + } + if($target eq 'gfx801') + { + $target_gfx801 = 1; + } + if($target eq 'gfx802') + { + $target_gfx802 = 1; + } + if($target eq 'gfx803') + { + $target_gfx803 = 1; + } + if($target eq 'gfx900') + { + $target_gfx900 = 1; + } } } if ($target_gfx701 eq 0 and $target_gfx801 eq 0 and $target_gfx802 eq 0 and $target_gfx803 eq 0 and $target_gfx900 eq 0) From 65bb22eefc620392cbcdf8f03beb61690acc71d3 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 15 Mar 2017 12:03:44 +0530 Subject: [PATCH 4/4] Disable broken tests on hcc path Change-Id: Id6234da576566faa32d5fdf42dca6d6267596823 --- tests/src/deviceLib/hipTestDeviceSymbol.cpp | 2 +- tests/src/kernel/hipDynamicShared.cpp | 2 +- tests/src/kernel/hipDynamicShared2.cpp | 2 +- tests/src/runtimeApi/memory/hipArray.cpp | 4 ++-- 4 files changed, 5 insertions(+), 5 deletions(-) diff --git a/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/tests/src/deviceLib/hipTestDeviceSymbol.cpp index e58aa58877..9e188e9f17 100644 --- a/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s EXCLUDE_HIP_PLATFORM all * RUN: %t * HIT_END */ diff --git a/tests/src/kernel/hipDynamicShared.cpp b/tests/src/kernel/hipDynamicShared.cpp index 522572a191..ba19fcaa0d 100644 --- a/tests/src/kernel/hipDynamicShared.cpp +++ b/tests/src/kernel/hipDynamicShared.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM hcc * RUN: %t EXCLUDE_HIP_PLATFORM nvcc * HIT_END */ diff --git a/tests/src/kernel/hipDynamicShared2.cpp b/tests/src/kernel/hipDynamicShared2.cpp index 0f6ebb4927..95e70a9956 100644 --- a/tests/src/kernel/hipDynamicShared2.cpp +++ b/tests/src/kernel/hipDynamicShared2.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM hcc * RUN: %t EXCLUDE_HIP_PLATFORM nvcc * HIT_END */ diff --git a/tests/src/runtimeApi/memory/hipArray.cpp b/tests/src/runtimeApi/memory/hipArray.cpp index 8f831bf5e0..b31973e3d2 100644 --- a/tests/src/runtimeApi/memory/hipArray.cpp +++ b/tests/src/runtimeApi/memory/hipArray.cpp @@ -21,8 +21,8 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc - * RUN: %t EXCLUDE_HIP_PLATFORM + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all + * RUN: %t * HIT_END */