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