From 409a10d21f6dc3d54a02b014a3e5406fb4bcab7f Mon Sep 17 00:00:00 2001 From: sumanthtg Date: Thu, 17 Sep 2020 23:44:30 +0530 Subject: [PATCH] [dtest] Tests for hipMemset3D/hipMemset3DAsync APIs Negative, Functional and Regression scenarios for - hipMemset3D - hipMemset3DAsync SWDEV-238517 for enhancing hip unit tests Change-Id: Idc5604f728ca1a96ec13876e006120f7a3d69acf [ROCm/hip commit: f692064d1041c0931db1ffa04cfa664f0dc28818] --- .../memory/hipMemset3DFunctional.cpp | 489 ++++++++++++++++++ .../runtimeApi/memory/hipMemset3DNegative.cpp | 251 +++++++++ .../memory/hipMemset3DRegressMultiThread.cpp | 327 ++++++++++++ projects/hip/tests/src/test_common.h | 5 +- 4 files changed, 1071 insertions(+), 1 deletion(-) create mode 100644 projects/hip/tests/src/runtimeApi/memory/hipMemset3DFunctional.cpp create mode 100644 projects/hip/tests/src/runtimeApi/memory/hipMemset3DNegative.cpp create mode 100644 projects/hip/tests/src/runtimeApi/memory/hipMemset3DRegressMultiThread.cpp diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset3DFunctional.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset3DFunctional.cpp new file mode 100644 index 0000000000..3d4e63b850 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset3DFunctional.cpp @@ -0,0 +1,489 @@ +/* +Copyright (c) 2020-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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + + (TestCase 1):: + 1) Passing width as 0 in extent, verify hipMemset3D api returns success and + doesn't modify the buffer passed. + 2) Passing width as 0 in extent, verify hipMemset3DAsync api returns success + and doesn't modify the buffer passed. + + 3) Passing height as 0 in extent, verify hipMemset3D api returns success and + doesn't modify the buffer passed. + 4) Passing height as 0 in extent, verify hipMemset3DAsync api returns success + and doesn't modify the buffer passed. + + 5) Passing depth as 0 in extent, verify hipMemset3D api returns success and + doesn't modify the buffer passed. + 6) Passing depth as 0 in extent, verify hipMemset3DAsync api returns success + and doesn't modify the buffer passed. + + 7) When extent passed with width, height and depth all as zeroes, verify + hipMemset3D api returns success and doesn't modify the buffer passed. + 8) When extent passed with width, height and depth all as zeroes, verify + hipMemset3DAsync api returns success and doesn't modify the buffer passed. + + (TestCase 2):: + 9) Validate data after performing memory set operation with max memset value + for hipMemset3D api. + 10) Validate data after performing memory set operation with max memset value + for hipMemset3DAsync api. + + (TestCase 3):: + 11) Select random slice of 3d array and Memset complete slice with + hipMemset3D api. + 12) Select random slice of 3d array and Memset complete slice with + hipMemset3DAsync api. + + (TestCase 4):: + 13) Seek device pitched ptr to desired portion of 3d array and memset the + portion with hipMemset3D api. + 14) Seek device pitched ptr to desired portion of 3d array and memset the + portion with hipMemset3DAsync api. + +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 + * TEST: %t --tests 1 + * TEST: %t --tests 2 + * TEST: %t --tests 3 + * TEST: %t --tests 4 + * HIT_END + */ + +#include "test_common.h" + +/* + * Defines + */ +#define MEMSETVAL 1 +#define TESTVAL 2 +#define NUMH_EXT 256 +#define NUMW_EXT 100 +#define DEPTH_EXT 10 +#define NUMH_MAX 256 +#define NUMW_MAX 256 +#define DEPTH_MAX 10 +#define ZSIZE_S 32 +#define YSIZE_S 32 +#define XSIZE_S 32 +#define ZSIZE_P 30 +#define YSIZE_P 30 +#define XSIZE_P 30 +#define ZPOS_START 10 +#define ZSET_LEN 10 +#define ZPOS_END 19 +#define YPOS_START 10 +#define YSET_LEN 10 +#define YPOS_END 19 +#define XPOS_START 10 +#define XSET_LEN 10 +#define XPOS_END 19 + + +/** + * Memset with extent passed and verify data to be intact + */ +bool testMemsetWithExtent(bool bAsync, hipExtent tstExtent) { + hipPitchedPtr devPitchedPtr; + bool testPassed = true; + hipError_t ret; + char *A_h; + size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT; + size_t width = numW * sizeof(char); + hipExtent extent = make_hipExtent(width, numH, depth); + + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + + A_h = reinterpret_cast(malloc(sizeElements)); + HIPASSERT(A_h != NULL); + memset(A_h, 0, sizeElements); + + HIPCHECK(hipMalloc3D(&devPitchedPtr, extent)); + if (bAsync) { + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + if ((ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream)) + != hipSuccess) { + printf("testMemsetWithExtent(%zu,%zu,%zu) Async: Expected to return" + " success but returned Error: '%s'(%d)\n", extent.width, + extent.height, extent.depth, hipGetErrorString(ret), ret); + testPassed &= false; + } + + if ((ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream)) + != hipSuccess) { + printf("testMemsetWithExtent(%zu,%zu,%zu) Async: Expected to return" + " success but returned Error: '%s'(%d)\n", tstExtent.width, + tstExtent.height, tstExtent.depth, hipGetErrorString(ret), ret); + testPassed &= false; + } + + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipStreamDestroy(stream)); + } else { + if ((ret = hipMemset3D(devPitchedPtr, MEMSETVAL, extent)) + != hipSuccess) { + printf("testMemsetWithExtent(%zu,%zu,%zu) : Expected to return" + " success but returned Error: '%s'(%d)\n", extent.width, + extent.height, extent.depth, hipGetErrorString(ret), ret); + testPassed &= false; + } + + if ((ret = hipMemset3D(devPitchedPtr, TESTVAL, tstExtent)) + != hipSuccess) { + printf("testMemsetWithExtent(%zu,%zu,%zu) : Expected to return" + " success but returned Error: '%s'(%d)\n", tstExtent.width, + tstExtent.height, tstExtent.depth, hipGetErrorString(ret), ret); + testPassed &= false; + } + } + + if (testPassed) { + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.srcPtr = devPitchedPtr; + myparms.extent = extent; + #ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); + #else + myparms.kind = hipMemcpyDeviceToHost; + #endif + + HIPCHECK(hipMemcpy3D(&myparms)); + + for (int i = 0; i < elements; i++) { + if (A_h[i] != MEMSETVAL) { + testPassed = false; + printf("testMemsetWithExtent: mismatch at index:%d computed:%02x, " + "memsetval:%02x\n", i, static_cast(A_h[i]), + static_cast(MEMSETVAL)); + break; + } + } + } + + HIPCHECK(hipFree(devPitchedPtr.ptr)); + free(A_h); + return testPassed; +} + +/** + * Validates data after performing memory set operation with max memset value + */ +bool testMemsetMaxValue(bool bAsync) { + hipPitchedPtr devPitchedPtr; + bool testPassed = true; + unsigned char *A_h; + int memsetval = std::numeric_limits::max(); + size_t numH = NUMH_MAX, numW = NUMW_MAX, depth = DEPTH_MAX; + size_t width = numW * sizeof(unsigned char); + hipExtent extent = make_hipExtent(width, numH, depth); + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + + A_h = reinterpret_cast (malloc(sizeElements)); + HIPASSERT(A_h != NULL); + memset(A_h, 0, sizeElements); + + HIPCHECK(hipMalloc3D(&devPitchedPtr, extent)); + if (bAsync) { + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + HIPCHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipStreamDestroy(stream)); + } else { + HIPCHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); + } + + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.srcPtr = devPitchedPtr; + myparms.extent = extent; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + HIPCHECK(hipMemcpy3D(&myparms)); + + for (int i = 0; i < elements; i++) { + if (A_h[i] != memsetval) { + testPassed = false; + printf("testMemsetMaxValue: mismatch at index:%d computed:%02x, " + "memsetval:%02x\n", i, static_cast(A_h[i]), + static_cast(memsetval)); + break; + } + } + HIPCHECK(hipFree(devPitchedPtr.ptr)); + free(A_h); + return testPassed; +} + +/** + * Function seeks device ptr to random slice and performs Memset operation + * on the slice selected. + */ +bool seekAndSet3DArraySlice(bool bAsync) { + char array3D[ZSIZE_S][YSIZE_S][XSIZE_S] = {0}; + bool testPassed = true; + dim3 arr_dimensions = dim3(ZSIZE_S, YSIZE_S, XSIZE_S); + hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x, + arr_dimensions.y, arr_dimensions.z); + hipPitchedPtr devicePitchedPointer; + int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL; + + HIPCHECK(hipMalloc3D(&devicePitchedPointer, extent)); + HIPCHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); + + // select random slice for memset + unsigned int seed = time(NULL); + int slice_index = rand_r(&seed) % ZSIZE_S; + + printf("memset3d for sliceindex %d\n", slice_index); + + // Get attributes from device pitched pointer + size_t pitch = devicePitchedPointer.pitch; + size_t slicePitch = pitch * extent.height; + + // Point devptr to selected slice + char *devPtrSlice = (reinterpret_cast(devicePitchedPointer.ptr)) + + slice_index * slicePitch; + hipExtent extentSlice = make_hipExtent(sizeof(char) * arr_dimensions.x, + arr_dimensions.y, 1); + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrSlice, pitch, + arr_dimensions.x, arr_dimensions.y); + + if (bAsync) { + // Memset selected slice (Async) + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + HIPCHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, + extentSlice, stream)); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipStreamDestroy(stream)); + } else { + // Memset selected slice + HIPCHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, extentSlice)); + } + + // Copy result back to host buffer + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x, + arr_dimensions.x, arr_dimensions.y); + myparms.srcPtr = devicePitchedPointer; + myparms.extent = extent; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + HIPCHECK(hipMemcpy3D(&myparms)); + + for (int z = 0; z < ZSIZE_S; z++) { + for (int y = 0; y < YSIZE_S; y++) { + for (int x = 0; x < XSIZE_S; x++) { + if (z == slice_index) { + if (array3D[z][y][x] != memsetval4seeked) { + testPassed = false; + printf("seekAndSet3DArray Slice: mismatch at index: Arr(%d,%d,%d)" + " computed:%02x, memsetval:%02x\n", z, y, x, + array3D[z][y][x], memsetval4seeked); + break; + } + } else { + if (array3D[z][y][x] != memsetval) { + testPassed = false; + printf("seekAndSet3DArray Slice: mismatch at index: Arr(%d,%d,%d)" + " computed:%02x, memsetval:%02x\n", z, y, x, + array3D[z][y][x], memsetval); + break; + } + } + } + } + } + + HIPCHECK(hipFree(devicePitchedPointer.ptr)); + return testPassed; +} + +/** + * Function seeks device ptr to selected portion of 3d array + * and performs Memset operation on the portion. + */ +bool seekAndSet3DArrayPortion(bool bAsync) { + char array3D[ZSIZE_P][YSIZE_P][XSIZE_P] = {0}; + bool testPassed = true; + dim3 arr_dimensions = dim3(ZSIZE_P, YSIZE_P, XSIZE_P); + hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x, + arr_dimensions.y, arr_dimensions.z); + hipPitchedPtr devicePitchedPointer; + int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL; + + HIPCHECK(hipMalloc3D(&devicePitchedPointer, extent)); + HIPCHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); + + // For memsetting extent/size(10,10,10) in the mid portion of cube(30,30,30), + // seek device ptr to (10,10,10) and then memset 10 bytes across x,y,z axis. + size_t pitch = devicePitchedPointer.pitch; + size_t slicePitch = pitch * extent.height; + int slice_index = ZPOS_START, y = YPOS_START, x = XPOS_START; + + // Select 10th slice + char *devPtrSlice = (reinterpret_cast(devicePitchedPointer.ptr)) + + slice_index * slicePitch; + + // Now select row at height as 10 + char *current_row = reinterpret_cast(devPtrSlice + y * pitch); + + // Now select index of selected row as 10 + char *devPtrIndexed = ¤t_row[x]; + + // Make dev Pitchedptr, extent + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrIndexed, pitch, + arr_dimensions.x, arr_dimensions.y); + hipExtent setExtent = make_hipExtent(sizeof(char) * XSET_LEN, YSET_LEN, + ZSET_LEN); + + if (bAsync) { + // Memset selected portion (Async) + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + HIPCHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, + setExtent, stream)); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipStreamDestroy(stream)); + } else { + // Memset selected portion + HIPCHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, setExtent)); + } + + // Copy result back to host buffer + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x, + arr_dimensions.x, arr_dimensions.y); + myparms.srcPtr = devicePitchedPointer; + myparms.extent = extent; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + HIPCHECK(hipMemcpy3D(&myparms)); + + for (int z = 0; z < ZSIZE_P; z++) { + for (int y = 0; y < YSIZE_P; y++) { + for (int x = 0; x < XSIZE_P; x++) { + if ((z >= ZPOS_START && z <= ZPOS_END) && + (y >= YPOS_START && y <= YPOS_END) && + (x >= XPOS_START && x <= XPOS_END)) { + if (array3D[z][y][x] != memsetval4seeked) { + testPassed = false; + printf("seekAndSet3DArray Portion: mismatch at index: Arr(%d,%d,%d)" + " computed:%02x, memsetval:%02x\n", z, y, x, + array3D[z][y][x], memsetval4seeked); + break; + } + } else { + if (array3D[z][y][x] != memsetval) { + testPassed = false; + printf("seekAndSet3DArray Portion: mismatch at index: Arr(%d,%d,%d)" + " computed:%02x, memsetval:%02x\n", z, y, x, + array3D[z][y][x], memsetval); + break; + } + } + } + } + } + + HIPCHECK(hipFree(devicePitchedPointer.ptr)); + return testPassed; +} + + +int main(int argc, char *argv[]) { + HipTest::parseStandardArguments(argc, argv, true); + bool TestPassed = true; + + if (p_tests == 1) { + hipExtent testExtent; + size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT; + + // Memset with extent width(0) and verify data to be intact + testExtent = make_hipExtent(0, numH, depth); + TestPassed &= testMemsetWithExtent(0, testExtent); + TestPassed &= testMemsetWithExtent(1, testExtent); + + // Memset with extent height(0) and verify data to be intact + testExtent = make_hipExtent(numW, 0, depth); + TestPassed &= testMemsetWithExtent(0, testExtent); + TestPassed &= testMemsetWithExtent(1, testExtent); + + // Memset with extent depth(0) and verify data to be intact + testExtent = make_hipExtent(numW, numH, 0); + TestPassed &= testMemsetWithExtent(0, testExtent); + TestPassed &= testMemsetWithExtent(1, testExtent); + + // Memset with extent width,height,depth as 0 and verify data to be intact + testExtent = make_hipExtent(0, 0, 0); + TestPassed &= testMemsetWithExtent(0, testExtent); + TestPassed &= testMemsetWithExtent(1, testExtent); + } else if (p_tests == 2) { + // Memset with max unsigned char and verify memset is success + TestPassed &= testMemsetMaxValue(0); + TestPassed &= testMemsetMaxValue(1); + } else if (p_tests == 3) { + // Seek and set random slice of 3d array + TestPassed &= seekAndSet3DArraySlice(0); + TestPassed &= seekAndSet3DArraySlice(1); + } else if (p_tests == 4) { + // Memset selected portion of 3d array + TestPassed &= seekAndSet3DArrayPortion(0); + TestPassed &= seekAndSet3DArrayPortion(1); + } else { + printf("Didnt receive any valid option. Try options 1 to 4\n"); + TestPassed = false; + } + + if (TestPassed) { + passed(); + } else { + failed("hipMemset3DFunctional validation Failed!"); + } +} diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset3DNegative.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset3DNegative.cpp new file mode 100644 index 0000000000..7dcf083bd5 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset3DNegative.cpp @@ -0,0 +1,251 @@ +/* +Copyright (c) 2020-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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + + (TestCase 1):: + 1) Test hipMemset3D() with uninitialized devPitchedPtr. + 2) Test hipMemset3DAsync() with uninitialized devPitchedPtr. + + (TestCase 2):: + 3) Reset devPitchedPtr to zero and check return value for hipMemset3D(). + 4) Reset devPitchedPtr to zero and check return value for hipMemset3DAsync(). + + (TestCase 3) + 5) Test hipMemset3D() with extent.width as max size_t and keeping height, + depth as valid values. + 6) Test hipMemset3DAsync() with extent.width as max size_t and keeping height, + depth as valid values. + 7) Test hipMemset3D() with extent.height as max size_t and keeping width, + depth as valid values. + 8) Test hipMemset3DAsync() with extent.height as max size_t and keeping width, + depth as valid values. + 9) Test hipMemset3D() with extent.depth as max size_t and keeping height, + width as valid values. +10) Test hipMemset3DAsync() with extent.depth as max size_t and keeping height, + width as valid values. + + (TestCase 4) +11) Device Ptr out bound and extent(0) passed for hipMemset3D(). +12) Device Ptr out bound and extent(0) passed for hipMemset3DAsync(). + + (TestCase 5) +13) Device Ptr out bound and valid extent passed for hipMemset3D(). +14) Device Ptr out bound and valid extent passed for hipMemset3DAsync(). +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 + * TEST: %t --tests 1 + * TEST: %t --tests 2 + * TEST: %t --tests 3 + * TEST: %t --tests 4 + * TEST: %t --tests 5 + * HIT_END + */ + +#include "test_common.h" + +int main(int argc, char *argv[]) { + HipTest::parseStandardArguments(argc, argv, true); + hipStream_t stream; + hipError_t ret; + hipPitchedPtr devPitchedPtr; + bool TestPassed = true; + int memsetval = 1; + size_t numH = 256; + size_t numW = 256; + size_t depth = 10; + size_t width = numW * sizeof(char); + hipExtent extent = make_hipExtent(width, numH, depth); + + HIPCHECK(hipStreamCreate(&stream)); + HIPCHECK(hipMalloc3D(&devPitchedPtr, extent)); + + if (p_tests == 1) { + // Use uninitialized devpitched ptr + hipPitchedPtr devPitchedUnPtr; + + if ((ret = hipMemset3D(devPitchedUnPtr, memsetval, extent)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "uninit devpitched ptr. Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipMemset3DAsync(devPitchedUnPtr, memsetval, extent, stream)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "uninit devpitched ptr(Async). Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + } else if (p_tests == 2) { + // Reset devPitchedPtr to zero + hipPitchedPtr rdevPitchedPtr = {0}; + + if ((ret = hipMemset3D(rdevPitchedPtr, memsetval, extent)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "rdevPitchedPtr(0). Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipMemset3DAsync(rdevPitchedPtr, memsetval, extent, stream)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "rdevPitchedPtr(0). Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + } else if (p_tests == 3) { + // Pass extent fields as max size_t + hipExtent extMW = make_hipExtent(std::numeric_limits::max(), + numH, + depth); + hipExtent extMH = make_hipExtent(width, + std::numeric_limits::max(), + depth); + hipExtent extMD = make_hipExtent(width, + numH, + std::numeric_limits::max()); + + if ((ret = hipMemset3D(devPitchedPtr, memsetval, extMW)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "extent.width max(size_t). Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMW, stream)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "extent.width max(size_t) Async. Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipMemset3D(devPitchedPtr, memsetval, extMH)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "extent.height max(size_t). Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMH, stream)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "extent.height max(size_t) Async. Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipMemset3D(devPitchedPtr, memsetval, extMD)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "extent.depth max(size_t). Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMD, stream)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned for " + "extent.depth max(size_t) Async. Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + } else if (p_tests == 4) { + // Device Ptr out bound and extent(0) passed for memset + + size_t pitch = devPitchedPtr.pitch; + size_t slicePitch = pitch * extent.height; + + // Point devptr to end of allocated memory + char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) + + depth * slicePitch; + + // Advance devptr further to go out of boundary + devPtrMod = devPtrMod + 10; + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, + numW * sizeof(char), numH); + hipExtent extent0 = {0}; + if ((ret = hipMemset3D(modDevPitchedPtr, memsetval, extent0)) + != hipSuccess) { + printf("ArgValidation : Inappropriate error value returned when " + "deviceptr goes out of boundary. Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent0, stream)) + != hipSuccess) { + printf("ArgValidation : Inappropriate error value returned when " + "deviceptr goes out of boundary Async. Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + } else if (p_tests == 5) { + // Device Ptr out bound and valid extent passed for memset + + size_t pitch = devPitchedPtr.pitch; + size_t slicePitch = pitch * extent.height; + + // Point devptr to end of allocated memory + char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) + + depth * slicePitch; + + // Advance devptr further to go out of boundary + devPtrMod = devPtrMod + 10; + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, + numW * sizeof(char), numH); + if ((ret = hipMemset3D(modDevPitchedPtr, memsetval, extent)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned when " + "deviceptr goes out of boundary. Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent, stream)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropriate error value returned when " + "deviceptr goes out of boundary Async. Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } + } else { + printf("Didnt receive any valid option. Try options 1 to 5\n"); + TestPassed = false; + } + + HIPCHECK(hipStreamDestroy(stream)); + HIPCHECK(hipFree(devPitchedPtr.ptr)); + + if (TestPassed) { + passed(); + } else { + failed("hipMemset3DNegative validation Failed!"); + } +} diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset3DRegressMultiThread.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset3DRegressMultiThread.cpp new file mode 100644 index 0000000000..b06a77ef48 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset3DRegressMultiThread.cpp @@ -0,0 +1,327 @@ +/* +Copyright (c) 2020-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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + + (TestCase 1):: + 1) Validate Async behavior of hipMemset3DAsync with commands queued + concurrently from multiple threads. + 2) Validate hipMemset3DAsync behavior when api is queued along with kernel + function operating on same memory. + + (TestCase 2):: + 3) Perform regression of hipMemset3D api in loop with device memory allocated + on different gpus. + 4) Perform regression of hipMemset3DAsync api in loop with device memory + allocated on different gpus. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 + * TEST: %t --tests 1 + * HIT_END + */ + +#include +#include +#include +#include +#include "test_common.h" + + +/* + * Defines + */ +#define MAX_REGRESS_ITERS 20 + +/** + * kernel function sets device memory with value passed + */ +__global__ void func_set_value(hipPitchedPtr devicePitchedPointer, + hipExtent extent, + unsigned char val) { + // Index Calculation + int x = threadIdx.x + blockDim.x * blockIdx.x; + int y = threadIdx.y + blockDim.y * blockIdx.y; + int z = threadIdx.z + blockDim.z * blockIdx.z; + + // Get attributes from device pitched pointer + char *devicePointer = reinterpret_cast(devicePitchedPointer.ptr); + size_t pitch = devicePitchedPointer.pitch; + size_t slicePitch = pitch * extent.height; + + // Loop over the device buffer + if (z < extent.depth) { + char *current_slice_index = devicePointer + z * slicePitch; + if (y < extent.height) { + // Get data array containing all elements from the current row + char *current_row = reinterpret_cast(current_slice_index + + y * pitch); + if (x < extent.width) { + current_row[x] = val; + } + } + } +} + +/** + * Fetches Gpu device count + */ +void getDeviceCount(int *pdevCnt) { +#ifdef __linux__ + int fd[2], val = 0; + pid_t childpid; + + // create pipe descriptors + pipe(fd); + + // disable visible_devices env from shell + unsetenv("ROCR_VISIBLE_DEVICES"); + unsetenv("HIP_VISIBLE_DEVICES"); + + childpid = fork(); + + if (childpid > 0) { // Parent + close(fd[1]); + // parent will wait to read the device cnt + read(fd[0], &val, sizeof(val)); + + // close the read-descriptor + close(fd[0]); + + // wait for child exit + wait(NULL); + + *pdevCnt = val; + } else if (!childpid) { // Child + int devCnt = 1; + // writing only, no need for read-descriptor + close(fd[0]); + + HIPCHECK(hipGetDeviceCount(&devCnt)); + // send the value on the write-descriptor: + write(fd[1], &devCnt, sizeof(devCnt)); + + // close the write descriptor: + close(fd[1]); + exit(0); + } else { // failure + *pdevCnt = 1; + return; + } + +#else + HIPCHECK(hipGetDeviceCount(pdevCnt)); +#endif +} + +/** + * Performs api regression in loop + */ +bool loopRegression(bool bAsync) { + bool testPassed = true; + char *A_h; + int memsetval = 1, numGpu = 0, hasPeerAccess = 0; + size_t numH = 256, numW = 100, depth = 10; + size_t width = numW * sizeof(char); + hipExtent extent = make_hipExtent(width, numH, depth); + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + std::vector devPitchedPtrlist; + hipPitchedPtr pitchedPtr, devpPtr; + + A_h = reinterpret_cast(malloc(sizeElements)); + HIPASSERT(A_h != NULL); + memset(A_h, 0, sizeElements); + + // Populate hipMemcpy3D parameters + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.extent = extent; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + getDeviceCount(&numGpu); + + // Alloc 3D arrays in all GPUs + for (int j = 0; j < numGpu; j++) { + HIPCHECK(hipSetDevice(j)); + HIPCHECK(hipMalloc3D(&pitchedPtr, extent)); + devPitchedPtrlist.push_back(pitchedPtr); + } + + for (int itern = 0; itern < MAX_REGRESS_ITERS; itern++) { + // Validate hipMemset3D data consistency in multiple iters + for (int i = 0; i < numGpu; i++) { + for (int j = 0; j < numGpu; j++) { + HIPCHECK(hipDeviceCanAccessPeer(&hasPeerAccess, i, j)); + if (!hasPeerAccess) { + // Skip and continue if no peer access + continue; + } + HIPCHECK(hipSetDevice(i)); + devpPtr = devPitchedPtrlist[j]; + HIPCHECK(hipMemset3D(devpPtr, 0, extent)); + + if (bAsync) { + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + HIPCHECK(hipMemset3DAsync(devpPtr, memsetval, extent, stream)); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipStreamDestroy(stream)); + } else { + HIPCHECK(hipMemset3D(devpPtr, memsetval, extent)); + } + + myparms.srcPtr = devpPtr; + memset(A_h, 0, sizeElements); + HIPCHECK(hipMemcpy3D(&myparms)); + + for (int indx = 0; indx < elements; indx++) { + if (A_h[indx] != memsetval) { + testPassed = false; + printf("RegressIter : mismatch at index:%d computed:%02x, " + "memsetval:%02x\n", indx, static_cast(A_h[indx]), + static_cast(memsetval)); + break; + } + } + } + } + } + + for (int j = 0; j < numGpu; j++) { + HIPCHECK(hipFree(devPitchedPtrlist[j].ptr)); + } + + free(A_h); + return testPassed; +} + + +/** + * Thread function queues kernel function and memset cmds + */ +void threadFunc(hipStream_t stream, hipPitchedPtr devpPtr, int memsetval, + int testval, hipExtent extent, hipMemcpy3DParms myparms) { + // Kernel Launch Configuration + dim3 threadsPerBlock = dim3(8, 8, 8); + dim3 blocks; + blocks = dim3((extent.width + threadsPerBlock.x - 1) / threadsPerBlock.x, + (extent.height + threadsPerBlock.y - 1) / threadsPerBlock.y, + (extent.depth + threadsPerBlock.z - 1) / threadsPerBlock.z); + + hipLaunchKernelGGL(func_set_value, dim3(blocks), dim3(threadsPerBlock), 0, + stream, devpPtr, extent, memsetval); + HIPCHECK(hipMemset3DAsync(devpPtr, testval, extent, stream)); + HIPCHECK(hipMemcpy3DAsync(&myparms, stream)); +} + +/** + * Async commands queued concurrently and executed + */ +bool validateAsyncConcurrencyMthread() { + bool testPassed = true; + char *A_h; + int memsetval = 1, numGpu = 0, testval = 2; + size_t numH = 256, numW = 100, depth = 10; + size_t width = numW * sizeof(char); + hipExtent extent = make_hipExtent(width, numH, depth); + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + hipPitchedPtr devpPtr; + hipStream_t stream; + + HIPCHECK(hipStreamCreate(&stream)); + HIPCHECK(hipMalloc3D(&devpPtr, extent)); + + A_h = reinterpret_cast(malloc(sizeElements)); + HIPASSERT(A_h != NULL); + memset(A_h, 0, sizeElements); + + // Populate hipMemcpy3D parameters + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.srcPtr = devpPtr; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.extent = extent; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + std::vector threadlist; + + // Queue cmds concurrently from multiple threads on same stream + for (int i = 0; i < MAX_THREADS; i++) { + threadlist.push_back(std::thread(threadFunc, stream, devpPtr, memsetval, + testval, extent, myparms)); + } + + for (auto &t : threadlist) { + t.join(); + } + + HIPCHECK(hipStreamSynchronize(stream)); + + for (int k = 0 ; k < elements ; k++) { + if (A_h[k] != testval) { + printf("validateAsyncConcurrencyMthread: Test failed\n"); + testPassed = false; + break; + } + } + + HIPCHECK(hipStreamDestroy(stream)); + free(A_h); + HIPCHECK(hipFree(devpPtr.ptr)); + return testPassed; +} + + +int main(int argc, char *argv[]) { + HipTest::parseStandardArguments(argc, argv, true); + bool TestPassed = true; + + if (p_tests == 1) { + TestPassed = validateAsyncConcurrencyMthread(); + } else if (p_tests == 2) { + /* TODO : Loop regression test auto execution in HIT is currently disabled. + To be enabled back after HIP API fix */ + TestPassed &= loopRegression(0); + TestPassed &= loopRegression(1); + } else { + printf("Didnt receive any valid option. Try options 1 to 2\n"); + TestPassed = false; + } + + if (TestPassed) { + passed(); + } else { + failed("hipMemset3DRegressMultiThread() validation Failed!"); + } +} diff --git a/projects/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h index 21a4c45ac8..ded14d6fad 100755 --- a/projects/hip/tests/src/test_common.h +++ b/projects/hip/tests/src/test_common.h @@ -50,10 +50,13 @@ THE SOFTWARE. #define KCYN "\x1B[36m" #define KWHT "\x1B[37m" - // HIP Skip Return code set at cmake +// HIP Skip Return code set at cmake #define HIP_SKIP_RETURN_CODE 127 #define HIP_ENABLE_SKIP_TESTS 0 +// Recommended thresholds for Tests +#define MAX_THREADS 100 + inline bool hip_skip_tests_enabled() { return HIP_ENABLE_SKIP_TESTS; }