From 9cadc3b875ba97bd041bc78198cab73e4e18b6bb Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Fri, 8 Dec 2023 07:24:30 +0000 Subject: [PATCH] EXSWHTEC-306 - Implement tests for memory device functions memcpy/memset (#279) Change-Id: I609efbd4f6ebaa48571e14e11ff97340dca1ab95 --- catch/hipTestMain/config/config_amd_windows | 3 + .../config/config_nvidia_linux.json | 5 +- catch/unit/CMakeLists.txt | 2 +- catch/unit/compileAndCaptureOutput.py | 4 +- catch/unit/device_memory/CMakeLists.txt | 56 ++++ catch/unit/device_memory/memcpy.cc | 249 ++++++++++++++++++ .../device_memory/memcpy_negative_kernels.cc | 31 +++ .../memcpy_negative_kernels_rtc.hh | 32 +++ catch/unit/device_memory/memset.cc | 172 ++++++++++++ .../device_memory/memset_negative_kernels.cc | 31 +++ .../memset_negative_kernels_rtc.hh | 32 +++ 11 files changed, 613 insertions(+), 4 deletions(-) create mode 100644 catch/unit/device_memory/CMakeLists.txt create mode 100644 catch/unit/device_memory/memcpy.cc create mode 100644 catch/unit/device_memory/memcpy_negative_kernels.cc create mode 100644 catch/unit/device_memory/memcpy_negative_kernels_rtc.hh create mode 100644 catch/unit/device_memory/memset.cc create mode 100644 catch/unit/device_memory/memset_negative_kernels.cc create mode 100644 catch/unit/device_memory/memset_negative_kernels_rtc.hh diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index e7d51ba0e2..f9080bed55 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -332,6 +332,9 @@ "Unit_Printf_length_Sanity_Positive", "Unit_Printf_Negative", "Unit_Device__hip_hc_8pk_Negative", + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/279 ===", + "Unit_Device_memcpy_Negative", + "Unit_Device_memset_Negative", #endif "End of json" ] diff --git a/catch/hipTestMain/config/config_nvidia_linux.json b/catch/hipTestMain/config/config_nvidia_linux.json index b6ea489d58..5daf1caa1b 100644 --- a/catch/hipTestMain/config/config_nvidia_linux.json +++ b/catch/hipTestMain/config/config_nvidia_linux.json @@ -54,6 +54,9 @@ "Unit_atomicExch_system_Positive_Host_And_GPU - unsigned long long", "Unit_atomicExch_system_Positive_Host_And_GPU - float", "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/274 ===", - "Unit_Printf_Negative" + "Unit_Printf_Negative", + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/279 ===", + "Unit_Device_memcpy_Negative", + "Unit_Device_memset_Negative" ] } diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 7845a3c6f9..5557f5c3cc 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -37,7 +37,7 @@ add_subdirectory(errorHandling) add_subdirectory(cooperativeGrps) add_subdirectory(warp) add_subdirectory(context) -add_subdirectory(warp) +add_subdirectory(device_memory) add_subdirectory(dynamicLoading) add_subdirectory(g++) add_subdirectory(module) diff --git a/catch/unit/compileAndCaptureOutput.py b/catch/unit/compileAndCaptureOutput.py index c06feb5858..a8a7fb506a 100644 --- a/catch/unit/compileAndCaptureOutput.py +++ b/catch/unit/compileAndCaptureOutput.py @@ -80,7 +80,7 @@ if __name__ == '__main__': CompileAndCapture.platform = sys.argv[2] except IndexError: CompileAndCapture.platform = None - + try: CompileAndCapture.hip_path = sys.argv[3] except IndexError: @@ -95,7 +95,7 @@ if __name__ == '__main__': CompileAndCapture.expected_error_count = int(sys.argv[5]) except IndexError: CompileAndCapture.expected_error_count = 0 - + try: CompileAndCapture.expected_warning_count = int(sys.argv[6]) except IndexError: diff --git a/catch/unit/device_memory/CMakeLists.txt b/catch/unit/device_memory/CMakeLists.txt new file mode 100644 index 0000000000..54ff74e275 --- /dev/null +++ b/catch/unit/device_memory/CMakeLists.txt @@ -0,0 +1,56 @@ +# Copyright (c) 2022 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 + +set(TEST_SRC + memcpy.cc + memset.cc +) + +if(HIP_PLATFORM MATCHES "nvidia") + set(LINKER_LIBS nvrtc) +elseif(HIP_PLATFORM MATCHES "amd") + set(LINKER_LIBS hiprtc) +endif() + +if(HIP_PLATFORM MATCHES "amd") + hip_add_exe_to_target(NAME DeviceMemoryTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS ${LINKER_LIBS} + PROPERTY CXX_STANDARD 17) +elseif (HIP_PLATFORM MATCHES "nvidia") + hip_add_exe_to_target(NAME DeviceMemoryTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS ${LINKER_LIBS} + COMPILE_OPTIONS -std=c++17) +endif() + +add_test(NAME Unit_Device_memcpy_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + memcpy_negative_kernels.cc 4) + +add_test(NAME Unit_Device_memset_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + memset_negative_kernels.cc 4) diff --git a/catch/unit/device_memory/memcpy.cc b/catch/unit/device_memory/memcpy.cc new file mode 100644 index 0000000000..efb3bea567 --- /dev/null +++ b/catch/unit/device_memory/memcpy.cc @@ -0,0 +1,249 @@ +/* +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 "memcpy_negative_kernels_rtc.hh" + +#include +#include +#include + +#include + +/** + * @addtogroup memcpy memcpy + * @{ + * @ingroup DeviceLanguageTest + * `memcpy(void* dst, const void* src, size_t size)` - + * copies device accessible data inside a kernel + */ + +template using kernel_sig = void (*)(T*, T*, const size_t); + +template +__global__ void memcpy_at_once_kernel(T* dst, T* src, const size_t alloc_size) { + memcpy(dst, src, alloc_size); +} + +template __global__ void memcpy_one_by_one_kernel(T* dst, T* src, const size_t N) { + const auto tid = cooperative_groups::this_grid().thread_rank(); + const auto stride = cooperative_groups::this_grid().size(); + + for (auto i = tid; i < N; i += stride) { + memcpy(dst + tid, src + tid, sizeof(T)); + } +} + +template void MemcpyDeviceToDeviceCommon(kernel_sig memcpy_kernel) { + const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); + const auto element_count = allocation_size / sizeof(T); + + LinearAllocGuard input(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard src_allocation(LinearAllocs::hipMalloc, allocation_size); + LinearAllocGuard dst_allocation(LinearAllocs::hipMalloc, allocation_size); + + /* fill input data */ + for (auto i = 0; i < element_count; i++) { + input.host_ptr()[i] = static_cast(i); + } + + /* Copy input data to device memory */ + HIP_CHECK( + hipMemcpy(src_allocation.ptr(), input.host_ptr(), allocation_size, hipMemcpyHostToDevice)); + + /* Launch appropriate kernel*/ + if (memcpy_kernel == &memcpy_at_once_kernel) { + memcpy_at_once_kernel<<<1, 1>>>(dst_allocation.ptr(), src_allocation.ptr(), allocation_size); + } else { + constexpr auto thread_count = 1024; + const auto block_count = element_count / thread_count + 1; + memcpy_one_by_one_kernel + <<>>(dst_allocation.ptr(), src_allocation.ptr(), element_count); + } + + /* Copy filled device memory to result */ + HIP_CHECK( + hipMemcpy(result.host_ptr(), dst_allocation.ptr(), allocation_size, hipMemcpyDeviceToHost)); + + ArrayMismatch(input.host_ptr(), result.host_ptr(), element_count); +} + +template void MemcpyPinnedToDeviceCommon(kernel_sig memcpy_kernel) { + const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); + const auto element_count = allocation_size / sizeof(T); + + LinearAllocGuard input(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard dst_allocation(LinearAllocs::hipMalloc, allocation_size); + + /* fill input data */ + for (auto i = 0; i < element_count; i++) { + input.host_ptr()[i] = static_cast(i); + } + + /* Launch appropriate kernel*/ + if (memcpy_kernel == &memcpy_at_once_kernel) { + memcpy_at_once_kernel<<<1, 1>>>(dst_allocation.ptr(), input.host_ptr(), allocation_size); + } else { + constexpr auto thread_count = 1024; + const auto block_count = element_count / thread_count + 1; + memcpy_one_by_one_kernel + <<>>(dst_allocation.ptr(), input.host_ptr(), element_count); + } + + /* Copy filled device memory to result */ + HIP_CHECK( + hipMemcpy(result.host_ptr(), dst_allocation.ptr(), allocation_size, hipMemcpyDeviceToHost)); + + ArrayMismatch(input.host_ptr(), result.host_ptr(), element_count); +} + +template void MemcpyDeviceToPinnedCommon(kernel_sig memcpy_kernel) { + const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); + const auto element_count = allocation_size / sizeof(T); + + LinearAllocGuard input(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard src_allocation(LinearAllocs::hipMalloc, allocation_size); + + /* fill input data */ + for (auto i = 0; i < element_count; i++) { + input.host_ptr()[i] = static_cast(i); + } + + /* Copy input data to device memory */ + HIP_CHECK( + hipMemcpy(src_allocation.ptr(), input.host_ptr(), allocation_size, hipMemcpyHostToDevice)); + + /* Launch appropriate kernel*/ + if (memcpy_kernel == &memcpy_at_once_kernel) { + memcpy_at_once_kernel<<<1, 1>>>(result.host_ptr(), src_allocation.ptr(), allocation_size); + } else { + constexpr auto thread_count = 1024; + const auto block_count = element_count / thread_count + 1; + memcpy_one_by_one_kernel + <<>>(result.host_ptr(), src_allocation.ptr(), element_count); + } + + HIP_CHECK(hipStreamSynchronize(nullptr)); + + ArrayMismatch(input.host_ptr(), result.host_ptr(), element_count); +} + +template void MemcpyPinnedToPinnedCommon(kernel_sig memcpy_kernel) { + const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); + const auto element_count = allocation_size / sizeof(T); + + LinearAllocGuard input(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size); + + /* fill input data */ + for (auto i = 0; i < element_count; i++) { + input.host_ptr()[i] = static_cast(i); + } + + /* Launch appropriate kernel*/ + if (memcpy_kernel == &memcpy_at_once_kernel) { + memcpy_at_once_kernel<<<1, 1>>>(result.host_ptr(), input.host_ptr(), allocation_size); + } else { + constexpr auto thread_count = 1024; + const auto block_count = element_count / thread_count + 1; + memcpy_one_by_one_kernel + <<>>(result.host_ptr(), input.host_ptr(), element_count); + } + + HIP_CHECK(hipStreamSynchronize(nullptr)); + + ArrayMismatch(input.host_ptr(), result.host_ptr(), element_count); +} + +template void DeviceMemcpyCommon(kernel_sig memcpy_kernel) { + SECTION("Device to Device memory") { MemcpyDeviceToDeviceCommon(memcpy_kernel); } + + SECTION("Pinned to Device memory") { MemcpyPinnedToDeviceCommon(memcpy_kernel); } + + SECTION("Device to Pinned memory") { MemcpyDeviceToPinnedCommon(memcpy_kernel); } + + SECTION("Pinned to Pinned memory") { MemcpyPinnedToPinnedCommon(memcpy_kernel); } +} + + +/** + * Test Description + * ------------------------ + * - Verifies basic test cases for copying device/pinned memory inside a kernel using various data + * types and memory sizes: + * -# Copies whole memory buffer in one thread + * -# Copies memory buffer elements one by one in multiple threads/blocks + * Test source + * ------------------------ + * - unit/device_memory/memcpy.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_memcpy_Positive", "", char, int, unsigned int, long, unsigned long, + long long, unsigned long long, float, double) { + SECTION("Memcpy whole buffer in one thread") { + DeviceMemcpyCommon(memcpy_at_once_kernel); + } + SECTION("Memcpy buffer in multiple threads/blocks") { + DeviceMemcpyCommon(memcpy_one_by_one_kernel); + } +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass combinations of arguments of invalid types for memcpy + * Test source + * ------------------------ + * - unit/device_memory/memcpy.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_memcpy_Negative_Parameters_RTC") { + hiprtcProgram program{}; + + const auto program_source = kMemcpyParam; + + HIPRTC_CHECK( + hiprtcCreateProgram(&program, program_source, "memcpy_negative.cc", 0, nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log and count compiler error messages + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + + int expected_error_count{4}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_error_count); +} diff --git a/catch/unit/device_memory/memcpy_negative_kernels.cc b/catch/unit/device_memory/memcpy_negative_kernels.cc new file mode 100644 index 0000000000..84621dca8b --- /dev/null +++ b/catch/unit/device_memory/memcpy_negative_kernels.cc @@ -0,0 +1,31 @@ +/* +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 + +struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +/*void* memcpy(void* dst, const void* src, size_t size)*/ +__global__ void memcpy_n1(int* dst, const int src, size_t size) { memcpy(dst, src, size); } +__global__ void memcpy_n2(int dst, const int* src, size_t size) { memcpy(dst, src, size); } +__global__ void memcpy_n3(int* dst, const int* src, size_t* size) { memcpy(dst, src, size); } +__global__ void memcpy_n4(int* dst, const int* src, Dummy size) { memcpy(dst, src, size); } \ No newline at end of file diff --git a/catch/unit/device_memory/memcpy_negative_kernels_rtc.hh b/catch/unit/device_memory/memcpy_negative_kernels_rtc.hh new file mode 100644 index 0000000000..ac86844ece --- /dev/null +++ b/catch/unit/device_memory/memcpy_negative_kernels_rtc.hh @@ -0,0 +1,32 @@ +/* +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. +*/ + +#pragma once + +static constexpr auto kMemcpyParam{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void memcpy_n1(int* dst, const int src, size_t size) { memcpy(dst, src, size); } + __global__ void memcpy_n2(int dst, const int* src, size_t size) { memcpy(dst, src, size); } + __global__ void memcpy_n3(int* dst, const int* src, size_t* size) { memcpy(dst, src, size); } + __global__ void memcpy_n8(int* dst, const int* src, Dummy size) { memcpy(dst, src, size); } + )"}; \ No newline at end of file diff --git a/catch/unit/device_memory/memset.cc b/catch/unit/device_memory/memset.cc new file mode 100644 index 0000000000..902a45111e --- /dev/null +++ b/catch/unit/device_memory/memset.cc @@ -0,0 +1,172 @@ +/* +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 "memset_negative_kernels_rtc.hh" + +#include +#include +#include + +#include + +/** + * @addtogroup memset memset + * @{ + * @ingroup DeviceLanguageTest + * `memset(void* ptr, int val, size_t size)` - + * sets device accessible data inside a kernel + */ + +template using kernel_sig = void (*)(T*, int, const size_t); + +template +__global__ void memset_at_once_kernel(T* dst, int value, const size_t alloc_size) { + memset(dst, value, alloc_size); +} + +template __global__ void memset_one_by_one_kernel(T* dst, int value, const size_t N) { + const auto tid = cooperative_groups::this_grid().thread_rank(); + const auto stride = cooperative_groups::this_grid().size(); + + for (auto i = tid; i < N; i += stride) { + memset(dst + tid, value, sizeof(T)); + } +} + +template void MemsetDeviceCommon(kernel_sig memset_kernel) { + const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); + const auto element_count = allocation_size / sizeof(T); + + LinearAllocGuard reference(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard dst_allocation(LinearAllocs::hipMalloc, allocation_size); + + constexpr auto thread_count = 1024; + const auto block_count = element_count / thread_count + 1; + constexpr auto expected_value = 42; + + memset(reference.host_ptr(), expected_value, allocation_size); + + if (memset_kernel == &memset_at_once_kernel) { + memset_at_once_kernel<<<1, 1>>>(dst_allocation.ptr(), expected_value, allocation_size); + } else { + memset_one_by_one_kernel + <<>>(dst_allocation.ptr(), expected_value, element_count); + } + + HIP_CHECK( + hipMemcpy(result.host_ptr(), dst_allocation.ptr(), allocation_size, hipMemcpyDeviceToHost)); + + ArrayMismatch(reference.host_ptr(), result.host_ptr(), element_count); +} + +template void MemsetPinnedCommon(kernel_sig memset_kernel) { + const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); + const auto element_count = allocation_size / sizeof(T); + + LinearAllocGuard reference(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size); + + constexpr auto thread_count = 1024; + const auto block_count = element_count / thread_count + 1; + constexpr auto expected_value = 42; + + memset(reference.host_ptr(), expected_value, allocation_size); + + if (memset_kernel == &memset_at_once_kernel) { + memset_at_once_kernel<<<1, 1>>>(result.host_ptr(), expected_value, allocation_size); + } else { + memset_one_by_one_kernel + <<>>(result.host_ptr(), expected_value, element_count); + } + + HIP_CHECK(hipStreamSynchronize(nullptr)); + + ArrayMismatch(reference.host_ptr(), result.host_ptr(), element_count); +} + +template void DeviceMemsetCommon(kernel_sig memset_kernel) { + SECTION("Set Device memory") { MemsetDeviceCommon(memset_kernel); } + + SECTION("Set Pinned memory") { MemsetPinnedCommon(memset_kernel); } +} + +/** + * Test Description + * ------------------------ + * - Verifies basic test cases for setting device/pinned memory inside a kernel using various data + * types and memory sizes: + * -# Set whole memory buffer in one thread + * -# Set memory buffer elements one by one in multiple threads/blocks + * Test source + * ------------------------ + * - unit/device_memory/memset.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_memset_Positive", "", char, int, unsigned int, long, unsigned long, + long long, unsigned long long, float, double) { + SECTION("Memset whole buffer in one thread") { + DeviceMemsetCommon(memset_at_once_kernel); + } + SECTION("Memset buffer in multiple threads/blocks") { + DeviceMemsetCommon(memset_one_by_one_kernel); + } +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass combinations of arguments of invalid types for memset + * Test source + * ------------------------ + * - unit/device_memory/memset.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_memset_Negative_Parameters_RTC") { + hiprtcProgram program{}; + + const auto program_source = kMemsetParam; + + HIPRTC_CHECK( + hiprtcCreateProgram(&program, program_source, "memset_negative.cc", 0, nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log and count compiler error messages + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + + int expected_error_count{4}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_error_count); +} \ No newline at end of file diff --git a/catch/unit/device_memory/memset_negative_kernels.cc b/catch/unit/device_memory/memset_negative_kernels.cc new file mode 100644 index 0000000000..e6d1c19e04 --- /dev/null +++ b/catch/unit/device_memory/memset_negative_kernels.cc @@ -0,0 +1,31 @@ +/* +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 + +struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +/*void* memset(void* ptr, int val, size_t size)*/ +__global__ void memset_n1(int* ptr, int* val, size_t size) { memset(ptr, val, size); } +__global__ void memset_n2(int ptr, int val, size_t size) { memset(ptr, val, size); } +__global__ void memset_n3(int* ptr, int val, size_t* size) { memset(ptr, val, size); } +__global__ void memset_n4(int* ptr, int val, Dummy size) { memset(ptr, val, size); } \ No newline at end of file diff --git a/catch/unit/device_memory/memset_negative_kernels_rtc.hh b/catch/unit/device_memory/memset_negative_kernels_rtc.hh new file mode 100644 index 0000000000..43d3c513b0 --- /dev/null +++ b/catch/unit/device_memory/memset_negative_kernels_rtc.hh @@ -0,0 +1,32 @@ +/* +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. +*/ + +#pragma once + +static constexpr auto kMemsetParam{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void memset_n1(int* ptr, int* val, size_t size) { memset(ptr, val, size); } + __global__ void memset_n2(int ptr, int val, size_t size) { memset(ptr, val, size); } + __global__ void memset_n3(int* ptr, int val, size_t* size) { memset(ptr, val, size); } + __global__ void memset_n4(int* ptr, int val, Dummy size) { memset(ptr, val, size); } + )"}; \ No newline at end of file