SWDEV-404445 - [catch2][dtest] Adding test cases for uniform/non-uniform workgroup feature (#355)
Change-Id: I6d0764c1a42e08034158b9138695fee7bdc3f93e
Tento commit je obsažen v:
ROCm CI Service Account
odevzdal
GitHub
GitHub
rodič
3e1169eb51
revize
dce969d702
@@ -36,6 +36,7 @@ add_subdirectory(compiler)
|
||||
add_subdirectory(errorHandling)
|
||||
add_subdirectory(cooperativeGrps)
|
||||
add_subdirectory(context)
|
||||
add_subdirectory(module)
|
||||
if(HIP_PLATFORM STREQUAL "amd")
|
||||
add_subdirectory(callback)
|
||||
#add_subdirectory(clock)
|
||||
|
||||
@@ -0,0 +1,48 @@
|
||||
# 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.
|
||||
|
||||
# Common Tests - Test independent of all platforms
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set(TEST_SRC
|
||||
hipExtModuleLaunchKernel.cc
|
||||
)
|
||||
|
||||
# Note to pass arch use format like -DOFFLOAD_ARCH_STR="--offload-arch=gfx900 --offload-arch=gfx906"
|
||||
# having space at the start/end of OFFLOAD_ARCH_STR can cause build failures
|
||||
|
||||
add_custom_target(copyKernel.code
|
||||
COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --genco ${OFFLOAD_ARCH_STR}
|
||||
${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
|
||||
-o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernel.code
|
||||
-I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/
|
||||
-I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH})
|
||||
|
||||
add_custom_target(copyKernel.s
|
||||
COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 -S ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
|
||||
-o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernel.s
|
||||
-I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/
|
||||
-I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH})
|
||||
|
||||
hip_add_exe_to_target(NAME ModuleTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
|
||||
|
||||
add_dependencies(build_tests copyKernel.code copyKernel.s)
|
||||
endif()
|
||||
@@ -0,0 +1,29 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
extern "C" __global__ void copy_ker(int* Ad, int *Bd, size_t size) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if (myId < size) {
|
||||
Bd[myId] = Ad[myId];
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,239 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipExtModuleLaunchKernel hipExtModuleLaunchKernel
|
||||
* @{
|
||||
* @ingroup KernelTest
|
||||
* `hipError_t hipExtModuleLaunchKernel (hipFunction_t f,
|
||||
* uint32_t globalWorkSizeX,
|
||||
* uint32_t globalWorkSizeY,
|
||||
* uint32_t globalWorkSizeZ,
|
||||
* uint32_t localWorkSizeX,
|
||||
* uint32_t localWorkSizeY,
|
||||
* uint32_t localWorkSizeZ,
|
||||
* size_t sharedMemBytes,
|
||||
* hipStream_t hStream,
|
||||
* void ** kernelParams,
|
||||
* void ** extra,
|
||||
* hipEvent_t startEvent = nullptr,
|
||||
* hipEvent_t stopEvent = nullptr,
|
||||
* uint32_t flags = 0
|
||||
* )` -
|
||||
* Launches kernel with parameters and shared memory on stream with arguments
|
||||
* passed to kernel params or extra arguments.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_defgroups.hh>
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include "hip/hip_ext.h"
|
||||
#include <regex> // NOLINT
|
||||
|
||||
static constexpr auto totalWorkGroups{1024};
|
||||
static constexpr auto localWorkSize{512};
|
||||
static constexpr auto lastWorkSizeEven{256};
|
||||
static constexpr auto lastWorkSizeOdd{257};
|
||||
|
||||
#define fileName "copyKernel.code"
|
||||
#define kernel_name "copy_ker"
|
||||
|
||||
/**
|
||||
Local Function to search a string in file.
|
||||
*/
|
||||
static bool searchRegExpr(const std::regex& expr, const char* filename) {
|
||||
std::ifstream assemblyfile(filename, std::ifstream::binary);
|
||||
REQUIRE(true == assemblyfile.is_open());
|
||||
|
||||
// Copy the contents of the file to buffer
|
||||
assemblyfile.seekg(0, assemblyfile.end);
|
||||
int len = assemblyfile.tellg();
|
||||
assemblyfile.seekg(0, assemblyfile.beg);
|
||||
char *fbuf = new char[len + 1];
|
||||
assemblyfile.read(fbuf, len);
|
||||
fbuf[len] = '\0';
|
||||
|
||||
// Search for uniform_work_group_size
|
||||
std::smatch pattern;
|
||||
std::string assemblyStr = fbuf;
|
||||
bool bfound = std::regex_search(assemblyStr, pattern, expr);
|
||||
delete[] fbuf;
|
||||
assemblyfile.close();
|
||||
return bfound;
|
||||
}
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Parse the Code Object Assembly file copyKernel.s and verify
|
||||
* if .uniform_work_group_size metadata is available.
|
||||
* ------------------------
|
||||
* - catch\unit\module\hipExtModuleLaunchKernel.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEST_CASE("Unit_hipExtModuleLaunchKernel_CheckCodeObjAttr") {
|
||||
// Open copyKernel.s and read the file
|
||||
const std::regex regexp("uniform_work_group_size\\s*:\\s*[0-1]");
|
||||
REQUIRE(true == searchRegExpr(regexp, "copyKernel.s"));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Precondition: .uniform_work_group_size = 1. Which means uniform workgroup
|
||||
* is enforced.
|
||||
* - Create a buffer of size globalWorkSizeX = (non multiple of localWorkSizeX)
|
||||
* and launch a kernel to perform some operations on it. If uniform work grouping
|
||||
* is enforced then hipExtModuleLaunchKernel returns hipErrorInvalidValue.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\module\hipExtModuleLaunchKernel.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEST_CASE("Unit_hipExtModuleLaunchKernel_NonUniformWorkGroup") {
|
||||
// first check if uniform_work_group_size = 1.
|
||||
const std::regex regexp("uniform_work_group_size\\s*:\\s*1");
|
||||
if (false == searchRegExpr(regexp, "copyKernel.s")) {
|
||||
HipTest::HIP_SKIP_TEST("uniform_work_group_size != 1. Skipping test ...");
|
||||
return;
|
||||
}
|
||||
REQUIRE(true == searchRegExpr(regexp, "copyKernel.s"));
|
||||
auto isEven = GENERATE(0, 1);
|
||||
// Calculate size
|
||||
auto lastWorkSize = isEven ? lastWorkSizeEven : lastWorkSizeOdd;
|
||||
size_t arraylength =
|
||||
(totalWorkGroups - 1)*localWorkSize + lastWorkSize;
|
||||
size_t sizeBytes{arraylength * sizeof(int)};
|
||||
// Get module and function from module
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileName));
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
// Allocate resources
|
||||
int *A = new int[arraylength];
|
||||
REQUIRE(A != nullptr);
|
||||
int *B = new int[arraylength];
|
||||
REQUIRE(B != nullptr);
|
||||
// Inititialize data
|
||||
for (size_t i = 0; i < arraylength; i++) {
|
||||
A[i] = i;
|
||||
}
|
||||
int *Ad, *Bd;
|
||||
HIP_CHECK(hipMalloc(&Ad, sizeBytes));
|
||||
HIP_CHECK(hipMalloc(&Bd, sizeBytes));
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
size_t buffersize;
|
||||
} args;
|
||||
|
||||
args._Ad = Ad;
|
||||
args._Bd = Bd;
|
||||
args.buffersize = arraylength;
|
||||
size_t size = sizeof(args);
|
||||
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
// Memcpy from A to Ad
|
||||
HIP_CHECK(hipMemcpy(Ad, A, sizeBytes, hipMemcpyDefault));
|
||||
REQUIRE(hipErrorInvalidValue == hipExtModuleLaunchKernel(Function,
|
||||
arraylength, 1, 1, localWorkSize, 1, 1, 0, 0, NULL,
|
||||
reinterpret_cast<void**>(&config), 0));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipFree(Ad));
|
||||
HIP_CHECK(hipFree(Bd));
|
||||
delete[] A;
|
||||
delete[] B;
|
||||
HIP_CHECK(hipModuleUnload(Module));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Create a buffer of size globalWorkSizeX = (multiple of localWorkSizeX)
|
||||
* and launch a kernel to perform some operations on it. Verify the output. This
|
||||
* operation should be allowed for both uniform_work_group_size = true/false.
|
||||
*
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\module\hipExtModuleLaunchKernel.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") {
|
||||
size_t arraylength = totalWorkGroups * localWorkSize;
|
||||
size_t sizeBytes{arraylength * sizeof(int)};
|
||||
// Get module and function from module
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileName));
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
// Allocate resources
|
||||
int *A = new int[arraylength];
|
||||
REQUIRE(A != nullptr);
|
||||
int *B = new int[arraylength];
|
||||
REQUIRE(B != nullptr);
|
||||
// Inititialize data
|
||||
for (size_t i = 0; i < arraylength; i++) {
|
||||
A[i] = i;
|
||||
}
|
||||
int *Ad, *Bd;
|
||||
HIP_CHECK(hipMalloc(&Ad, sizeBytes));
|
||||
HIP_CHECK(hipMalloc(&Bd, sizeBytes));
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
size_t buffersize;
|
||||
} args;
|
||||
|
||||
args._Ad = Ad;
|
||||
args._Bd = Bd;
|
||||
args.buffersize = arraylength;
|
||||
size_t size = sizeof(args);
|
||||
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
// Memcpy from A to Ad
|
||||
HIP_CHECK(hipMemcpy(Ad, A, sizeBytes, hipMemcpyDefault));
|
||||
HIP_CHECK(hipExtModuleLaunchKernel(Function, arraylength, 1, 1,
|
||||
localWorkSize, 1, 1, 0, 0, NULL,
|
||||
reinterpret_cast<void**>(&config), 0));
|
||||
// Memcpy results back to host
|
||||
HIP_CHECK(hipMemcpy(B, Bd, sizeBytes, hipMemcpyDefault));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
// Verify results
|
||||
for (size_t i = 0; i < arraylength; i++) {
|
||||
REQUIRE(B[i] == i);
|
||||
}
|
||||
HIP_CHECK(hipFree(Ad));
|
||||
HIP_CHECK(hipFree(Bd));
|
||||
delete[] A;
|
||||
delete[] B;
|
||||
HIP_CHECK(hipModuleUnload(Module));
|
||||
}
|
||||
Odkázat v novém úkolu
Zablokovat Uživatele