From 31e62c5852010c60941f63cbb75cb2bad01ea7ad Mon Sep 17 00:00:00 2001 From: Sudheer Kumar Date: Mon, 13 Jul 2020 16:55:41 +0530 Subject: [PATCH] [dtest] Enhanced tests to cover multiple hipmemset and hipmemsetAsync apis SWDEV-238517 for enhancing hip unit tests Change-Id: Iba2c419a4487955f34b4f19abe174ef427d289d8 [ROCm/clr commit: 7d77902198b439b6f20bafb29711b4b07356fd17] --- .../src/runtimeApi/memory/hipMemset2D.cpp | 303 ++++++++++++------ .../hipMemset2DAsyncMultiThreadAndKernel.cpp | 173 ++++++++++ .../memory/hipMemsetAsyncAndKernel.cpp | 191 +++++++++++ .../memory/hipMemsetAsyncMultiThread.cpp | 247 ++++++++++++++ .../runtimeApi/memory/hipMemsetInvalidPtr.cpp | 97 ++++++ 5 files changed, 917 insertions(+), 94 deletions(-) create mode 100644 projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2DAsyncMultiThreadAndKernel.cpp create mode 100644 projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncAndKernel.cpp create mode 100644 projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncMultiThread.cpp create mode 100644 projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetInvalidPtr.cpp diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2D.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2D.cpp index e7fe932b3d..0e54c4ec67 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2D.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2D.cpp @@ -1,119 +1,234 @@ /* -Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved. + * 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. + */ -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. +// Test for hipMemset2D functionality for different width and height values /* HIT_START - * BUILD: %t %s ../../test_common.cpp - * TEST: %t + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST_NAMED: %t hipMemset2D-basic + * TEST_NAMED: %t hipMemset2D-dim1 --width2D 10 --height2D 10 --memsetWidth 4 --memsetHeight 4 + * TEST_NAMED: %t hipMemset2D-dim2 --width2D 100 --height2D 100 --memsetWidth 20 --memsetHeight 40 + * TEST_NAMED: %t hipMemset2D-dim3 --width2D 256 --height2D 256 --memsetWidth 39 --memsetHeight 19 + * TEST_NAMED: %t hipMemset2D-zeroH --width2D 100 --height2D 100 --memsetWidth 20 --memsetHeight 0 + * TEST_NAMED: %t hipMemset2D-zeroW --width2D 100 --height2D 100 --memsetWidth 0 --memsetHeight 20 + * TEST_NAMED: %t hipMemset2D-zeroW*H --width2D 100 --height2D 100 --memsetWidth 0 --memsetHeight 0 * 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; +// Check hipMemset2D functionality +bool testhipMemset2D(int memsetval, int p_gpuDevice) { + bool testResult = true; + 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; + HIPCHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width , + numH)); + A_h = reinterpret_cast(malloc(sizeElements)); + HIPASSERT(A_h != NULL); - 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(A_h[i]), static_cast(memsetval)); + break; } - HIPCHECK ( hipMemset2D(A_d, pitch_A, memsetval, numW, numH) ); - HIPCHECK ( hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost)); + } - for (int i=0; i(&A_d), &pitch_A, + width , numH)); + A_h = reinterpret_cast(malloc(sizeElements)); + HIPASSERT(A_h != NULL); - printf ("testhipMemset2DAsync memsetval=%2x device=%d\n", memsetval, p_gpuDevice); - char *A_d; - char *A_h; - bool testResult = true; + for (size_t i=0; i < elements; i++) { + A_h[i] = 1; + } - HIPCHECK (hipMallocPitch((void**)&A_d, &pitch_A, width , numH)); - A_h = (char*)malloc(sizeElements); - HIPASSERT(A_h != NULL); - for (size_t i=0; i(A_h[i]), static_cast(memsetval)); + break; } - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - HIPCHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream) ); - HIPCHECK(hipStreamSynchronize(stream)); - HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost)); + } - for (int i=0; i= argc || !HipTest::parseInt(argv[i], &width2D)) { + failed("Bad width2D argument"); + } + } else if (!strcmp(arg, "--height2D")) { + if (++i >= argc || !HipTest::parseInt(argv[i], &height2D)) { + failed("Bad height2D argument"); + } + } else if (!strcmp(arg, "--memsetWidth")) { + if (++i >= argc || !HipTest::parseInt(argv[i], &memsetWidth)) { + failed("Bad memsetWidth argument"); + } + } else if (!strcmp(arg, "--memsetHeight")) { + if (++i >= argc || !HipTest::parseInt(argv[i], &memsetHeight)) { + failed("Bad memsetHeight argument"); + } + } else { + failed("Bad argument"); + } + } + return i; +} + +// Memset random dimensions +bool testMemset2DPartial(int memsetval, int p_gpuDevice) { + bool testResult = true; + size_t NUM_H = height2D; + size_t NUM_W = width2D; + size_t Nbytes = N*sizeof(char); + size_t pitch_A; + size_t width = NUM_W * sizeof(char); + size_t sizeElements = width * NUM_H; + size_t elements = NUM_W * NUM_H; + char *A_d; + char *A_h; + printf("testhipMemset2DPartial memsetval=%2x device=%d\n", memsetval, + p_gpuDevice); + + HIPCHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, + width, NUM_H)); + hipError_t e; + int index; + + A_h = reinterpret_cast(malloc(sizeElements)); + HIPASSERT(A_h != NULL); + + for (index = 0; index < sizeElements; index++) { + A_h[0] = 'c'; + } + + printf("2D Dimension: %zuX%zu, MemsetWidth:%d, memsetHeight:%d\n", + NUM_W, NUM_H, memsetWidth, memsetHeight); + e = hipMemset2D(A_d, pitch_A, memsetval, memsetWidth, memsetHeight); + HIPASSERT(e == hipSuccess); + + HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, NUM_W, NUM_H, + hipMemcpyDeviceToHost)); + + for (int row = 0; row < memsetHeight; row++) { + for (int column = 0; column < memsetWidth; column++) { + if (A_h[(row * width) + column] != memsetval) { + printf("A_h[%d][%d] did not match %d", row, column, memsetval); + testResult = false; + } + } + } + hipFree(A_d); + free(A_h); + return testResult; +} + +int main(int argc, char *argv[]) { + int extraArgs = 0; + bool testResult = true; + + HIPCHECK(hipSetDevice(p_gpuDevice)); + extraArgs = HipTest::parseStandardArguments(argc, argv, false); + parseExtraArguments(extraArgs, argv); + + if (extraArgs == 1) { testResult &= testhipMemset2D(memsetval, p_gpuDevice); - testResult &= testhipMemset2DAsync(memsetval, p_gpuDevice); - if(testResult){ - passed(); + if (!(testResult)) { + printf("hipMemset2D failed\n"); } + testResult &= testhipMemset2DAsync(memsetval, p_gpuDevice); + if (!(testResult)) { + printf("hipMemset2DAsync failed\n"); + } + } else if (extraArgs == 9) { + testResult &= testMemset2DPartial(memsetval, p_gpuDevice); + if (!(testResult)) { + printf("hipMemset2D at random dimensions failed\n"); + } + } else { + failed("Wrong Arguments for test\n"); + } + + if (testResult) { + passed(); + } else { + failed("one or more hipMemset2D tests failed"); + } } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2DAsyncMultiThreadAndKernel.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2DAsyncMultiThreadAndKernel.cpp new file mode 100644 index 0000000000..6bad4c054c --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2DAsyncMultiThreadAndKernel.cpp @@ -0,0 +1,173 @@ +/* + * 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. +*/ + +// +// Test to verify +// a) Order of execution of device kernel and hipMemset2DAsync api +// b) hipMemSet2DAsync execution in multiple threads +// + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#define NUM_THREADS 1000 +#define ITER 100 +#define NUM_H 256 +#define NUM_W 256 + +unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); +hipStream_t stream; + +bool testResult = true; +char *A_d, *A_h, *B_d, *B_h, *C_d; +int validateCount; + +size_t pitch_A, pitch_B, pitch_C; +size_t width = NUM_W * sizeof(char); +size_t sizeElements = width * NUM_H; +size_t elements = NUM_W * NUM_H; + +/* + * Square each element in the array B and write to array C. + */ + +__global__ void +vector_square(char* B_d, char* C_d, size_t elements) { + for (int i=0 ; i < elements ; i++) { + C_d[i] = B_d[i] * B_d[i]; + } +} + +void memAllocate() { + HIPCHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, NUM_H)); + HIPCHECK(hipMallocPitch(reinterpret_cast(&B_d), &pitch_B, width, NUM_H)); + A_h = reinterpret_cast(malloc(sizeElements)); + HIPASSERT(A_h != NULL); + B_h = reinterpret_cast(malloc(sizeElements)); + HIPASSERT(B_h != NULL); + HIPCHECK(hipMallocPitch(reinterpret_cast(&C_d), &pitch_C, width, NUM_H)); + + for (int i = 0 ; i < elements ; i++) { + B_h[i] = i; + } + HIPCHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H, + hipMemcpyHostToDevice)); + HIPCHECK(hipStreamCreate(&stream)); +} + +void memDeallocate() { + HIPCHECK(hipFree(A_d)); HIPCHECK(hipFree(B_d)); HIPCHECK(hipFree(C_d)); + free(A_h); free(B_h); + HIPCHECK(hipStreamDestroy(stream)); +} + +void queueJobsForhipMemset2DAsync(char* A_d, char* A_h, size_t pitch, + size_t width) { + HIPCHECK(hipMemset2DAsync(A_d, pitch, memsetval, NUM_W, NUM_H, stream)); + HIPCHECK(hipMemcpy2DAsync(A_h, width, A_d, pitch, NUM_W, NUM_H, + hipMemcpyDeviceToHost, stream)); +} + +bool testhipMemset2DAsyncWithKernel() { + validateCount = 0; + memAllocate(); + printf("info: Launching vector_square kernel and hipMemset2DAsync " + "simultaneously\n"); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, + stream, B_d, C_d, elements); + + HIPCHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream)); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, + hipMemcpyDeviceToHost)); + + for (int p = 0 ; p < elements ; p++) { + if (A_h[p] == memsetval) { + validateCount+= 1; + } + } + } + + testResult = (validateCount == (ITER * elements)) ? true : false; + memDeallocate(); + return testResult; +} + +bool testhipMemset2DAsyncMultiThread() { + validateCount = 0; + std::thread t[NUM_THREADS]; + + memAllocate(); + + printf("info: Queueing up hipMemset2DAsync jobs over multiple threads\n"); + for (int i = 0 ; i < ITER ; i++) { + for (int k = 0 ; k < NUM_THREADS ; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, A_h, pitch_A, + width); + } else { + t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, B_h, pitch_A, + width); + } + } + for (int j = 0 ; j < NUM_THREADS ; j++) { + t[j].join(); + } + + HIPCHECK(hipStreamSynchronize(stream)); + for (int k = 0 ; k < elements ; k++) { + if ((A_h[k] == memsetval) && (B_h[k] == memsetval)) { + validateCount+= 1; + } + } + } + memDeallocate(); + testResult = (validateCount == (ITER * elements)) ? true : false; + return testResult; +} + +int main() { + bool testResult = true; + + testResult &= testhipMemset2DAsyncWithKernel(); + if (testResult) { + printf("Kernel and hipMemset2DAsync executed in correct order!\n"); + } else { + printf("Kernel and hipMemset2DAsync order of execution failed\n"); + } + + testResult &= testhipMemset2DAsyncMultiThread(); + if (testResult) { + printf("hipMemset2DAsync jobs on all threads finished successfully!\n"); + passed(); + } else { + printf("hipMemset2DAsync failed in multi thread scenario\n"); + } + + if (testResult) { + passed(); + } else { + failed("One or more tests failed\n"); + } +} diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncAndKernel.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncAndKernel.cpp new file mode 100644 index 0000000000..564b804a9f --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncAndKernel.cpp @@ -0,0 +1,191 @@ +/* + * 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. +*/ + +/* + * Test for checking order of execution of device kernel and + * hipMemsetAsync apis on all gpus + */ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#define ITER 10 +#define N 1024 * 1024 + +unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + +template +__global__ void vector_square(T* B_d, T* C_d, size_t M) { + for (int i=0 ; i < M ; i++) { + C_d[i] = B_d[i] * B_d[i]; + } +} + +template +class MemSetTest { + 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)); + HIPASSERT(A_h != NULL); + HIPCHECK(hipMalloc(&B_d , Nbytes)); + B_h = reinterpret_cast(malloc(Nbytes)); + HIPASSERT(B_h != NULL); + HIPCHECK(hipMalloc(&C_d , Nbytes)); + + for (int i = 0 ; i < N ; i++) { + B_h[i] = i; + } + HIPCHECK(hipMemcpy(B_d , B_h , Nbytes , hipMemcpyHostToDevice)); + HIPCHECK(hipStreamCreate(&stream)); + } + + void memDeallocate() { + HIPCHECK(hipFree(B_d)); HIPCHECK(hipFree(C_d)); + free(B_h); free(A_h); + HIPCHECK(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; + } +}; + +bool testhipMemsetAsyncWithKernel() { + MemSetTest obj; + obj.memAllocate(memsetval); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, + obj.stream, obj.B_d, obj.C_d, N); + HIPCHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream)); + HIPCHECK(hipStreamSynchronize(obj.stream)); + HIPCHECK(hipMemcpy(obj.A_h , obj.C_d , obj.Nbytes , hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD32AsyncWithKernel() { + MemSetTest obj; + obj.memAllocate(memsetD32val); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, + obj.stream, obj.B_d, obj.C_d, N); + HIPCHECK(hipMemsetD32Async(obj.C_d , obj.memSetVal , N , obj.stream)); + HIPCHECK(hipStreamSynchronize(obj.stream)); + HIPCHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD16AsyncWithKernel() { + MemSetTest obj; + obj.memAllocate(memsetD16val); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, + obj.stream, obj.B_d, obj.C_d, N); + HIPCHECK(hipMemsetD16Async(obj.C_d , obj.memSetVal , N , obj.stream)); + HIPCHECK(hipStreamSynchronize(obj.stream)); + HIPCHECK(hipMemcpy(obj.A_h , obj.C_d, obj.Nbytes , hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD8AsyncWithKernel() { + MemSetTest obj; + obj.memAllocate(memsetD8val); + for (int k = 0; k < ITER; k++) { + hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, + obj.stream, obj.B_d, obj.C_d, N); + HIPCHECK(hipMemsetD8Async(obj.C_d, obj.memSetVal, N, obj.stream)); + HIPCHECK(hipStreamSynchronize(obj.stream)); + HIPCHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +int main() { + bool testResult = true; + int numDevices = 0; + HIPCHECK(hipGetDeviceCount(&numDevices)); + printf("total number of gpus in the system: %d\n", numDevices); + + for (int i = 0; i < numDevices; i++) { + HIPCHECK(hipSetDevice(i)); + printf("test running on gpu %d\n", i); + + testResult &= testhipMemsetAsyncWithKernel(); + if (!(testResult)) { + printf("Mismatch in order of execution of hipMemsetAsync and kernel\n"); + } + + testResult &= testhipMemsetD32AsyncWithKernel(); + if (!(testResult)) { + printf("Mismatch in order of execution of hipMemsetD32Async and kernel\n"); + } + + testResult &= testhipMemsetD16AsyncWithKernel(); + if (!(testResult)) { + printf("Mismatch in order of execution of hipMemsetD16Async and kernel\n"); + } + + testResult &= testhipMemsetD8AsyncWithKernel(); + if (!(testResult)) { + printf("Mismatch in order of execution of hipMemsetD8Async and kernel\n"); + } + } + + if (testResult) { + printf("Execution order of Kernel and hipMemsetAsync apis on " + "all gpus is correct!\n"); + passed(); + } else { + failed("One or more hipMemsetAsync tests failed\n"); + } +} diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncMultiThread.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncMultiThread.cpp new file mode 100644 index 0000000000..b4ae8658f7 --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncMultiThread.cpp @@ -0,0 +1,247 @@ +/* + * 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. +*/ + +/* + * Test that validates functionality of hipmemsetAsync apis over multi threads + */ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#define NUM_THREADS 50 +#define ITER 50 + +unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + +template +class MemSetTest { + 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)); + HIPASSERT(A_h != NULL); + + HIPCHECK(hipMalloc(&A_d, Nbytes)); + B_h = reinterpret_cast(malloc(Nbytes)); + HIPASSERT(B_h != NULL); + + HIPCHECK(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() { + HIPCHECK(hipFree(A_d)); + free(A_h); + free(B_h); + HIPCHECK(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(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(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(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() { + MemSetTest obj; + 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(); + } + + HIPCHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD32AsyncWithMultiThread() { + MemSetTest obj; + 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(); + } + + HIPCHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD16AsyncWithMultiThread() { + MemSetTest obj; + 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(); + } + + HIPCHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD8AsyncWithMultiThread() { + MemSetTest obj; + 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(); + } + + HIPCHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +int main() { + bool testResult = true; + printf("Queueing up hipMemSetAsync jobs on multiple threads" + "and checking results\n"); + + testResult &= testhipMemsetAsyncWithMultiThread(); + if (!(testResult)) { + printf("Thread execution did not complete for hipMemsetAsync\n"); + } + + testResult &= testhipMemsetD32AsyncWithMultiThread(); + if (!(testResult)) { + printf("Thread execution did not complete for hipMemsetD32Async\n"); + } + + testResult &= testhipMemsetD16AsyncWithMultiThread(); + if (!(testResult)) { + printf("Thread execution did not complete for hipMemsetD16Async\n"); + } + testResult &= testhipMemsetD8AsyncWithMultiThread(); + if (!(testResult)) { + printf("Thread execution did not complete for hipMemsetD8Async\n"); + } + + if (testResult) { + printf("All threads ran successfully for all hipMemsetAsync apis\n"); + passed(); + } else { + failed("One or more tests failed\n"); + } +} diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetInvalidPtr.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetInvalidPtr.cpp new file mode 100644 index 0000000000..6945c73a43 --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetInvalidPtr.cpp @@ -0,0 +1,97 @@ +/* + * 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. +*/ + +// * To test invalid pointer to hipMemset* apis + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#define N 50 +#define MEMSETVAL 0x42 +#define NUM_H 256 +#define NUM_W 256 + +int main() { + size_t Nbytes = N*sizeof(char); + size_t pitch_A; + size_t width = NUM_W * sizeof(char); + size_t sizeElements = width * NUM_H; + size_t elements = NUM_W * NUM_H; + char *A_d; + + HIPCHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width , NUM_H)); + + hipError_t e; + + e = hipMemset(NULL , MEMSETVAL , Nbytes); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemsetD32(NULL , MEMSETVAL , Nbytes); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemsetD16(NULL , MEMSETVAL , Nbytes); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemsetD8(NULL , MEMSETVAL , Nbytes); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemsetAsync(NULL , MEMSETVAL , Nbytes , 0); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemsetD32Async(NULL , MEMSETVAL , Nbytes, 0); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemsetD16Async(NULL , MEMSETVAL , Nbytes, 0); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemsetD8Async(NULL , MEMSETVAL , Nbytes, 0); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemset2D(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemset2DAsync(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H, 0); + HIPASSERT(e == hipErrorInvalidValue); + + /* Passing host pointer to hipMemset.Ticket SWDEV-243206 is open for this. + * Disabling this test until the ticket is closed + * + char *A_h; + A_h = (char*)malloc(Nbytes); + e = hipMemset(A_h, MEMSETVAL , Nbytes); + HIPASSERT(e == hipErrorInvalidValue); + */ + + /* Passing invalid pitch to hipMemset2D.Ticket SWDEV-243104 is open for this. + * Disabling this test until the ticket is closed + * + e = hipMemset2D(A_d, 0, MEMSETVAL, NUM_W, NUM_H); + HIPASSERT(e == hipErrorInvalidValue); + + e = hipMemset2DAsync(A_d, 0, MEMSETVAL, NUM_W, NUM_H,0); + HIPASSERT(e == hipErrorInvalidValue); + */ + + hipFree(A_d); + passed(); +}