diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index d9e688d167..093b6d8ba9 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -33,6 +33,10 @@ set(TEST_SRC hipMemoryAllocateCoherent.cc hipMallocManaged_MultiScenario.cc hipManagedKeyword.cc + hipMemsetInvalidPtr.cc + hipMemset.cc + hipMemsetAsyncMultiThread.cc + hipMemsetAsyncAndKernel.cc ) else() set(TEST_SRC @@ -66,6 +70,10 @@ set(TEST_SRC hipMemoryAllocateCoherent.cc hipMallocManaged_MultiScenario.cc hipManagedKeyword.cc + hipMemsetInvalidPtr.cc + hipMemset.cc + hipMemsetAsyncMultiThread.cc + hipMemsetAsyncAndKernel.cc ) endif() # Create shared lib of all tests diff --git a/catch/unit/memory/hipMemset.cc b/catch/unit/memory/hipMemset.cc new file mode 100644 index 0000000000..600a008436 --- /dev/null +++ b/catch/unit/memory/hipMemset.cc @@ -0,0 +1,281 @@ +/* + * Copyright (c) 2021 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 : + 1) Test hipMemset small size buffers with unique memset values. + 2) Test hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32 apis with unique + number of elements and memset values. + 3) Test hipMemsetAsync, hipMemsetD8Async, hipMemsetD16Async, hipMemsetD32Async + apis with unique number of elements and memset values. + 4) Test two memset async operations at the same time. +*/ + + +#include + + +// Table with unique number of elements and memset values. +// (N, memsetval, memsetD32val, memsetD16val, memsetD8val) +typedef std::tuple tupletype; +static constexpr std::initializer_list tableItems { + std::make_tuple((4*1024*1024), 0x42, 0xDEADBEEF, 0xDEAD, 0xDE), + std::make_tuple((10) , 0x42, 0x101 , 0x10, 0x1), + std::make_tuple((10013) , 0x5a, 0xDEADBEEF, 0xDEAD, 0xDE), + std::make_tuple((256*1024*1024), 0xa6, 0xCAFEBABE, 0xCAFE, 0xCA) + }; + +enum MemsetType { + hipMemsetTypeDefault, + hipMemsetTypeD8, + hipMemsetTypeD16, + hipMemsetTypeD32 +}; + +template +static bool testhipMemset(T *A_h, T *A_d, T memsetval, enum MemsetType type, + size_t numElements) { + size_t Nbytes = numElements * sizeof(T); + bool testResult = true; + constexpr auto MAX_OFFSET = 3; // To memset on unaligned ptr. + + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + A_h = reinterpret_cast (malloc(Nbytes)); + REQUIRE(A_h != nullptr); + + for (int offset = MAX_OFFSET; offset >= 0; offset --) { + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset(A_d + offset, memsetval, numElements - offset)); + + } else if (type == hipMemsetTypeD8) { + HIP_CHECK(hipMemsetD8((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset)); + + } else if (type == hipMemsetTypeD16) { + HIP_CHECK(hipMemsetD16((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset)); + + } else if (type == hipMemsetTypeD32) { + HIP_CHECK(hipMemsetD32((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset)); + } + + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (size_t i = offset; i < numElements; i++) { + if (A_h[i] != memsetval) { + testResult = false; + CAPTURE(i, A_h[i], memsetval); + break; + } + } + } + + HIP_CHECK(hipFree(A_d)); + free(A_h); + return testResult; +} + + +template +static bool testhipMemsetAsync(T *A_h, T *A_d, T memsetval, + enum MemsetType type, size_t numElements) { + size_t Nbytes = numElements * sizeof(T); + bool testResult = true; + constexpr auto MAX_OFFSET = 3; // To memset on unaligned ptr. + hipStream_t stream; + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + A_h = reinterpret_cast (malloc(Nbytes)); + REQUIRE(A_h != nullptr); + + for (int offset = MAX_OFFSET; offset >= 0; offset --) { + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemsetAsync(A_d + offset, memsetval, numElements - offset, + stream)); + + } else if (type == hipMemsetTypeD8) { + HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset, stream)); + + } else if (type == hipMemsetTypeD16) { + HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset, stream)); + + } else if (type == hipMemsetTypeD32) { + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset, stream)); + } + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (size_t i = offset; i < numElements; i++) { + if (A_h[i] != memsetval) { + testResult = false; + CAPTURE(i, A_h[i], memsetval); + break; + } + } + } + + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + free(A_h); + return testResult; +} + + +/** + * Test hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32 apis with unique + * number of elements and memset values. + */ +TEST_CASE("Unit_hipMemset_SetMemoryWithOffset") { + char memsetval; + int memsetD32val; + int16_t memsetD16val; + char memsetD8val; + size_t N; + bool ret; + + std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) = + GENERATE(table(tableItems)); + + + SECTION("Memset with hipMemsetTypeDefault") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemset(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD32") { + int32_t *iA_d{nullptr}, *iA_h{nullptr}; + ret = testhipMemset(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD16") { + int16_t *siA_d{nullptr}, *siA_h{nullptr}; + ret = testhipMemset(siA_h, siA_d, memsetD16val, hipMemsetTypeD16, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD8") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemset(cA_h, cA_d, memsetD8val, hipMemsetTypeD8, N); + REQUIRE(ret == true); + } +} + + +/** + * Test hipMemsetAsync, hipMemsetD8Async, hipMemsetD16Async, hipMemsetD32Async + * apis with unique number of elements and memset values. + */ +TEST_CASE("Unit_hipMemsetAsync_SetMemoryWithOffset") { + char memsetval; + int memsetD32val; + int16_t memsetD16val; + char memsetD8val; + size_t N; + bool ret; + + std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) = + GENERATE(table(tableItems)); + + + SECTION("Memset with hipMemsetTypeDefault") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemsetAsync(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD32") { + int32_t *iA_d{nullptr}, *iA_h{nullptr}; + ret = testhipMemsetAsync(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD16") { + int16_t *siA_d{nullptr}, *siA_h{nullptr}; + ret = testhipMemsetAsync(siA_h, siA_d, memsetD16val, hipMemsetTypeD16, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD8") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemsetAsync(cA_h, cA_d, memsetD8val, hipMemsetTypeD8, N); + REQUIRE(ret == true); + } +} + +/** + * Test hipMemset small size buffers with unique memset values. + */ +TEST_CASE("Unit_hipMemset_SmallBufferSizes") { + char *A_d, *A_h; + constexpr int memsetval = 0x24; + + auto numElements = GENERATE(range(1, 4)); + int numBytes = numElements * sizeof(char); + + HIP_CHECK(hipMalloc(&A_d, numBytes)); + A_h = reinterpret_cast (malloc(numBytes)); + + HIP_CHECK(hipMemset(A_d, memsetval, numBytes)); + HIP_CHECK(hipMemcpy(A_h, A_d, numBytes, hipMemcpyDeviceToHost)); + + for (int i = 0; i < numBytes; i++) { + if (A_h[i] != memsetval) { + INFO("Mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(false); + } + } + + HIP_CHECK(hipFree(A_d)); + free(A_h); +} + + +/** + * Test two memset async operations at the same time. + */ +TEST_CASE("Unit_hipMemset_2AsyncOperations") { + std::vector v; + v.resize(2048); + float* p2, *p3; + hipMalloc(reinterpret_cast(&p2), 4096 + 4096*2); + p3 = p2+2048; + hipStream_t s; + hipStreamCreate(&s); + hipMemsetAsync(p2, 0, 32*32*4, s); + hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s); + hipStreamSynchronize(s); + for (int i = 0; i < 256; ++i) { + hipMemsetAsync(p2, 0, 32*32*4, s); + hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s); + } + hipStreamSynchronize(s); + hipDeviceSynchronize(); + hipMemcpy(&v[0], p2, 1024, hipMemcpyDeviceToHost); + hipMemcpy(&v[1024], p3, 1024, hipMemcpyDeviceToHost); + + REQUIRE(v[0] == 0); + REQUIRE(v[1024] == 1.75f); +} diff --git a/catch/unit/memory/hipMemsetAsyncAndKernel.cc b/catch/unit/memory/hipMemsetAsyncAndKernel.cc new file mode 100644 index 0000000000..adbd4a3964 --- /dev/null +++ b/catch/unit/memory/hipMemsetAsyncAndKernel.cc @@ -0,0 +1,193 @@ +/* + * Copyright (c) 2021 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. +*/ + +/* + * Test for checking order of execution of device kernel and + * hipMemsetAsync apis on all gpus + */ + +#include +#include +#include + +#define ITER 6 +#define N 1024 * 1024 + +constexpr auto blocksPerCU = 6; // to hide latency +constexpr auto threadsPerBlock = 256; +static unsigned blocks = 0; + + +template +class MemSetKernelTest { + public: + T *A_h, *B_d, *B_h, *C_d; + T memSetVal; + size_t Nbytes; + bool testResult = true; + int validateCount = 0; + hipStream_t stream; + + void memAllocate(T memSetValue) { + memSetVal = memSetValue; + Nbytes = N * sizeof(T); + + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(A_h != nullptr); + HIP_CHECK(hipMalloc(&B_d , Nbytes)); + B_h = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(B_h != nullptr); + HIP_CHECK(hipMalloc(&C_d , Nbytes)); + + for (int i = 0 ; i < N ; i++) { + B_h[i] = i; + } + HIP_CHECK(hipMemcpy(B_d , B_h , Nbytes , hipMemcpyHostToDevice)); + HIP_CHECK(hipStreamCreate(&stream)); + } + + void memDeallocate() { + HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); + free(B_h); free(A_h); + HIP_CHECK(hipStreamDestroy(stream)); + } + + void validateExecutionOrder() { + for (int p = 0 ; p < N ; p++) { + if (A_h[p] == memSetVal) { + validateCount+= 1; + } + } + } + + bool resultAfterAllIterations() { + testResult = (validateCount == (ITER * N)) ? true : false; + memDeallocate(); + return testResult; + } +}; + +static bool testhipMemsetAsyncWithKernel() { + MemSetKernelTest obj; + constexpr char memsetval = 0x42; + + obj.memAllocate(memsetval); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); + HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +static bool testhipMemsetD32AsyncWithKernel() { + MemSetKernelTest obj; + constexpr int memsetD32val = 0xDEADBEEF; + + obj.memAllocate(memsetD32val); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)obj.C_d , obj.memSetVal, + N, obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +static bool testhipMemsetD16AsyncWithKernel() { + MemSetKernelTest obj; + constexpr int16_t memsetD16val = 0xDEAD; + + obj.memAllocate(memsetD16val); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); + HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)obj.C_d , obj.memSetVal, + N, obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +static bool testhipMemsetD8AsyncWithKernel() { + MemSetKernelTest obj; + constexpr char memsetD8val = 0xDE; + + obj.memAllocate(memsetD8val); + for (int k = 0; k < ITER; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); + HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)obj.C_d, obj.memSetVal, + N, obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + + +/* + * Test for checking order of execution of device kernel and + * hipMemsetAsync apis on all gpus + */ +TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") { + int numDevices = 0; + bool ret; + + blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGetDeviceCount(&numDevices)); + REQUIRE(numDevices > 0); + + auto devNum = GENERATE_COPY(range(0, numDevices)); + HIP_CHECK(hipSetDevice(devNum)); + + SECTION("hipMemsetAsync With Kernel") { + ret = testhipMemsetAsyncWithKernel(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD32Async With Kernel") { + ret = testhipMemsetD32AsyncWithKernel(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD16Async With Kernel") { + ret = testhipMemsetD16AsyncWithKernel(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD8Async With Kernel") { + ret = testhipMemsetD8AsyncWithKernel(); + REQUIRE(ret == true); + } +} diff --git a/catch/unit/memory/hipMemsetAsyncMultiThread.cc b/catch/unit/memory/hipMemsetAsyncMultiThread.cc new file mode 100644 index 0000000000..68deacda92 --- /dev/null +++ b/catch/unit/memory/hipMemsetAsyncMultiThread.cc @@ -0,0 +1,243 @@ +/* + * Copyright (c) 2021 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. + */ + +/* + * Test that validates functionality of hipmemsetAsync apis over multi threads + */ + +#include +#include + + +#define NUM_THREADS 20 +#define ITER 10 +#define N (4*1024*1024) + + +template +class MemSetAsyncMthreadTest { + public: + T *A_h, *A_d, *B_h; + T memSetVal; + size_t Nbytes; + bool testResult = true; + int validateCount = 0; + hipStream_t stream; + + void memAllocate(T memSetValue) { + memSetVal = memSetValue; + Nbytes = N * sizeof(T); + + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(A_h != nullptr); + + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + B_h = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(B_h != nullptr); + + HIP_CHECK(hipStreamCreate(&stream)); + } + + void threadCompleteStatus() { + for (int k = 0 ; k < N ; k++) { + if ((A_h[k] == memSetVal) && (B_h[k] == memSetVal)) { + validateCount+= 1; + } + } + } + + bool resultAfterAllIterations() { + memDeallocate(); + testResult = (validateCount == (ITER * N)) ? true: false; + return testResult; + } + + void memDeallocate() { + HIP_CHECK(hipFree(A_d)); + free(A_h); + free(B_h); + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + +template +void queueJobsForhipMemsetAsync(T* A_d, T* A_h, T memSetVal, size_t Nbytes, + hipStream_t stream) { + HIPCHECK(hipMemsetAsync(A_d, memSetVal, N, stream)); + HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); +} + +template +void queueJobsForhipMemsetD32Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes, + hipStream_t stream) { + HIPCHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memSetVal, N, stream)); + HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); +} + +template +void queueJobsForhipMemsetD16Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes, + hipStream_t stream) { + HIPCHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memSetVal, N, stream)); + HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); +} + +template +void queueJobsForhipMemsetD8Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes, + hipStream_t stream) { + HIPCHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memSetVal, N, stream)); + HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); +} + +/* Queue hipMemsetAsync jobs on multiple threads and verify they all + * finished on all threads successfully + */ +bool testhipMemsetAsyncWithMultiThread() { + MemSetAsyncMthreadTest obj; + constexpr char memsetval = 0x42; + obj.memAllocate(memsetval); + std::thread t[NUM_THREADS]; + + for (int i = 0 ; i < ITER ; i++) { + for (int k = 0 ; k < NUM_THREADS ; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemsetAsync, obj.A_d, obj.A_h, + obj.memSetVal, obj.Nbytes, obj.stream); + } else { + t[k] = std::thread(queueJobsForhipMemsetAsync, obj.A_d, obj.B_h, + obj.memSetVal, obj.Nbytes, obj.stream); + } + } + + for (int j = 0 ; j < NUM_THREADS ; j++) { + t[j].join(); + } + + HIP_CHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD32AsyncWithMultiThread() { + MemSetAsyncMthreadTest obj; + constexpr int memsetD32val = 0xDEADBEEF; + obj.memAllocate(memsetD32val); + std::thread t[NUM_THREADS]; + + for (int i = 0 ; i < ITER ; i++) { + for (int k = 0 ; k < NUM_THREADS ; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemsetD32Async, obj.A_d, + obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream); + } else { + t[k] = std::thread(queueJobsForhipMemsetD32Async, obj.A_d, + obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream); + } + } + + for (int j = 0 ; j < NUM_THREADS ; j++) { + t[j].join(); + } + + HIP_CHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD16AsyncWithMultiThread() { + MemSetAsyncMthreadTest obj; + constexpr int16_t memsetD16val = 0xDEAD; + obj.memAllocate(memsetD16val); + std::thread t[NUM_THREADS]; + + for (int i = 0 ; i < ITER ; i++) { + for (int k = 0 ; k < NUM_THREADS ; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemsetD16Async, obj.A_d, + obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream); + } else { + t[k] = std::thread(queueJobsForhipMemsetD16Async, obj.A_d, + obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream); + } + } + + for (int j = 0 ; j < NUM_THREADS ; j++) { + t[j].join(); + } + + HIP_CHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD8AsyncWithMultiThread() { + MemSetAsyncMthreadTest obj; + constexpr char memsetD8val = 0xDE; + obj.memAllocate(memsetD8val); + std::thread t[NUM_THREADS]; + + for (int i = 0 ; i < ITER ; i++) { + for (int k = 0 ; k < NUM_THREADS ; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemsetD8Async, obj.A_d, + obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream); + } else { + t[k] = std::thread(queueJobsForhipMemsetD8Async, obj.A_d, + obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream); + } + } + for (int j = 0 ; j < NUM_THREADS ; j++) { + t[j].join(); + } + + HIP_CHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + + +/* + * Test that validates functionality of hipmemsetAsync apis over multi threads + */ +TEST_CASE("Unit_hipMemsetAsync_QueueJobsMultithreaded") { + bool ret; + + SECTION("hipMemsetAsync With MultiThread") { + ret = testhipMemsetAsyncWithMultiThread(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD32Async With MultiThread") { + ret = testhipMemsetD32AsyncWithMultiThread(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD16Async With MultiThread") { + ret = testhipMemsetD16AsyncWithMultiThread(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD8Async With MultiThread") { + ret = testhipMemsetD8AsyncWithMultiThread(); + REQUIRE(ret == true); + } +} diff --git a/catch/unit/memory/hipMemsetInvalidPtr.cc b/catch/unit/memory/hipMemsetInvalidPtr.cc new file mode 100644 index 0000000000..8c87c66ea8 --- /dev/null +++ b/catch/unit/memory/hipMemsetInvalidPtr.cc @@ -0,0 +1,127 @@ +/* + * Copyright (c) 2021 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 : + 1) Test hipMemset apis with invalid pointer and invalid 2D pitch. + 2) Test hipMemsetAsync apis with invalid pointer and invalid 2D pitch. +*/ + + +#include + +#define N 50 +#define MEMSETVAL 0x42 + +/** + * Testcase validates hipMemset apis behavior with + * invalid pointer and invalid 2D pitch value. + */ +TEST_CASE("Unit_hipMemset_InvalidPtrTests") { + hipError_t ret; + constexpr int Nbytes = N*sizeof(char); + char *A_d; + + SECTION("hipMemset with null") { + ret = hipMemset(NULL, MEMSETVAL , Nbytes); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemset with hostptr") { + char *A_h; + A_h = reinterpret_cast(malloc(Nbytes)); + + ret = hipMemset(A_h, MEMSETVAL, Nbytes); + REQUIRE(ret != hipSuccess); + + free(A_h); + } + + SECTION("hipMemsetD32 with null") { + ret = hipMemsetD32(NULL, MEMSETVAL , Nbytes); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD16 with null") { + ret = hipMemsetD16(NULL, MEMSETVAL , Nbytes); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD8 with null") { + ret = hipMemsetD8(NULL, MEMSETVAL , Nbytes); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemset2D with null") { + constexpr size_t NUM_H = 256, NUM_W = 256; + size_t pitch_A; + size_t width = NUM_W * sizeof(char); + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, + width , NUM_H)); + ret = hipMemset2D(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H); + REQUIRE(ret != hipSuccess); + + hipFree(A_d); + } +} + + +/** + * Testcase validates hipMemsetAsync apis behavior with + * invalid pointer and invalid 2D pitch value. + */ +TEST_CASE("Unit_hipMemsetAsync_InvalidPtrTests") { + hipError_t ret; + constexpr int Nbytes = N*sizeof(char); + char *A_d; + + SECTION("hipMemsetAsync with null") { + ret = hipMemsetAsync(NULL, MEMSETVAL, Nbytes , 0); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD32Async with null") { + ret = hipMemsetD32Async(NULL, MEMSETVAL , Nbytes, 0); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD16Async with null") { + ret = hipMemsetD16Async(NULL, MEMSETVAL , Nbytes, 0); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD8Async with null") { + ret = hipMemsetD8Async(NULL, MEMSETVAL , Nbytes, 0); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemset2DAsync with null") { + constexpr size_t NUM_H = 256, NUM_W = 256; + size_t pitch_A; + size_t width = NUM_W * sizeof(char); + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, + width , NUM_H)); + ret = hipMemset2DAsync(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H, 0); + REQUIRE(ret != hipSuccess); + + hipFree(A_d); + } +}