From a1346c114a55f0b8b45d3976fe893d5bfadb43f9 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Sat, 8 Jul 2023 20:54:42 +0530 Subject: [PATCH] SWDEV-388834 - [catch2][dtest] Kernel tests migrated from direct to catch2 (#345) Change-Id: I8d1d7c6d5db018301cd76f2e38b5997ae91c15db [ROCm/hip-tests commit: a461ae2fc986945b5c4667a257999d2387d6c639] --- .../catch/include/hip_test_defgroups.hh | 9 +- .../catch/unit/kernel/CMakeLists.txt | 18 ++ .../catch/unit/kernel/hipPrintfKernel.cc | 67 +++++ .../catch/unit/kernel/hipTestConstant.cc | 76 ++++++ .../unit/kernel/hipTestGlobalVariable.cc | 101 +++++++ .../catch/unit/kernel/hipTestMemKernel.cc | 248 ++++++++++++++++++ .../catch/unit/kernel/inline_asm_vadd.cc | 136 ++++++++++ .../catch/unit/kernel/inline_asm_vmac.cc | 128 +++++++++ .../catch/unit/kernel/launch_bounds.cc | 102 +++++++ .../catch/unit/kernel/printf_common.h | 99 +++++++ 10 files changed, 983 insertions(+), 1 deletion(-) create mode 100644 projects/hip-tests/catch/unit/kernel/hipPrintfKernel.cc create mode 100644 projects/hip-tests/catch/unit/kernel/hipTestConstant.cc create mode 100644 projects/hip-tests/catch/unit/kernel/hipTestGlobalVariable.cc create mode 100644 projects/hip-tests/catch/unit/kernel/hipTestMemKernel.cc create mode 100644 projects/hip-tests/catch/unit/kernel/inline_asm_vadd.cc create mode 100644 projects/hip-tests/catch/unit/kernel/inline_asm_vmac.cc create mode 100644 projects/hip-tests/catch/unit/kernel/launch_bounds.cc create mode 100644 projects/hip-tests/catch/unit/kernel/printf_common.h diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 172f228f84..f4ceedbc1b 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 - 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2021 - 2023 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 @@ -100,3 +100,10 @@ THE SOFTWARE. * This section describes the stream management types & functions of HIP runtime API. * @} */ + +/** + * @defgroup KernelTest Kernel Functions Management + * @{ + * This section describes the various kernel functions invocation. + * @} + */ diff --git a/projects/hip-tests/catch/unit/kernel/CMakeLists.txt b/projects/hip-tests/catch/unit/kernel/CMakeLists.txt index 22fef5b29f..77a48363ed 100644 --- a/projects/hip-tests/catch/unit/kernel/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/kernel/CMakeLists.txt @@ -23,7 +23,25 @@ set(TEST_SRC hipMemFaultStackAllocation.cc hipLaunchBounds.cc hipShflTests.cc + hipTestConstant.cc + hipTestGlobalVariable.cc + hipTestMemKernel.cc + launch_bounds.cc + inline_asm_vadd.cc ) +if(UNIX) + set(TEST_SRC ${TEST_SRC} + hipPrintfKernel.cc) +endif() +string(FIND "${OFFLOAD_ARCH_STR}" "gfx1030" RESULT) +if ("${RESULT}" EQUAL "-1") + set(AMD_TEST_SRC + inline_asm_vmac.cc) +endif() + +if(HIP_PLATFORM MATCHES "amd") + set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC}) +endif() # only for AMD if(HIP_PLATFORM MATCHES "amd") diff --git a/projects/hip-tests/catch/unit/kernel/hipPrintfKernel.cc b/projects/hip-tests/catch/unit/kernel/hipPrintfKernel.cc new file mode 100644 index 0000000000..da81c767fa --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipPrintfKernel.cc @@ -0,0 +1,67 @@ +/* +Copyright (c) 2023 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. +*/ +#include +#include +#include +#include "../kernel/printf_common.h" + +#define HIP_ENABLE_PRINTF + +__global__ void run_printf() { + printf("Hello World\n"); +} +/** +* @addtogroup hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Test case to check printf function via kernel call. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipPrintfKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_kernel_ChkPrintf") { + int device_count = 0; + CaptureStream capture(stdout); + HIP_CHECK(hipGetDeviceCount(&device_count)); + std::string st = "Hello World"; + const char * check = st.c_str(); + for (int i = 0; i < device_count; ++i) { + HIP_CHECK(hipSetDevice(i)); + hipLaunchKernelGGL(run_printf, dim3(1), dim3(1), 0, 0); + HIP_CHECK(hipDeviceSynchronize()); + char* data = new char[st.size()];; + std::ifstream CapturedData = capture.getCapturedData(); + CapturedData.getline(data, st.size()+1); + int result = strcmp(data, check); + REQUIRE(result == 0); + delete [] data; + } +} diff --git a/projects/hip-tests/catch/unit/kernel/hipTestConstant.cc b/projects/hip-tests/catch/unit/kernel/hipTestConstant.cc new file mode 100644 index 0000000000..911457af0f --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipTestConstant.cc @@ -0,0 +1,76 @@ +/* +Copyright (c) 2023 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. +*/ +#include +#include + +#define LEN 512 +#define SIZE 2048 + +__constant__ int Value[LEN]; + +static __global__ void Get(int* Ad) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tid] = Value[tid]; +} +/** +* @addtogroup hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Test case to check const variable via kernel call. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipTestConstant.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_kernel_chkConstantViaKernel") { + int *A, *B, *Ad; + A = new int[LEN]; + B = new int[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = -1 * i; + B[i] = 0; + } + + HIP_CHECK(hipMalloc(reinterpret_cast(&Ad), SIZE)); + + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(Value), A, SIZE, 0, + hipMemcpyHostToDevice)); + hipLaunchKernelGGL(Get, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + REQUIRE(A[i] == B[i]); + } + delete [] A; + delete [] B; + HIP_CHECK(hipFree(Ad)); +} + diff --git a/projects/hip-tests/catch/unit/kernel/hipTestGlobalVariable.cc b/projects/hip-tests/catch/unit/kernel/hipTestGlobalVariable.cc new file mode 100644 index 0000000000..a2d99fa8b7 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipTestGlobalVariable.cc @@ -0,0 +1,101 @@ +/* +Copyright (c) 2023 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. +*/ +#include +#include + +#define LEN 512 +#define SIZE 2048 +/** +* @addtogroup hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Test case to check constant global variable and global array via kernel call. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipTestGlobalVariable.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +__constant__ int ConstantGlobalVar = 123; + +static __global__ void kernel(int* Ad) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tid] = ConstantGlobalVar; +} + +void runTestConstantGlobalVar() { + int *A, *Ad; + A = new int[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + HIP_CHECK(hipMalloc(reinterpret_cast(&Ad), SIZE)); + hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + REQUIRE(123 == A[i]); + } + delete [] A; + HIP_CHECK(hipFree(Ad)); +} + +__device__ int GlobalArray[LEN]; +static __global__ void kernelWrite() { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + GlobalArray[tid] = tid; +} +static __global__ void kernelRead(int* Ad) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tid] = GlobalArray[tid]; +} + +void runTestGlobalArray() { + int *A, *Ad; + A = new int[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + + HIP_CHECK(hipMalloc(reinterpret_cast(&Ad), SIZE)); + hipLaunchKernelGGL(kernelWrite, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0); + hipLaunchKernelGGL(kernelRead, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + REQUIRE(i == A[i]); + } + delete [] A; + HIP_CHECK(hipFree(Ad)); +} + +TEST_CASE("Unit_kernel_chkGlobalArrAndGlobalVaribleViaKernelFn") { + runTestConstantGlobalVar(); + runTestGlobalArray(); +} diff --git a/projects/hip-tests/catch/unit/kernel/hipTestMemKernel.cc b/projects/hip-tests/catch/unit/kernel/hipTestMemKernel.cc new file mode 100644 index 0000000000..d97a5698e0 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipTestMemKernel.cc @@ -0,0 +1,248 @@ +/* +Copyright (c) 2023 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. +*/ +#include +#include + +#define LEN8 8 * 4 +#define LEN9 9 * 4 +#define LEN10 10 * 4 +#define LEN11 11 * 4 +#define LEN12 12 * 4 + +__global__ void MemCpy8(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 8, In + tid * 8, 8); +} + +__global__ void MemCpy9(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 9, In + tid * 9, 9); +} + +__global__ void MemCpy10(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 10, In + tid * 10, 10); +} + +__global__ void MemCpy11(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 11, In + tid * 11, 11); +} + +__global__ void MemCpy12(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 12, In + tid * 12, 12); +} + +__global__ void MemSet8(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 8, 1, 8); +} + +__global__ void MemSet9(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 9, 1, 9); +} + +__global__ void MemSet10(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 10, 1, 10); +} + +__global__ void MemSet11(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 11, 1, 11); +} + +__global__ void MemSet12(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 12, 1, 12); +} +/** +* @addtogroup hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Test case to check memcpy and memset via kernel call. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipTestMemKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_kernel_MemoryOperationsViaKernels") { + uint8_t *A, *Ad, *B, *Bd, *C, *Cd; + A = new uint8_t[LEN8]; + B = new uint8_t[LEN8]; + C = new uint8_t[LEN8]; + for (uint32_t i = 0; i < LEN8; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN8)); + HIP_CHECK(hipMalloc(&Bd, LEN8)); + HIP_CHECK(hipMalloc(&Cd, LEN8)); + HIP_CHECK(hipMemcpy(Ad, A, LEN8, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy8, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet8, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy(B, Bd, LEN8, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(C, Cd, LEN8, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN8; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + + SECTION("MemCpySet1") { + A = new uint8_t[LEN9]; + B = new uint8_t[LEN9]; + C = new uint8_t[LEN9]; + for (uint32_t i = 0; i < LEN9; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN9)); + HIP_CHECK(hipMalloc(&Bd, LEN9)); + HIP_CHECK(hipMalloc(&Cd, LEN9)); + HIP_CHECK(hipMemcpy(Ad, A, LEN9, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy9, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet9, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy(B, Bd, LEN9, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(C, Cd, LEN9, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN9; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + } + + SECTION("MemCpySet2") { + A = new uint8_t[LEN10]; + B = new uint8_t[LEN10]; + C = new uint8_t[LEN10]; + for (uint32_t i = 0; i < LEN10; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN10)); + HIP_CHECK(hipMalloc(&Bd, LEN10)); + HIP_CHECK(hipMalloc(&Cd, LEN10)); + HIP_CHECK(hipMemcpy(Ad, A, LEN10, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy10, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet10, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy(B, Bd, LEN10, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(C, Cd, LEN10, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN10; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + } + + SECTION("MemCpySet3") { + A = new uint8_t[LEN11]; + B = new uint8_t[LEN11]; + C = new uint8_t[LEN11]; + for (uint32_t i = 0; i < LEN11; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN11)); + HIP_CHECK(hipMalloc(&Bd, LEN11)); + HIP_CHECK(hipMalloc(&Cd, LEN11)); + HIP_CHECK(hipMemcpy(Ad, A, LEN11, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy11, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet11, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy(B, Bd, LEN11, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(C, Cd, LEN11, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN11; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + } + + SECTION("MemCpySet4") { + A = new uint8_t[LEN12]; + B = new uint8_t[LEN12]; + C = new uint8_t[LEN12]; + for (uint32_t i = 0; i < LEN12; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN12)); + HIP_CHECK(hipMalloc(&Bd, LEN12)); + HIP_CHECK(hipMalloc(&Cd, LEN12)); + HIP_CHECK(hipMemcpy(Ad, A, LEN12, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy12, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet12, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy(B, Bd, LEN12, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(C, Cd, LEN12, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN12; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + } +} diff --git a/projects/hip-tests/catch/unit/kernel/inline_asm_vadd.cc b/projects/hip-tests/catch/unit/kernel/inline_asm_vadd.cc new file mode 100644 index 0000000000..29a4d8fe29 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/inline_asm_vadd.cc @@ -0,0 +1,136 @@ +/* +Copyright (c) 2023 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. +*/ +#include +#include + +#define NUM 1024 +#define THREADS_PER_BLOCK_X 4 + +// Device (Kernel) function, it must be void +__global__ void vadd_asm(float* out, float* in) { + int i = blockDim.x * blockIdx.x + threadIdx.x; + + #ifdef __HIP_PLATFORM_NVIDIA__ + asm volatile("add.f32 %0,%1,%2;" : "=f"(out[i]) : "f"(in[i]), "f"(out[i])); + #endif + + #ifdef __HIP_PLATFORM_AMD__ + asm volatile("v_add_f32_e32 %0, %1, %2" : "=v"(out[i]) : "v"(in[i]), + "v"(out[i])); + #endif +} + +// CPU implementation of Vector Result +void addCPUReference(float* output, float* input) { + for (unsigned int j = 0; j < NUM; j++) { + output[j] = input[j] + output[j]; + } +} +/** +* @addtogroup hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Test case to check inline asm vadd instruction via kernel call. + + * Test source + * ------------------------ + * - catch/unit/kernel/inline_asm_vadd.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_kernel_inline_asm_vadd_Functional") { + float* VectorA; + float* ResultVector; + float* VectorB; + + float* gpuVector; + float* gpuResultVector; + + int i; + int errors; + + VectorA = reinterpret_cast(malloc(NUM * sizeof(float))); + ResultVector = reinterpret_cast(malloc(NUM * sizeof(float))); + VectorB = reinterpret_cast(malloc(NUM * sizeof(float))); + + // initialize the input data + for (i = 0; i < NUM; i++) { + VectorA[i] = static_cast(i * 10.0f); + VectorB[i] = static_cast(i * 30.0f); + } + + // allocate the memory on the device side + HIP_CHECK(hipMalloc(reinterpret_cast(&gpuVector), + NUM * sizeof(float))); + HIP_CHECK(hipMalloc(reinterpret_cast(&gpuResultVector), + NUM * sizeof(float))); + + // Memory transfer from host to device + HIP_CHECK(hipMemcpy(gpuVector, VectorA, NUM * sizeof(float), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(gpuResultVector, VectorB, NUM * sizeof(float), + hipMemcpyHostToDevice)); + + // Lauching kernel from host + hipLaunchKernelGGL(vadd_asm, dim3(NUM / THREADS_PER_BLOCK_X), + dim3(THREADS_PER_BLOCK_X), 0, 0, + gpuResultVector, gpuVector); + + // Memory transfer from device to host + HIP_CHECK(hipMemcpy(ResultVector, gpuResultVector, NUM * sizeof(float), + hipMemcpyDeviceToHost)); + + // CPU Result computation + addCPUReference(VectorB, VectorA); + + // verify the results + errors = 0; + double eps = 1.0E-3; + for (i = 0; i < NUM; i++) { + if (std::abs(ResultVector[i] - VectorB[i]) > eps) { + errors++; + } + } + if (errors != 0) { + REQUIRE(false); + } else { + REQUIRE(true); + } + + // free the resources on device side + HIP_CHECK(hipFree(gpuVector)); + HIP_CHECK(hipFree(gpuResultVector)); + + HIP_CHECK(hipDeviceReset()); + + // free the resources on host side + free(VectorA); + free(ResultVector); + free(VectorB); +} diff --git a/projects/hip-tests/catch/unit/kernel/inline_asm_vmac.cc b/projects/hip-tests/catch/unit/kernel/inline_asm_vmac.cc new file mode 100644 index 0000000000..e9fbebda37 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/inline_asm_vmac.cc @@ -0,0 +1,128 @@ +/* +Copyright (c) 2023 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. +*/ +#include +#include + +#define NUM 1024 +#define THREADS_PER_BLOCK_X 4 + +// Device (Kernel) function, it must be void +__global__ void vmac_asm(float* out, float* in, float a) { + int i = blockDim.x * blockIdx.x + threadIdx.x; + asm volatile("v_mac_f32_e32 %0, %2, %3" : "=v"(out[i]) : + "0"(out[i]), "v"(a), "v"(in[i])); +} + +// CPU implementation of saxpy +void addCPUReference(float* output, float* input, float a) { + for (unsigned int j = 0; j < NUM; j++) { + output[j] = a * input[j] + output[j]; + } +} +/** +* @addtogroup hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Test case to check inline asm vmac instruction via kernel call. + + * Test source + * ------------------------ + * - catch/unit/kernel/inline_asm_vmac.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_kernel_inline_asm_vmac_Functional") { + float* VectorA; + float* ResultVector; + float* VectorB; + + float* gpuVector; + float* gpuResultVector; + + const float a = 10.0f; + int i; + int errors; + + VectorA = reinterpret_cast(malloc(NUM * sizeof(float))); + ResultVector = reinterpret_cast(malloc(NUM * sizeof(float))); + VectorB = reinterpret_cast(malloc(NUM * sizeof(float))); + + // initialize the input data + for (i = 0; i < NUM; i++) { + VectorA[i] = static_cast(i * 10.0f); + VectorB[i] = static_cast(i * 30.0f); + } + + // allocate the memory on the device side + HIP_CHECK(hipMalloc(reinterpret_cast(&gpuVector), + NUM * sizeof(float))); + HIP_CHECK(hipMalloc(reinterpret_cast(&gpuResultVector), + NUM * sizeof(float))); + + // Memory transfer from host to device + HIP_CHECK(hipMemcpy(gpuVector, VectorA, NUM * sizeof(float), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(gpuResultVector, VectorB, NUM * sizeof(float), + hipMemcpyHostToDevice)); + + // Lauching kernel from host + hipLaunchKernelGGL(vmac_asm, dim3(NUM / THREADS_PER_BLOCK_X), + dim3(THREADS_PER_BLOCK_X), 0, 0, + gpuResultVector, gpuVector, a); + + // Memory transfer from device to host + HIP_CHECK(hipMemcpy(ResultVector, gpuResultVector, NUM * sizeof(float), + hipMemcpyDeviceToHost)); + + // CPU Result computation + addCPUReference(VectorB, VectorA, a); + + // verify the results + errors = 0; + double eps = 1.0E-3; + for (i = 0; i < NUM; i++) { + if (std::abs(ResultVector[i] - VectorB[i]) > eps) { + errors++; + } + } + if (errors != 0) { + REQUIRE(false); + } else { + REQUIRE(true); + } + + // free the resources on device side + HIP_CHECK(hipFree(gpuVector)); + HIP_CHECK(hipFree(gpuResultVector)); + HIP_CHECK(hipDeviceReset()); + + // free the resources on host side + free(VectorA); + free(ResultVector); + free(VectorB); +} diff --git a/projects/hip-tests/catch/unit/kernel/launch_bounds.cc b/projects/hip-tests/catch/unit/kernel/launch_bounds.cc new file mode 100644 index 0000000000..59b1132898 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/launch_bounds.cc @@ -0,0 +1,102 @@ +/* +Copyright (c) 2023 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. +*/ +#include +#include + +constexpr size_t N = 1024; +int p_blockSize = 256; + +__global__ void +__launch_bounds__(256, 2) + myKern(int* C, const int* A, int N) { + int tid = (blockIdx.x * blockDim.x + threadIdx.x); + + if (tid < N) { + C[tid] = A[tid]; + } +} +/** +* @addtogroup hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Test case to check launch bounds via kernel call. + + * Test source + * ------------------------ + * - catch/unit/kernel/launch_bounds.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_kernel_LaunchBounds_Functional") { + size_t Nbytes = N * sizeof(int); + int *A_d, *C_d, *A_h, *C_h; + HIPCHECK(hipMalloc(&A_d, Nbytes)); + HIPCHECK(hipMalloc(&C_d, Nbytes)); + + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + + for (int i = 0; i < N; i++) { + A_h[i] = i * 10; + C_h[i] = 0x0; + } + int blocks = N / p_blockSize; + + HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK(hipGetLastError()); + hipLaunchKernelGGL(myKern, dim3(blocks), dim3(p_blockSize), 0, + 0, C_d, A_d, N); + +#ifdef __HIP_PLATFORM_NVIDIA__ + cudaFuncAttributes attrib; + cudaFuncGetAttributes(&attrib, myKern); + printf("binaryVersion = %d\n", attrib.binaryVersion); + printf("cacheModeCA = %d\n", attrib.cacheModeCA); + printf("constSizeBytes = %zu\n", attrib.constSizeBytes); + printf("localSizeBytes = %zud\n", attrib.localSizeBytes); + printf("maxThreadsPerBlock = %d\n", attrib.maxThreadsPerBlock); + printf("numRegs = %d\n", attrib.numRegs); + printf("ptxVersion = %d\n", attrib.ptxVersion); + printf("sharedSizeBytes = %zud\n", attrib.sharedSizeBytes); +#endif + + HIPCHECK(hipDeviceSynchronize()); + HIPCHECK(hipGetLastError()); + HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipDeviceSynchronize()); + + for (int i = 0; i < N; i++) { + int goldVal = i * 10; + REQUIRE(C_h[i] == goldVal); + } + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); + free(A_h); + free(C_h); +} diff --git a/projects/hip-tests/catch/unit/kernel/printf_common.h b/projects/hip-tests/catch/unit/kernel/printf_common.h new file mode 100644 index 0000000000..36c512d2a2 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/printf_common.h @@ -0,0 +1,99 @@ +/* +Copyright (c) 2023 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. +*/ + +#ifndef _STRESSTEST_PRINTF_COMMON_H_ +#define _STRESSTEST_PRINTF_COMMON_H_ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +struct CaptureStream { + int saved_fd; + int orig_fd; + int temp_fd; + + char tempname[13] = "mytestXXXXXX"; + + explicit CaptureStream(FILE *original) { + orig_fd = fileno(original); + saved_fd = dup(orig_fd); + + if ((temp_fd = mkstemp(tempname)) == -1) { + error(0, errno, "Error"); + assert(false); + } + + fflush(nullptr); + if (dup2(temp_fd, orig_fd) == -1) { + error(0, errno, "Error"); + assert(false); + } + if (close(temp_fd) != 0) { + error(0, errno, "Error"); + assert(false); + } + } + + void restoreStream() { + if (saved_fd == -1) + return; + fflush(nullptr); + if (dup2(saved_fd, orig_fd) == -1) { + error(0, errno, "Error"); + assert(false); + } + if (close(saved_fd) != 0) { + error(0, errno, "Error"); + assert(false); + } + saved_fd = -1; + } + + const char *getTempFilename() { + return (const char*)tempname; + } + + std::ifstream getCapturedData() { + restoreStream(); + std::ifstream temp(tempname); + return temp; + } + + ~CaptureStream() { + restoreStream(); + if (remove(tempname) != 0) { + error(0, errno, "Error"); + assert(false); + } + } +}; + +#endif // _STRESSTEST_PRINTF_COMMON_H_