From 7ff2cb905bc92de41e278d33aa8adfb7be127b17 Mon Sep 17 00:00:00 2001 From: Vladana Stojiljkovic Date: Wed, 28 Feb 2024 17:04:18 +0100 Subject: [PATCH] SWDEV-441595 - Add hipSetupArgument test Change-Id: I54aa28416986fd000ce0d238074d291d64c861cd [ROCm/hip-tests commit: 43ce72334c7c9589658661d546ae312adc79b444] --- .../catch/unit/kernel/CMakeLists.txt | 1 + .../catch/unit/kernel/hipSetupArgument.cc | 90 +++++++++++++++++++ 2 files changed, 91 insertions(+) create mode 100644 projects/hip-tests/catch/unit/kernel/hipSetupArgument.cc diff --git a/projects/hip-tests/catch/unit/kernel/CMakeLists.txt b/projects/hip-tests/catch/unit/kernel/CMakeLists.txt index 3ac8bb41f2..a32d48f4f2 100644 --- a/projects/hip-tests/catch/unit/kernel/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/kernel/CMakeLists.txt @@ -43,6 +43,7 @@ endif() if(HIP_PLATFORM MATCHES "amd") set(AMD_SRC hipExtLaunchKernelGGL.cc + hipSetupArgument.cc ) set(TEST_SRC ${TEST_SRC} ${AMD_SRC}) endif() diff --git a/projects/hip-tests/catch/unit/kernel/hipSetupArgument.cc b/projects/hip-tests/catch/unit/kernel/hipSetupArgument.cc new file mode 100644 index 0000000000..f6a139a4b5 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipSetupArgument.cc @@ -0,0 +1,90 @@ +/* +Copyright (c) 2024 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 + +__global__ void add_vectors(int* a, int* b, int* result, int size) { + int tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid < size) { + result[tid] = a[tid] + b[tid]; + } +} + +TEST_CASE("Unit_hipSetupArgument_Simple") { + dim3 grid_dim(1, 1, 1); + dim3 block_dim(1, 1, 1); + + HIP_CHECK(hipConfigureCall(grid_dim, block_dim, 0, 0)); + + int arg = 42; + HIP_CHECK(hipSetupArgument(&arg, sizeof(int), 0)); +} + +/** + * Test Description + * ------------------------ + * - Verifies that arguments sent to the kernel with hipSetupArgument are correct by executing + * kernel that calculates sum of two vectors, doing the same calculation on CPU and checking if the + * results are the same, which proves that the arguments used in kernel are the proper ones + * + * Test source + * ------------------------ + * - unit/kernel/hipSetupArgument.cc + */ +TEST_CASE("Unit_hipSetupArgument_Execute_Kernel_And_Check_Result") { + constexpr auto block_size = 32; + auto vec_size = 256; + size_t block_num = static_cast(std::ceil(static_cast(vec_size) / block_size)); + dim3 grid_dim(block_num, 1, 1); + dim3 block_dim(block_size, 1, 1); + + HIP_CHECK(hipConfigureCall(grid_dim, block_dim, 0, 0)); + + std::vector vec_a, vec_b, vec_c; + for (size_t i = 0; i < vec_size; i++) { + vec_a.push_back(getRandom()); + vec_b.push_back(getRandom()); + vec_c.push_back(vec_a[i] + vec_b[i]); + } + + size_t vec_size_in_bytes = sizeof(int) * vec_size; + int *dev_a, *dev_b, *dev_c; + HIP_CHECK(hipMalloc(&dev_a, vec_size_in_bytes)); + HIP_CHECK(hipMalloc(&dev_b, vec_size_in_bytes)); + HIP_CHECK(hipMalloc(&dev_c, vec_size_in_bytes)); + + HIP_CHECK(hipMemcpy(dev_a, vec_a.data(), vec_size_in_bytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(dev_b, vec_b.data(), vec_size_in_bytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipSetupArgument(&dev_a, sizeof(int*), 0)); + HIP_CHECK(hipSetupArgument(&dev_b, sizeof(int*), sizeof(int*))); + HIP_CHECK(hipSetupArgument(&dev_c, sizeof(int*), 2 * sizeof(int*))); + HIP_CHECK(hipSetupArgument(&vec_size, sizeof(int), 3 * sizeof(int*))); + + HIP_CHECK(hipLaunchByPtr((const void*)add_vectors)); + + std::vector vec_res(vec_size); + HIP_CHECK(hipMemcpy(vec_res.data(), dev_c, vec_size_in_bytes, hipMemcpyDeviceToHost)); + REQUIRE(vec_res == vec_c); + + HIP_CHECK(hipFree(dev_c)); + HIP_CHECK(hipFree(dev_b)); + HIP_CHECK(hipFree(dev_a)); +}