From 3cfb9c0d4061ba513d9a72e57754be06c6853532 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 17 Apr 2018 18:27:27 +0530 Subject: [PATCH] Added hipMemset2DAsync support --- include/hip/hcc_detail/hip_runtime_api.h | 21 ++++ include/hip/nvcc_detail/hip_runtime_api.h | 8 ++ src/hip_memory.cpp | 20 +++- tests/src/runtimeApi/memory/hipMemset2D.cpp | 119 ++++++++++++++++++++ 4 files changed, 167 insertions(+), 1 deletion(-) create mode 100644 tests/src/runtimeApi/memory/hipMemset2D.cpp diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 78d9e5386c..8674824c48 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -50,6 +50,13 @@ THE SOFTWARE. #define HIP_LAUNCH_PARAM_BUFFER_SIZE ((void*)0x02) #define HIP_LAUNCH_PARAM_END ((void*)0x03) +#ifdef __cplusplus + #define __dparm(x) \ + = x +#else + #define __dparm(x) +#endif + // Structure definitions: #ifdef __cplusplus extern "C" { @@ -1436,6 +1443,20 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height); +/** + * @brief Fills asynchronously the memory area pointed to by dst with the constant value. + * + * @param[in] dst Pointer to device memory + * @param[in] pitch - data size in bytes + * @param[in] value - constant value to be set + * @param[in] width + * @param[in] height + * @param[in] stream + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree + */ + +hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height,hipStream_t stream __dparm(0)); + /** * @brief Query memory info. * Return snapshot of free memory, and total allocatable memory on the device. diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 1ab3b258db..d391c31db8 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -625,6 +625,14 @@ inline static hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, s return hipCUResultTohipError(cuMemsetD8(dest, value, sizeBytes)); } +inline static hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) { + return hipCUDAErrorTohipError(cudaMemset2D(dst, pitch, value, width, height)); +} + +inline static hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream __dparm(0)) { + return hipCUDAErrorTohipError(cudaMemset2DAsync(dst, pitch, value, width, height, stream)); +} + inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int device) { struct cudaDeviceProp cdprop; cudaError_t cerror; diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index e6d6ebd09e..3073921f50 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -1552,7 +1552,7 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea { hipError_t e = hipSuccess; - if (stream) { + if (stream && (dst != NULL)) { if(copyDataType == ihipMemsetDataTypeChar){ if ((sizeBytes & 0x3) == 0) { // use a faster dword-per-workitem copy: @@ -1647,6 +1647,24 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t return ihipLogStatus(e); } +hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream ) +{ + HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, pitch, value, width, height, stream); + + hipError_t e = hipSuccess; + + stream = ihipSyncAndResolveStream(stream); + + if (stream) { + size_t sizeBytes = pitch * height; + e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar); + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +}; + hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes); diff --git a/tests/src/runtimeApi/memory/hipMemset2D.cpp b/tests/src/runtimeApi/memory/hipMemset2D.cpp new file mode 100644 index 0000000000..51477e9eb2 --- /dev/null +++ b/tests/src/runtimeApi/memory/hipMemset2D.cpp @@ -0,0 +1,119 @@ +/* +Copyright (c) 2015-present 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 memset. +// Also serves as a template for other tests. + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * RUN: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "test_common.h" + +bool testhipMemset2D(int memsetval,int p_gpuDevice) +{ + size_t numH = 256; + size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH; + size_t elements = numW* numH; + + + printf ("testhipMemset2D memsetval=%2x device=%d\n", memsetval, p_gpuDevice); + char *A_d; + char *A_h; + bool testResult = true; + + HIPCHECK ( hipMallocPitch((void**)&A_d, &pitch_A, width , numH) ); + A_h = (char*)malloc(sizeElements); + HIPASSERT(A_h != NULL); + for (size_t i=0; i