From da302c3e9304c4cc5033ebffeb4e753cbc834697 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 7 May 2018 10:24:30 +0530 Subject: [PATCH] Added hipMemset3D --- include/hip/hcc_detail/hip_runtime_api.h | 21 +++++ include/hip/nvcc_detail/hip_runtime_api.h | 8 ++ src/hip_memory.cpp | 36 ++++++++ tests/src/runtimeApi/memory/hipMemset3D.cpp | 98 +++++++++++++++++++++ 4 files changed, 163 insertions(+) create mode 100644 tests/src/runtimeApi/memory/hipMemset3D.cpp diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 8becab3c9c..ada0fac19a 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -1450,6 +1450,27 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height,hipStream_t stream __dparm(0)); +/** + * @brief Fills synchronously the memory area pointed to by pitchedDevPtr with the constant value. + * + * @param[in] pitchedDevPtr + * @param[in] value - constant value to be set + * @param[in] extent + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree + */ +hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ); + +/** + * @brief Fills asynchronously the memory area pointed to by pitchedDevPtr with the constant value. + * + * @param[in] pitchedDevPtr + * @param[in] value - constant value to be set + * @param[in] extent + * @param[in] stream + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree + */ +hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,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 6f222648a2..f8e4e39136 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -649,6 +649,14 @@ inline static hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, si return hipCUDAErrorTohipError(cudaMemset2DAsync(dst, pitch, value, width, height, stream)); } +inline static hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ){ + return hipCUDAErrorTohipError(cudaMemset3D(pitchedDevPtr, value, extent)); +} + +inline static hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream __dparm(0) ){ + return hipCUDAErrorTohipError(cudaMemset3DAsync(pitchedDevPtr, value, extent, 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 48a8da81ca..d9b0c7acab 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -1698,6 +1698,42 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes return ihipLogStatus(e); } +hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ) +{ + HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent); + hipError_t e = hipSuccess; + + hipStream_t stream = hipStreamNull; + // TODO - call an ihip memset so HIP_TRACE is correct. + stream = ihipSyncAndResolveStream(stream); + if (stream) { + size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth; + e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar); + stream->locked_wait(); + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} + +hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream ) +{ + HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent); + hipError_t e = hipSuccess; + + // TODO - call an ihip memset so HIP_TRACE is correct. + stream = ihipSyncAndResolveStream(stream); + if (stream) { + size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth; + e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar); + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} + hipError_t hipMemGetInfo(size_t* free, size_t* total) { HIP_INIT_API(free, total); diff --git a/tests/src/runtimeApi/memory/hipMemset3D.cpp b/tests/src/runtimeApi/memory/hipMemset3D.cpp new file mode 100644 index 0000000000..40f2f3e67f --- /dev/null +++ b/tests/src/runtimeApi/memory/hipMemset3D.cpp @@ -0,0 +1,98 @@ +/* +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 memset. +// Also serves as a template for other tests. + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * //Small copy + * RUN: %t -N 10 --memsetval 0x42 + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "test_common.h" + +bool testhipMemset3D(int memsetval,int p_gpuDevice) +{ + size_t numH = 256; + size_t numW = 256; + size_t depth = 1; + size_t pitch_A; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + + + printf ("testhipMemset3D memsetval=%2x device=%d\n", memsetval, p_gpuDevice); + char *A_d; + char *A_h; + bool testResult = true; + hipExtent extent = make_hipExtent(width, numH, depth); + hipPitchedPtr devPitchedPtr; + + HIPCHECK(hipMalloc3D(&devPitchedPtr, extent)); + A_h = (char*)malloc(sizeElements); + HIPASSERT(A_h != NULL); + for (size_t i=0; i