From 4c7a0c7bef17e6e3b99c5c68c627b0b83fbbcc34 Mon Sep 17 00:00:00 2001 From: milos-mozetic <118800401+milos-mozetic@users.noreply.github.com> Date: Tue, 18 Jul 2023 09:23:27 +0200 Subject: [PATCH 01/17] EXSWHTEC-231 - Implement custom Benchmarking tool for Performance tests (#117) - Introduce performance tests to project. - Improve benchmarking utilities. - Delete copy constructors from Timer and Benchmark classes. - Disable Catch2's benchmarking functionalities. - Address review comments and add progress bar/display output to the Benchmarking tool - Add flushing of the buffer on the benchmark name display - Introduce command line options. - Add allocation type to string helper function. - Add output modifier to Benchmark class. - Fix invalid calculation of deviation - Update performance_common.hh - Resolve build error on Windows by adding include for reduce and accumulate [ROCm/hip-tests commit: 5fbbdcae68c55b8f308c5ce8d7f9ba4dd0cf3b9c] --- projects/hip-tests/catch/CMakeLists.txt | 1 + projects/hip-tests/catch/hipTestMain/main.cc | 32 ++- .../hip-tests/catch/include/cmd_options.hh | 33 +++ .../catch/include/performance_common.hh | 247 ++++++++++++++++++ .../catch/performance/CMakeLists.txt | 21 ++ .../catch/performance/example/CMakeLists.txt | 28 ++ .../catch/performance/example/example.cc | 59 +++++ 7 files changed, 419 insertions(+), 2 deletions(-) create mode 100644 projects/hip-tests/catch/include/cmd_options.hh create mode 100644 projects/hip-tests/catch/include/performance_common.hh create mode 100644 projects/hip-tests/catch/performance/CMakeLists.txt create mode 100644 projects/hip-tests/catch/performance/example/CMakeLists.txt create mode 100644 projects/hip-tests/catch/performance/example/example.cc diff --git a/projects/hip-tests/catch/CMakeLists.txt b/projects/hip-tests/catch/CMakeLists.txt index 1b53c746d3..f9078da432 100644 --- a/projects/hip-tests/catch/CMakeLists.txt +++ b/projects/hip-tests/catch/CMakeLists.txt @@ -260,6 +260,7 @@ add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers) if(UNIX) add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc) endif() +add_subdirectory(performance ${CATCH_BUILD_DIR}/performance) cmake_policy(POP) diff --git a/projects/hip-tests/catch/hipTestMain/main.cc b/projects/hip-tests/catch/hipTestMain/main.cc index ea3bb9bd23..030267f04c 100644 --- a/projects/hip-tests/catch/hipTestMain/main.cc +++ b/projects/hip-tests/catch/hipTestMain/main.cc @@ -1,7 +1,10 @@ #define CATCH_CONFIG_RUNNER +#include #include #include +CmdOptions cmd_options; + int main(int argc, char** argv) { auto& context = TestContext::get(argc, argv); if (context.skipTest()) { @@ -9,8 +12,33 @@ int main(int argc, char** argv) { std::cout << "HIP_SKIP_THIS_TEST" << std::endl; return 0; } - int out = Catch::Session().run(argc, argv); + + Catch::Session session; + + using namespace Catch::clara; + // clang-format off + auto cli = session.cli() + | Opt(cmd_options.iterations, "iterations") + ["-I"]["--iterations"] + ("Number of iterations used for performance tests (default: 1000)") + | Opt(cmd_options.warmups, "warmups") + ["-W"]["--warmups"] + ("Number of warmup iterations used for performance tests (default: 100)") + | Opt(cmd_options.no_display) + ["-S"]["--no-display"] + ("Do not display the output of performance tests") + | Opt(cmd_options.progress) + ["-P"]["--progress"] + ("Show progress bar when running performance tests") + | Opt(cmd_options.extended_run) + ["-E"]["--extended-run"] + ("TODO: Description goes here") + ; + // clang-format on + + session.cli(cli); + + int out = session.run(argc, argv); TestContext::get().cleanContext(); return out; - } diff --git a/projects/hip-tests/catch/include/cmd_options.hh b/projects/hip-tests/catch/include/cmd_options.hh new file mode 100644 index 0000000000..5dbd2f300c --- /dev/null +++ b/projects/hip-tests/catch/include/cmd_options.hh @@ -0,0 +1,33 @@ +/* +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. +*/ + +#pragma once + +struct CmdOptions { + int iterations = 1000; + int warmups = 100; + bool no_display = false; + bool progress = false; + bool extended_run = false; +}; + +extern CmdOptions cmd_options; \ No newline at end of file diff --git a/projects/hip-tests/catch/include/performance_common.hh b/projects/hip-tests/catch/include/performance_common.hh new file mode 100644 index 0000000000..a40d6b3f5c --- /dev/null +++ b/projects/hip-tests/catch/include/performance_common.hh @@ -0,0 +1,247 @@ +/* +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. +*/ + +#pragma once + +#include +#include +#include +#include +#include +#include + +#include +#include +#include + +#if defined(_WIN32) +#if defined(_WIN64) +typedef __int64 ssize_t; +#else // !_WIN64 +typedef __int32 ssize_t; +#endif // !_WIN64 +#endif /*_WIN32*/ + +class Timer { + public: + Timer(const Timer&) = delete; + Timer& operator=(const Timer&) = delete; + + protected: + Timer(float& time, hipStream_t stream) : time_(time), stream_(stream) {} + + void Record(float time) { time_ += time; } + + hipStream_t GetStream() const { return stream_; } + + private: + float& time_; + hipStream_t stream_; +}; + +class EventTimer : public Timer { + public: + EventTimer(float& time, hipStream_t stream = nullptr) : Timer(time, stream) { + HIP_CHECK(hipEventCreate(&start_)); + HIP_CHECK(hipEventCreate(&stop_)); + HIP_CHECK(hipEventRecord(start_, GetStream())); + } + + ~EventTimer() { + hipError_t error; // to avoid compiler warnings + + error = hipEventRecord(stop_, GetStream()); + error = hipEventSynchronize(stop_); + + float ms; + error = hipEventElapsedTime(&ms, start_, stop_); + Record(ms); + + error = hipEventDestroy(start_); + error = hipEventDestroy(stop_); + } + + private: + hipEvent_t start_; + hipEvent_t stop_; +}; + +class CpuTimer : public Timer { + public: + CpuTimer(float& time, hipStream_t stream = nullptr) : Timer(time, stream) { + start_ = std::chrono::steady_clock::now(); + } + + ~CpuTimer() { + hipError_t error; // to avoid compiler warnings + error = hipStreamSynchronize(GetStream()); + + stop_ = std::chrono::steady_clock::now(); + + std::chrono::duration ms = stop_ - start_; + Record(ms.count()); + } + + private: + std::chrono::time_point start_; + std::chrono::time_point stop_; +}; + +template class Benchmark { + public: + Benchmark() + : iterations_(cmd_options.iterations), + warmups_(cmd_options.warmups), + display_output_(!cmd_options.no_display), + progress_bar_(cmd_options.progress) { + benchmark_name_ = Catch::getResultCapture().getCurrentTestName(); + } + + Benchmark(const Benchmark&) = delete; + Benchmark& operator=(const Benchmark&) = delete; + + static constexpr ssize_t kWarmup = -1; + + void Configure(size_t iterations, size_t warmups) { + iterations_ = iterations; + warmups_ = warmups; + } + + void AddSectionName(const std::string& section_name) { benchmark_name_ += "/" + section_name; } + + using ModifierSignature = std::function; + void RegisterModifier(const ModifierSignature& modifier) { modifier_ = modifier; } + + template std::tuple Run(Args&&... args) { + AddSectionName(std::to_string(iterations_)); + AddSectionName(std::to_string(warmups_)); + + auto& derived = static_cast(*this); + + current_ = kWarmup; + for (size_t i = 0u; i < warmups_; ++i) { + PrintProgress("warmup", static_cast(100.f * (i + 1) / warmups_)); + derived(args...); + } + time_ = .0; + + std::vector samples; + samples.reserve(iterations_); + + for (current_ = 0; current_ < iterations_; ++current_) { + PrintProgress("measurement", static_cast(100.f * (current_ + 1) / iterations_)); + derived(args...); + if (modifier_) time_ = modifier_(time_); + samples.push_back(time_); + time_ = .0; + } + + float sum = std::reduce(cbegin(samples), cend(samples)); + float mean = sum / samples.size(); + + float deviation = + std::accumulate(cbegin(samples), cend(samples), .0, + [mean](float sum, float next) { return sum + std::pow(next - mean, 2); }); + deviation = sqrt(deviation / samples.size()); + + float best = *std::min_element(cbegin(samples), cend(samples)); + float worst = *std::max_element(cbegin(samples), cend(samples)); + + PrintStats(mean, deviation, best, worst); + + return {mean, deviation, best, worst}; + } + + protected: + template + using TimerType = std::conditional_t; + + template + std::unique_ptr> GetTimer(hipStream_t stream = nullptr) { + return std::make_unique>(time_, stream); + } + + float time() const { return time_; } + + size_t iterations() const { return iterations_; } + + size_t warmups() const { return warmups_; } + + ssize_t current() const { return current_; } + + private: + std::string benchmark_name_; + float time_; + size_t iterations_; + size_t warmups_; + ssize_t current_; + bool display_output_; + bool progress_bar_; + + ModifierSignature modifier_; + + void Print(const std::string& out = "") { + if (!display_output_) return; + std::cout << "\r" << std::setw(110) << std::left << benchmark_name_ << "\t|\t" << out + << std::flush; + } + + void PrintProgress(const std::string& name, int progress) { + if (!(display_output_ && progress_bar_)) return; + Print(name + ": [" + std::to_string(progress) + "%]"); + } + + void PrintStats(float mean, float deviation, float best, float worst) { + if (!display_output_) return; + Print("Average time: " + std::to_string(mean) + " ms, Standard deviation: " + + std::to_string(deviation) + " ms, Fastest: " + std::to_string(best) + + " ms, Slowest: " + std::to_string(worst) + " ms\n"); + } +}; + +constexpr bool kTimerTypeCpu = false; +constexpr bool kTimerTypeEvent = true; + +#define TIMED_SECTION_STREAM(TIMER_TYPE, STREAM) \ + if (auto _ = this->template GetTimer(STREAM); true) +#define TIMED_SECTION(TIMER_TYPE) TIMED_SECTION_STREAM(TIMER_TYPE, nullptr) + +constexpr size_t operator"" _KB(unsigned long long int kb) { return kb << 10; } + +constexpr size_t operator"" _MB(unsigned long long int mb) { return mb << 20; } + +constexpr size_t operator"" _GB(unsigned long long int gb) { return gb << 30; } + +static std::string GetAllocationSectionName(LinearAllocs allocation_type) { + switch (allocation_type) { + case LinearAllocs::malloc: + return "host pageable"; + case LinearAllocs::hipHostMalloc: + return "host pinned"; + case LinearAllocs::hipMalloc: + return "device malloc"; + case LinearAllocs::hipMallocManaged: + return "managed"; + default: + return "unknown alloc type"; + } +} diff --git a/projects/hip-tests/catch/performance/CMakeLists.txt b/projects/hip-tests/catch/performance/CMakeLists.txt new file mode 100644 index 0000000000..5412636ebc --- /dev/null +++ b/projects/hip-tests/catch/performance/CMakeLists.txt @@ -0,0 +1,21 @@ +# 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. + +add_subdirectory(example) diff --git a/projects/hip-tests/catch/performance/example/CMakeLists.txt b/projects/hip-tests/catch/performance/example/CMakeLists.txt new file mode 100644 index 0000000000..001a9e7b1d --- /dev/null +++ b/projects/hip-tests/catch/performance/example/CMakeLists.txt @@ -0,0 +1,28 @@ +# 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. + +set(TEST_SRC + example.cc +) + +hip_add_exe_to_target(NAME ExamplePerformance + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++17) diff --git a/projects/hip-tests/catch/performance/example/example.cc b/projects/hip-tests/catch/performance/example/example.cc new file mode 100644 index 0000000000..2ced78b023 --- /dev/null +++ b/projects/hip-tests/catch/performance/example/example.cc @@ -0,0 +1,59 @@ +/* +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. +*/ + +#include +#include +#include + +class ExampleBenchmark : public Benchmark { + public: + void operator()(void* dst) { + const int value = 42; + const size_t kSize = 4_MB; + + TIMED_SECTION(kTimerTypeEvent) { // event based timing + HIP_CHECK(hipMemset(dst, value, kSize)); + } + + HIP_CHECK(hipMemset(dst, 0, kSize)); // not timed + + TIMED_SECTION(kTimerTypeCpu) { // cpu based timing + HIP_CHECK(hipMemset(dst, value, kSize)); + } + + // accessing properties + // std::cout << "Time recorded up until now: " << time() << std::endl; + // std::cout << "Number of iterations: " << iterations() << std::endl; + // std::cout << "Number of warmup iterations: " << warmups() << std::endl; + // std::cout << "Current iteration: " << current() << std::endl; + } +}; + +TEST_CASE("Performance_Example") { + ExampleBenchmark benchmark; + + // to override cmd options + // benchmark.Configure(10000 /* iterations */, 1000 /* warmups */); + + LinearAllocGuard dst(LinearAllocs::hipMalloc, 4_MB); + benchmark.Run(dst.ptr()); +} From d0abae43464f68607f97463f784481869f0d73b6 Mon Sep 17 00:00:00 2001 From: nives-vukovic <110852104+nives-vukovic@users.noreply.github.com> Date: Tue, 18 Jul 2023 09:25:00 +0200 Subject: [PATCH 02/17] EXSWHTEC-256 - Implement tests for grid_group APIs (#153) - Migrate basic Cooperative Groups tests and integrate to catch - Refactor basic Cooperative Groups tests - Rename tiled partition related files and fix minor bug - Add LaunchCooperativeKernal and LaunchCooperativeKernelMultiDevice tests - Refactor hipCGThreadBlockTileType to use common function - Fix updated file not added during merge - Add coalesced_group type tests - Add coalesced_group shuffle_up and shuffle_down tests - Add coalesced_group shuffle tests - test fails - Implement common code for cooperative group tests - Fixed compilation errror in cooperative_groups_common.hh - Implement busy wait device function - Reimplement tests for grid_group APIs - Add tests for grid_group member and non-member APIs - Refactor existing test for grid_group sync testing - Add thread and block dimensions generators - Add check of grid and block dimensions - Modify doxygen comments - Move cpu_grid.h and supporting functions to catch/include - Use warp_size from properties in grid/block dims generators - Fix condition for warp size 32 on AMD - Fix cpu_grid.h for warp function tests - Add missing include into cpu_grid.h - Code cleanup - Fix doxygen comments - Add missing include in utils header [ROCm/hip-tests commit: 22f3d9034ba458fc1e72f5c52da5d0f341d75621] --- projects/hip-tests/catch/include/cpu_grid.h | 154 ++++++++++ .../catch/include/hip_test_defgroups.hh | 7 + projects/hip-tests/catch/include/utils.hh | 15 + .../catch/unit/cooperativeGrps/CMakeLists.txt | 1 + .../cooperative_groups_common.hh | 71 +++++ .../catch/unit/cooperativeGrps/grid_group.cc | 285 ++++++++++++++++++ 6 files changed, 533 insertions(+) create mode 100644 projects/hip-tests/catch/include/cpu_grid.h create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/grid_group.cc diff --git a/projects/hip-tests/catch/include/cpu_grid.h b/projects/hip-tests/catch/include/cpu_grid.h new file mode 100644 index 0000000000..98e5521840 --- /dev/null +++ b/projects/hip-tests/catch/include/cpu_grid.h @@ -0,0 +1,154 @@ +/* +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 + +#include + +#include +#include + +struct CPUGrid { + CPUGrid() = default; + + CPUGrid(const dim3 grid_dim, const dim3 block_dim) + : grid_dim_{grid_dim}, + block_dim_{block_dim}, + block_count_{grid_dim.x * grid_dim.y * grid_dim.z}, + threads_in_block_count_{block_dim.x * block_dim.y * block_dim.z}, + thread_count_{block_count_ * threads_in_block_count_} {} + + inline std::optional thread_rank_in_block( + const unsigned int thread_rank_in_grid) const { + if (thread_rank_in_grid > thread_count_) { + return std::nullopt; + } + + return thread_rank_in_grid % threads_in_block_count_; + } + + inline std::optional block_idx(const unsigned int thread_rank_in_grid) const { + if (thread_rank_in_grid > thread_count_) { + return std::nullopt; + } + + dim3 block_idx; + const auto block_rank_in_grid = thread_rank_in_grid / threads_in_block_count_; + block_idx.x = block_rank_in_grid % grid_dim_.x; + block_idx.y = (block_rank_in_grid / grid_dim_.x) % grid_dim_.y; + block_idx.z = block_rank_in_grid / (grid_dim_.x * grid_dim_.y); + + return block_idx; + } + + inline std::optional thread_idx(const unsigned int thread_rank_in_grid) const { + if (thread_rank_in_grid > thread_count_) { + return std::nullopt; + } + + dim3 thread_idx; + const auto thread_rank_in_block = thread_rank_in_grid % threads_in_block_count_; + thread_idx.x = thread_rank_in_block % block_dim_.x; + thread_idx.y = (thread_rank_in_block / block_dim_.x) % block_dim_.y; + thread_idx.z = thread_rank_in_block / (block_dim_.x * block_dim_.y); + + return thread_idx; + } + + dim3 grid_dim_; + dim3 block_dim_; + unsigned int block_count_; + unsigned int threads_in_block_count_; + unsigned int thread_count_; +}; + +inline dim3 GenerateThreadDimensions() { + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + const auto multipliers = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9, 1.0, 1.1, 1.2, 1.3, + 1.4, 1.5, 1.6, 1.7, 1.8, 1.9, 2.0, 2.1, 2.2, 2.3, 2.4, 2.5}; + return GENERATE_COPY( + dim3(1, 1, 1), dim3(props.maxThreadsDim[0], 1, 1), dim3(1, props.maxThreadsDim[1], 1), + dim3(1, 1, props.maxThreadsDim[2]), + map([max = props.maxThreadsDim[0], warp_size = props.warpSize]( + double i) { return dim3(std::min(static_cast(i * warp_size), max), 1, 1); }, + values(multipliers)), + map([max = props.maxThreadsDim[1], warp_size = props.warpSize]( + double i) { return dim3(1, std::min(static_cast(i * warp_size), max), 1); }, + values(multipliers)), + map([max = props.maxThreadsDim[2], warp_size = props.warpSize]( + double i) { return dim3(1, 1, std::min(static_cast(i * warp_size), max)); }, + values(multipliers)), + dim3(16, 8, 8), dim3(32, 32, 1), dim3(64, 8, 2), dim3(16, 16, 3), dim3(props.warpSize - 1, 3, 3), + dim3(props.warpSize + 1, 3, 3)); +} + +inline dim3 GenerateBlockDimensions() { + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + const auto multipliers = {0.1, 0.5, 0.9, 1.0, 1.1, 1.5, 1.9, 2.0, 3.0, 4.0}; + return GENERATE_COPY(dim3(1, 1, 1), + map([sm = props.multiProcessorCount]( + double i) { return dim3(static_cast(i * sm), 1, 1); }, + values(multipliers)), + map([sm = props.multiProcessorCount]( + double i) { return dim3(1, static_cast(i * sm), 1); }, + values(multipliers)), + map([sm = props.multiProcessorCount]( + double i) { return dim3(1, 1, static_cast(i * sm)); }, + values(multipliers)), + dim3(5, 5, 5)); +} + +inline dim3 GenerateThreadDimensionsForShuffle() { + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + const auto multipliers = {0.5, 0.9, 1.0, 1.5, 2.0}; + return GENERATE_COPY( + dim3(1, 1, 1), dim3(props.maxThreadsDim[0], 1, 1), dim3(1, props.maxThreadsDim[1], 1), + dim3(1, 1, props.maxThreadsDim[2]), + map([max = props.maxThreadsDim[0], warp_size = props.warpSize]( + double i) { return dim3(std::min(static_cast(i * warp_size), max), 1, 1); }, + values(multipliers)), + map([max = props.maxThreadsDim[1], warp_size = props.warpSize]( + double i) { return dim3(1, std::min(static_cast(i * warp_size), max), 1); }, + values(multipliers)), + map([max = props.maxThreadsDim[2], warp_size = props.warpSize]( + double i) { return dim3(1, 1, std::min(static_cast(i * warp_size), max)); }, + values(multipliers)), + dim3(16, 8, 8), dim3(32, 32, 1), dim3(64, 8, 2), dim3(16, 16, 3), dim3(props.warpSize - 1, 3, 3), + dim3(props.warpSize + 1, 3, 3)); +} + +inline dim3 GenerateBlockDimensionsForShuffle() { + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + const auto multipliers = {0.5, 1.0}; + return GENERATE_COPY(dim3(1, 1, 1), + map([sm = props.multiProcessorCount]( + double i) { return dim3(static_cast(i * sm), 1, 1); }, + values(multipliers)), + map([sm = props.multiProcessorCount]( + double i) { return dim3(1, static_cast(i * sm), 1); }, + values(multipliers)), + map([sm = props.multiProcessorCount]( + double i) { return dim3(1, 1, static_cast(i * sm)); }, + values(multipliers)), + dim3(5, 5, 5)); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 590f680880..f0966ae93b 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -128,3 +128,10 @@ THE SOFTWARE. * This section describes the various kernel functions invocation. * @} */ + +/** + * @defgroup DeviceLanguageTest Device Language + * @{ + * This section describes tests for the Device Language API. + * @} + */ diff --git a/projects/hip-tests/catch/include/utils.hh b/projects/hip-tests/catch/include/utils.hh index 98f6676d6d..6dd8070b9c 100644 --- a/projects/hip-tests/catch/include/utils.hh +++ b/projects/hip-tests/catch/include/utils.hh @@ -20,6 +20,7 @@ THE SOFTWARE. #pragma once #include +#include #include #include @@ -54,6 +55,20 @@ void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_ele ArrayFindIfNot(array, array + num_elements, expected_value); } +template +static inline void ArrayAllOf(const T* arr, uint32_t count, F value_gen) { + for (auto i = 0u; i < count; ++i) { + const std::optional expected_val = value_gen(i); + if (!expected_val.has_value()) continue; + // Using require on every iteration leads to a noticeable performance loss on large arrays, + // even when the require passes. + if (arr[i] != expected_val.value()) { + INFO("Mismatch at index: " << i); + REQUIRE(arr[i] == expected_val.value()); + } + } +} + template void PitchedMemoryVerify(T* const ptr, const size_t pitch, const size_t width, const size_t height, const size_t depth, F expected_value_generator) { diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt index ce7632fb8b..ad31c0a22a 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt @@ -6,6 +6,7 @@ set(TEST_SRC hipCGMultiGridGroupType.cc hipCGMultiGridGroupTypeViaBaseType.cc hipCGMultiGridGroupTypeViaPublicApi.cc + grid_group.cc coalesced_groups_shfl_down.cc coalesced_groups_shfl_up.cc simple_coalesced_groups.cc diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh b/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh new file mode 100644 index 0000000000..b0c8598581 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh @@ -0,0 +1,71 @@ +/* +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 + +#include +#include + +namespace { +#if (!__GFX8__ && !__GFX9__) || HT_NVIDIA +constexpr size_t kWarpSize = 32; +#else +constexpr size_t kWarpSize = 64; +#endif +} // namespace + +constexpr int MaxGPUs = 8; + +__device__ inline unsigned int thread_rank_in_grid() { + const auto block_size = blockDim.x * blockDim.y * blockDim.z; + const auto block_rank_in_grid = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x; + const auto thread_rank_in_block = + (threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x; + return block_rank_in_grid * block_size + thread_rank_in_block; +} + +static __device__ void busy_wait(unsigned long long wait_period) { + unsigned long long time_diff = 0; + unsigned long long last_clock = clock64(); + while (time_diff < wait_period) { + unsigned long long cur_clock = clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + last_clock = cur_clock; + } +} + +template bool CheckDimensions(unsigned int device, T kernel, dim3 blocks, dim3 threads) { + hipDeviceProp_t props; + int max_blocks_per_sm = 0; + int num_sm = 0; + HIP_CHECK(hipSetDevice(device)); + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, kernel, + threads.x * threads.y * threads.z, 0)); + + HIP_CHECK(hipGetDeviceProperties(&props, device)); + num_sm = props.multiProcessorCount; + + if ((blocks.x * blocks.y * blocks.z) > max_blocks_per_sm * num_sm) { + return false; + } + + return true; +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/grid_group.cc b/projects/hip-tests/catch/unit/cooperativeGrps/grid_group.cc new file mode 100644 index 0000000000..ee7402341a --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/grid_group.cc @@ -0,0 +1,285 @@ +/* +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 "cooperative_groups_common.hh" + +#include +#include +#include + +/** + * @addtogroup grid_group grid_group + * @{ + * @ingroup DeviceLanguageTest + * Contains unit tests for all grid_group APIs + */ + +namespace cg = cooperative_groups; + +static __global__ void grid_group_size_getter(unsigned int* sizes) { + sizes[thread_rank_in_grid()] = cg::this_grid().size(); +} + +static __global__ void grid_group_thread_rank_getter(unsigned int* thread_ranks) { + thread_ranks[thread_rank_in_grid()] = cg::this_grid().thread_rank(); +} + +static __global__ void grid_group_is_valid_getter(unsigned int* is_valid_flags) { + is_valid_flags[thread_rank_in_grid()] = cg::this_grid().is_valid(); +} + +static __global__ void grid_group_non_member_size_getter(unsigned int* sizes) { + sizes[thread_rank_in_grid()] = cg::group_size(cg::this_grid()); +} + +static __global__ void grid_group_non_member_thread_rank_getter(unsigned int* thread_ranks) { + thread_ranks[thread_rank_in_grid()] = cg::thread_rank(cg::this_grid()); +} + +static __global__ void sync_kernel(unsigned int* atomic_val, unsigned int* array, + unsigned int loops) { + cg::grid_group grid = cg::this_grid(); + unsigned rank = grid.thread_rank(); + + int offset = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x; + for (int i = 0; i < loops; i++) { + // Make the last thread run way behind everyone else. + // If the sync below fails, then the other threads may hit the + // atomicInc instruction many times before the last thread ever gets to it. + // If the sync works, then it will likely contain "total number of blocks"*i + if (rank == (grid.size() - 1)) { + busy_wait(100000); + } + if (threadIdx.x == blockDim.x - 1 && threadIdx.y == blockDim.y - 1 && + threadIdx.z == blockDim.z - 1) { + array[offset] = atomicInc(&atomic_val[0], UINT_MAX); + } + grid.sync(); + offset += gridDim.x * gridDim.y * gridDim.z; + } +} + +/** + * Test Description + * ------------------------ + * - Launches kernels that write the return values of size, thread_rank and is_valid member + * functions to an output array that is validated on the host side. The kernels are run + * sequentially, reusing the output array, to avoid running out of device memory for large kernel + * launches. + * Test source + * ------------------------ + * - unit/cooperativeGrps/grid_group.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Device supports cooperative launch + */ +TEST_CASE("Unit_Grid_Group_Getters_Positive_Basic") { + int device; + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + if (!device_properties.cooperativeLaunch) { + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; + } + + const auto blocks = GenerateBlockDimensions(); + const auto threads = GenerateThreadDimensions(); + if (!CheckDimensions(device, grid_group_size_getter, blocks, threads)) return; + INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z); + INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z); + + const CPUGrid grid(blocks, threads); + + LinearAllocGuard uint_arr_dev(LinearAllocs::hipMalloc, + grid.thread_count_ * sizeof(unsigned int)); + LinearAllocGuard uint_arr(LinearAllocs::hipHostMalloc, + grid.thread_count_ * sizeof(unsigned int)); + + // Launch Kernel + unsigned int* uint_arr_dev_ptr = uint_arr_dev.ptr(); + void* params[1]; + params[0] = &uint_arr_dev_ptr; + + HIP_CHECK(hipLaunchCooperativeKernel(grid_group_size_getter, blocks, threads, params, 0, 0)); + + HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(), + grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK( + hipLaunchCooperativeKernel(grid_group_thread_rank_getter, blocks, threads, params, 0, 0)); + + // Verify grid_group.size() values + ArrayAllOf(uint_arr.ptr(), grid.thread_count_, + [size = grid.thread_count_](uint32_t) { return size; }); + + HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(), + grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipLaunchCooperativeKernel(grid_group_is_valid_getter, blocks, threads, params, 0, 0)); + + // Verify grid_group.thread_rank() values + ArrayAllOf(uint_arr.ptr(), grid.thread_count_, [](uint32_t i) { return i; }); + + HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(), + grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + // Verify grid_group.is_valid() values + ArrayAllOf(uint_arr.ptr(), grid.thread_count_, [](uint32_t i) { return 1; }); +} + +/** + * Test Description + * ------------------------ + * - Launches kernels that write the return values of size and thread_rank non-member functions + * to an output array that is validated on the host side. The kernels are run sequentially, reusing + * the output array, to avoid running out of device memory for large kernel launches. + * Test source + * ------------------------ + * - unit/cooperativeGrps/grid_group.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Device supports cooperative launch + */ +TEST_CASE("Unit_Grid_Group_Getters_Via_Non_Member_Functions_Positive_Basic") { + int device; + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + if (!device_properties.cooperativeLaunch) { + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; + } + + const auto blocks = GenerateBlockDimensions(); + const auto threads = GenerateThreadDimensions(); + if (!CheckDimensions(device, grid_group_non_member_size_getter, blocks, threads)) return; + INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z); + INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z); + + const CPUGrid grid(blocks, threads); + + LinearAllocGuard uint_arr_dev(LinearAllocs::hipMalloc, + grid.thread_count_ * sizeof(unsigned int)); + LinearAllocGuard uint_arr(LinearAllocs::hipHostMalloc, + grid.thread_count_ * sizeof(unsigned int)); + + // Launch Kernel + unsigned int* uint_arr_dev_ptr = uint_arr_dev.ptr(); + void* params[1]; + params[0] = &uint_arr_dev_ptr; + + HIP_CHECK( + hipLaunchCooperativeKernel(grid_group_non_member_size_getter, blocks, threads, params, 0, 0)); + + HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(), + grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipLaunchCooperativeKernel(grid_group_non_member_thread_rank_getter, blocks, threads, + params, 0, 0)); + + // Verify grid_group.size() values + ArrayAllOf(uint_arr.ptr(), grid.thread_count_, + [size = grid.thread_count_](uint32_t) { return size; }); + + HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(), + grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + // Verify grid_group.thread_rank() values + ArrayAllOf(uint_arr.ptr(), grid.thread_count_, [](uint32_t i) { return i; }); +} + +/** + * Test Description + * ------------------------ + * - Launches a kernel where the last thread in a block atomically increments a global variable + * within a work loop. The value returned from this atomic increment entirely depends on the order + * the threads arrive at the atomic instruction. Each thread then stores the result in the global + * array based on its block id. A wait loop is inserted into the last thread so that it runs behind + * all other threads. If the sync doesn't work, the other threads will increment the atomic variable + * many times before the last thread gets to it and it will read a very large value. If the sync + * works, each thread will increment the variable once per loop iteration and the last thread will + * contain total number of blocks * loop iteration. + * Test source + * ------------------------ + * - unit/cooperativeGrps/grid_group.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Device supports cooperative launch + */ +TEST_CASE("Unit_Grid_Group_Sync_Positive_Basic") { + int device; + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + if (!device_properties.cooperativeLaunch) { + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; + } + + auto loops = GENERATE(2, 4, 8, 16); + const auto blocks = GenerateBlockDimensions(); + const auto threads = GenerateThreadDimensions(); + if (!CheckDimensions(device, sync_kernel, blocks, threads)) return; + INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z); + INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z); + + const CPUGrid grid(blocks, threads); + unsigned int array_len = grid.block_count_ * loops; + + LinearAllocGuard uint_arr_dev(LinearAllocs::hipMalloc, + array_len * sizeof(unsigned int)); + LinearAllocGuard uint_arr(LinearAllocs::hipHostMalloc, + array_len * sizeof(unsigned int)); + LinearAllocGuard atomic_val(LinearAllocs::hipMalloc, sizeof(unsigned int)); + HIP_CHECK(hipMemset(atomic_val.ptr(), 0, sizeof(unsigned int))); + + // Launch Kernel + unsigned int* uint_arr_dev_ptr = uint_arr_dev.ptr(); + unsigned int* atomic_val_ptr = atomic_val.ptr(); + void* params[3]; + params[0] = reinterpret_cast(&atomic_val_ptr); + params[1] = reinterpret_cast(&uint_arr_dev_ptr); + params[2] = reinterpret_cast(&loops); + + HIP_CHECK(hipLaunchCooperativeKernel(sync_kernel, blocks, threads, params, 0, 0)); + + HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(), array_len * sizeof(*uint_arr.ptr()), + hipMemcpyDeviceToHost)); + + HIP_CHECK(hipDeviceSynchronize()); + + // Verify host buffer values + unsigned int max_in_this_loop = 0; + for (unsigned int i = 0; i < loops; i++) { + max_in_this_loop += grid.block_count_; + unsigned int j = 0; + for (j = 0; j < grid.block_count_ - 1; j++) { + REQUIRE(uint_arr.ptr()[i * grid.block_count_ + j] < max_in_this_loop); + } + REQUIRE(uint_arr.ptr()[i * grid.block_count_ + j] == max_in_this_loop - 1); + } +} \ No newline at end of file From 7dbb3a94ae9a2eefbd4e585f66085902aa5b24be Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Tue, 18 Jul 2023 00:29:03 -0700 Subject: [PATCH 03/17] SWDEV-396963 - skip Unit_hipMemcpy2DFromArrayAsync_Positive_Synchronization_Behavior (#243) Co-authored-by: Rakesh Roy <137397847+rakesroy@users.noreply.github.com> [ROCm/hip-tests commit: 26bc233996df0faefb5ee074710a1827d5e2b067] --- .../catch/hipTestMain/config/config_amd_linux_common.json | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index 445b274c88..092093cb9e 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -63,6 +63,8 @@ "Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag", "Disabling test tracked SWDEV-395683", "Unit_hipStreamPerThread_MultiThread", + "SWDEV-396963", + "Unit_hipMemcpy2DFromArrayAsync_Positive_Synchronization_Behavior", "Disabling tests tracked with SWDEV-389647..", "Unit_hipMemcpy2DToArrayAsync_Positive_Synchronization_Behavior", "Disabling test tracked SWDEV-391555", From 0d73b9c81f8ee1beae59f2628c54bf65c52848e0 Mon Sep 17 00:00:00 2001 From: Ioannis Assiouras <38722728+iassiour@users.noreply.github.com> Date: Tue, 18 Jul 2023 08:29:37 +0100 Subject: [PATCH 04/17] SWDEV-398646 - Skip Unit_hipStreamAddCallback_StrmSyncTiming (#261) Unit_hipStreamAddCallback_StrmSyncTiming fails on windows because of timing issues. [ROCm/hip-tests commit: b9fde9f27c027df4bcea6991798b15de7c548958] --- .../catch/hipTestMain/config/config_amd_windows_common.json | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json index 80401608ef..7471c65928 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -180,6 +180,7 @@ "Unit_hipMemcpyPeerAsync_Positive_ZeroSize", "Disabling test tracked SWDEV-391718", "Unit_hipMemRangeGetAttribute_TstCountParam", + "Unit_hipStreamAddCallback_StrmSyncTiming", "SWDEV-402082 - PAL Backend fails to reserve address on GPU except first one", "Unit_hipGraphInstantiateWithFlags_FlagAutoFreeOnLaunch_check", "SWDEV-398981 fails in stress test", From d660c0a84d9ae8d63e670f66787c0e608f2d3b83 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Tue, 18 Jul 2023 00:30:29 -0700 Subject: [PATCH 05/17] =?UTF-8?q?SWDEV-400049=20-=20disable=20Unit=5FhipMe?= =?UTF-8?q?msetDSync=20=E2=80=93=20init16=5Ft=20(#264)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit tdr in windows so disabling the test [ROCm/hip-tests commit: 0be7e3d12d56672ae9d5d21f468fa8ea721b1910] --- .../catch/hipTestMain/config/config_amd_windows_common.json | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json index 7471c65928..5d661665af 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -180,6 +180,8 @@ "Unit_hipMemcpyPeerAsync_Positive_ZeroSize", "Disabling test tracked SWDEV-391718", "Unit_hipMemRangeGetAttribute_TstCountParam", + "SWDEV-400049 tdr intermittently", + "Unit_hipMemsetDSync – init16_t", "Unit_hipStreamAddCallback_StrmSyncTiming", "SWDEV-402082 - PAL Backend fails to reserve address on GPU except first one", "Unit_hipGraphInstantiateWithFlags_FlagAutoFreeOnLaunch_check", From 64167b64c48a2cc611849086dd4d895a6839befd Mon Sep 17 00:00:00 2001 From: nives-vukovic <110852104+nives-vukovic@users.noreply.github.com> Date: Tue, 18 Jul 2023 09:30:51 +0200 Subject: [PATCH 06/17] EXSWHTEC-304 - Expand and refactor tests for clock device functions (#273) - Expand and refactor clock device functions - Add comments and format code - Fix doxygen test group comments [ROCm/hip-tests commit: f73d1b21ff4f6c310a3cd6ba3d5ad52a67ce568c] --- .../catch/include/hip_test_defgroups.hh | 6 + projects/hip-tests/catch/unit/CMakeLists.txt | 2 +- .../catch/unit/clock/hipClockCheck.cc | 209 ++++++++++++------ 3 files changed, 149 insertions(+), 68 deletions(-) diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index f0966ae93b..402625006a 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -135,3 +135,9 @@ THE SOFTWARE. * This section describes tests for the Device Language API. * @} */ + +/** + * @defgroup DeviceLanguageTest Device Language + * @{ + * This section describes tests for the Device Language API. + */ diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index e0b1f5a1b2..50e2ca0f58 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -44,7 +44,7 @@ add_subdirectory(executionControl) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(callback) -#add_subdirectory(clock) +add_subdirectory(clock) # Vulkan interop APIs currently undefined for Nvidia add_subdirectory(vulkan_interop) endif() diff --git a/projects/hip-tests/catch/unit/clock/hipClockCheck.cc b/projects/hip-tests/catch/unit/clock/hipClockCheck.cc index 4d817a3764..148ba143c8 100644 --- a/projects/hip-tests/catch/unit/clock/hipClockCheck.cc +++ b/projects/hip-tests/catch/unit/clock/hipClockCheck.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -21,40 +21,51 @@ THE SOFTWARE. #include #include -#define ONESECOND 1000 // in ms -#define HALFSECOND 500 // in ms +/** + * @addtogroup clock clock + * @{ + * @ingroup DeviceLanguageTest + * Contains unit tests for clock, clock64 and wall_clock64 APIs + */ -enum CLOCK_MODE { - CLOCK_MODE_CLOCK64, - CLOCK_MODE_WALL_CLOCK64 -}; - -__global__ void kernel_c(int clockRate, uint64_t wait_t) { - uint64_t start = clock64() / clockRate, cur = 0; // in ms - do { cur = clock64() / clockRate-start;} while (cur < wait_t); +__global__ void kernel_c64(int clock_rate, uint64_t wait_t) { + uint64_t start = clock64() / clock_rate, cur = 0; // in ms + do { + cur = clock64() / clock_rate - start; + } while (cur < wait_t); } -__global__ void kernel_w(int clockRate, uint64_t wait_t) { - uint64_t start = wall_clock64() / clockRate, cur = 0; // in ms - do { cur = wall_clock64() / clockRate-start;} while (cur < wait_t); +__global__ void kernel_c(int clock_rate, uint64_t wait_t) { + uint64_t start = clock() / clock_rate, cur = 0; // in ms + do { + cur = clock() / clock_rate - start; + } while (cur < wait_t); } -bool verifyTimeExecution(CLOCK_MODE m, float time1, float time2, - float expectedTime1, float expectedTime2) { - bool testStatus = false; - float ratio = m == CLOCK_MODE_CLOCK64 ? 0.5 : 0.01; +__global__ void kernel_wc64(int clock_rate, uint64_t wait_t) { + uint64_t start = wall_clock64() / clock_rate, cur = 0; // in ms + do { + cur = wall_clock64() / clock_rate - start; + } while (cur < wait_t); +} - if (fabs(time1 - expectedTime1) < ratio * expectedTime1 - && fabs(time2 - expectedTime2) < ratio * expectedTime2) { - WARN("Succeeded: Expected Vs Actual: Kernel1 - " << expectedTime1 << " Vs " << time1 - << ", Kernel2 - " << expectedTime2 << " Vs " << time2); - testStatus = true; +bool verify_time_execution(float ratio, float time1, float time2, float expected_time1, + float expected_time2) { + bool test_status = false; + + if (fabs(time1 - expected_time1) < ratio * expected_time1 && + fabs(time2 - expected_time2) < ratio * expected_time2) { + INFO("Succeeded: Expected Vs Actual: Kernel1 - " << expected_time1 << " Vs " << time1 + << ", Kernel2 - " << expected_time2 << " Vs " + << time2); + test_status = true; } else { - FAIL_CHECK("Failed: Expected Vs Actual: Kernel1 -" << expectedTime1 << " Vs " << time1 - << ", Kernel2 - " << expectedTime2 << " Vs " << time2); - testStatus = false; + INFO("Failed: Expected Vs Actual: Kernel1 -" << expected_time1 << " Vs " << time1 + << ", Kernel2 - " << expected_time2 << " Vs " + << time2); + test_status = false; } - return testStatus; + return test_status; } /* @@ -62,55 +73,119 @@ bool verifyTimeExecution(CLOCK_MODE m, float time1, float time2, * get the event elapsed time of each kernel using the start and * end events.The event elapsed time should return us the kernel * execution time for that particular kernel -*/ -bool kernelTimeExecution(CLOCK_MODE m, int clockRate, - uint64_t expectedTime1, uint64_t expectedTime2) { + */ +bool kernel_time_execution(void (*kernel)(int, uint64_t), int clock_rate, uint64_t expected_time1, + uint64_t expected_time2) { hipStream_t stream; hipEvent_t start_event1, end_event1, start_event2, end_event2; float time1 = 0, time2 = 0; - HIPCHECK(hipEventCreate(&start_event1)); - HIPCHECK(hipEventCreate(&end_event1)); - HIPCHECK(hipEventCreate(&start_event2)); - HIPCHECK(hipEventCreate(&end_event2)); - HIPCHECK(hipStreamCreate(&stream)); - hipExtLaunchKernelGGL( m == CLOCK_MODE_CLOCK64 ? kernel_c : kernel_w, - dim3(1), dim3(1), 0, stream, start_event1, end_event1, 0, clockRate, expectedTime1); - hipExtLaunchKernelGGL( m == CLOCK_MODE_CLOCK64 ? kernel_c : kernel_w, - dim3(1), dim3(1), 0, stream, start_event2, end_event2, 0, clockRate, expectedTime2); - HIPCHECK(hipStreamSynchronize(stream)); - HIPCHECK(hipEventElapsedTime(&time1, start_event1, end_event1)); - HIPCHECK(hipEventElapsedTime(&time2, start_event2, end_event2)); + HIP_CHECK(hipEventCreate(&start_event1)); + HIP_CHECK(hipEventCreate(&end_event1)); + HIP_CHECK(hipEventCreate(&start_event2)); + HIP_CHECK(hipEventCreate(&end_event2)); + HIP_CHECK(hipStreamCreate(&stream)); + hipExtLaunchKernelGGL(kernel, dim3(1), dim3(1), 0, stream, start_event1, end_event1, 0, + clock_rate, expected_time1); + hipExtLaunchKernelGGL(kernel, dim3(1), dim3(1), 0, stream, start_event2, end_event2, 0, + clock_rate, expected_time2); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipEventElapsedTime(&time1, start_event1, end_event1)); + HIP_CHECK(hipEventElapsedTime(&time2, start_event2, end_event2)); - HIPCHECK(hipStreamDestroy(stream)); - HIPCHECK(hipEventDestroy(start_event1)); - HIPCHECK(hipEventDestroy(end_event1)); - HIPCHECK(hipEventDestroy(start_event2)); - HIPCHECK(hipEventDestroy(end_event2)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipEventDestroy(start_event1)); + HIP_CHECK(hipEventDestroy(end_event1)); + HIP_CHECK(hipEventDestroy(start_event2)); + HIP_CHECK(hipEventDestroy(end_event2)); - return verifyTimeExecution(m, time1, time2, expectedTime1, expectedTime2); + float ratio = kernel == kernel_wc64 ? 0.01 : 0.5; + + return verify_time_execution(ratio, time1, time2, expected_time1, expected_time2); } -TEST_CASE("Unit_hipClock64_Check") { - HIPCHECK(hipSetDevice(0)); - int clockRate = 0; // in KHz - HIPCHECK(hipDeviceGetAttribute(&clockRate, hipDeviceAttributeClockRate, 0)); +/** + * Test Description + * ------------------------ + * - Launches two kernels that run for a specified amount of time passed as a kernel argument by + * using device function clock64. Kernel execution time is calculated through elapsed time between + * the start and end event, and calculated time is compared with passed time values. + * Test source + * ------------------------ + * - catch/unit/clock/hipClockCheck.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipClock64_Positive_Basic") { + HIP_CHECK(hipSetDevice(0)); + int clock_rate = 0; // in kHz + HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeClockRate, 0)); - SECTION("Verify kernel execution time via clock64()") { - CHECK(kernelTimeExecution(CLOCK_MODE_CLOCK64, clockRate, ONESECOND, HALFSECOND)); - } -} - -TEST_CASE("Unit_hipWallClock64_Check") { - HIPCHECK(hipSetDevice(0)); - int clockRate = 0; // in KHz - HIPCHECK(hipDeviceGetAttribute(&clockRate, hipDeviceAttributeWallClockRate, 0)); - - if(!clockRate) { - INFO("hipDeviceAttributeWallClockRate has not been supported. Skipped"); + if (IsGfx11()) { + HipTest::HIP_SKIP_TEST("Issue with clock64() function on gfx11 devices!"); return; } - SECTION("Verify kernel execution time via wall_clock64()") { - CHECK(kernelTimeExecution(CLOCK_MODE_WALL_CLOCK64, clockRate, ONESECOND, HALFSECOND)); - } + const auto expected_time1 = GENERATE(1000, 1500, 2000); + const auto expected_time2 = expected_time1 / 2; + + REQUIRE(kernel_time_execution(kernel_c64, clock_rate, expected_time1, expected_time2)); +} + +/** + * Test Description + * ------------------------ + * - Launches two kernels that run for a specified amount of time passed as a kernel argument by + * using device function clock. Kernel execution time is calculated through elapsed time between + * the start and end event, and calculated time is compared with passed time values. + * Test source + * ------------------------ + * - catch/unit/clock/hipClockCheck.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipClock_Positive_Basic") { + HIP_CHECK(hipSetDevice(0)); + int clock_rate = 0; // in kHz + HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeClockRate, 0)); + + if (IsGfx11()) { + HipTest::HIP_SKIP_TEST("Issue with clock() function on gfx11 devices!"); + return; + } + + const auto expected_time1 = GENERATE(1000, 1500, 2000); + const auto expected_time2 = expected_time1 / 2; + + REQUIRE(kernel_time_execution(kernel_c, clock_rate, expected_time1, expected_time2)); +} + +/** + * Test Description + * ------------------------ + * - Launches two kernels that run for a specified amount of time passed as a kernel argument by + * using device function wall_clock64. Kernel execution time is calculated through elapsed time + * between the start and end event, and calculated time is compared with passed time values. + * Test source + * ------------------------ + * - catch/unit/clock/hipClockCheck.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipWallClock64_Positive_Basic") { + HIP_CHECK(hipSetDevice(0)); + int clock_rate = 0; // in kHz + HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeWallClockRate, 0)); + + if (!clock_rate) { + HipTest::HIP_SKIP_TEST("hipDeviceAttributeWallClockRate is not supported"); + return; + } + + const auto expected_time1 = GENERATE(1000, 1500, 2000); + const auto expected_time2 = expected_time1 / 2; + + REQUIRE(kernel_time_execution(kernel_wc64, clock_rate, expected_time1, expected_time2)); } From 4dfba013ff212bc3bf26372e5b80cdec2490db53 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Tue, 18 Jul 2023 00:31:07 -0700 Subject: [PATCH 07/17] SWDEV-402054 - Disable Unit_hipEventDestroy_WithWaitingStream (#278) skip Unit_hipEventDestroy_WithWaitingStream in both linux/windows [ROCm/hip-tests commit: 3c11353740f99d7f149cfb52e6362dfa481d70be] --- .../catch/hipTestMain/config/config_amd_linux_common.json | 2 ++ .../catch/hipTestMain/config/config_amd_windows_common.json | 2 ++ 2 files changed, 4 insertions(+) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index 092093cb9e..1d2c9e458b 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -79,6 +79,8 @@ "Unit_hipMemset2DSync", "SWDEV-398981 fails in stress test", "Unit_hipStreamCreateWithPriority_MulthreadDefaultflag", + "SWDEV-402054 fails in external github build", + "Unit_hipEventDestroy_WithWaitingStream", "=== Below tests fail in stress test on 23/06/23 ===", "Unit_hipIpcMemAccess_ParameterValidation", "Unit_hipMemcpy2DFromArrayAsync_Positive_Synchronization_Behavior", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json index 5d661665af..15e35fef94 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -187,6 +187,8 @@ "Unit_hipGraphInstantiateWithFlags_FlagAutoFreeOnLaunch_check", "SWDEV-398981 fails in stress test", "Unit_hipStreamCreateWithPriority_MulthreadDefaultflag", + "SWDEV-402054 fails in external github build", + "Unit_hipEventDestroy_WithWaitingStream", "Note: UUID returned empty on some windows nodes", "Unit_hipDeviceGetUuid_Positive", "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/96 ===", From 8b44d2d773210f7b1f4934c6ee3142523cf4b9b5 Mon Sep 17 00:00:00 2001 From: milos-mozetic <118800401+milos-mozetic@users.noreply.github.com> Date: Tue, 18 Jul 2023 09:32:03 +0200 Subject: [PATCH 08/17] EXSWHTEC-315 - Implement unit tests for dim3 (#298) - Implement unit tests for dim3 - Add doxygen comments [ROCm/hip-tests commit: 8412bcb1b3d11c8854f1940347c8f6eaccfde178] --- .../catch/include/hip_test_defgroups.hh | 7 + projects/hip-tests/catch/unit/CMakeLists.txt | 1 + .../catch/unit/vector_types/CMakeLists.txt | 29 ++ .../hip-tests/catch/unit/vector_types/dim3.cc | 261 ++++++++++++++++++ 4 files changed, 298 insertions(+) create mode 100644 projects/hip-tests/catch/unit/vector_types/CMakeLists.txt create mode 100644 projects/hip-tests/catch/unit/vector_types/dim3.cc diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 402625006a..92857a12bc 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -36,6 +36,13 @@ THE SOFTWARE. * @} */ +/** + * @defgroup VectorTypeTest Vector types + * @{ + * This section describes tests for the Vector type functions and operators. + * @} + */ + /** * @defgroup DeviceTest Device Management * @{ diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index 50e2ca0f58..135d3a12ab 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -48,3 +48,4 @@ add_subdirectory(clock) # Vulkan interop APIs currently undefined for Nvidia add_subdirectory(vulkan_interop) endif() +add_subdirectory(vector_types) diff --git a/projects/hip-tests/catch/unit/vector_types/CMakeLists.txt b/projects/hip-tests/catch/unit/vector_types/CMakeLists.txt new file mode 100644 index 0000000000..49619275f3 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/CMakeLists.txt @@ -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. + +# Common Tests - Test independent of all platforms + +set(TEST_SRC + dim3.cc +) + +hip_add_exe_to_target(NAME VectorTypesTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) diff --git a/projects/hip-tests/catch/unit/vector_types/dim3.cc b/projects/hip-tests/catch/unit/vector_types/dim3.cc new file mode 100644 index 0000000000..2b1db91ded --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/dim3.cc @@ -0,0 +1,261 @@ +/* +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 + +/** + * @addtogroup dim3 dim3 + * @{ + * @ingroup VectorTypeTest + */ + +__global__ void Dim3VectorKernel(dim3* vector) { *vector = dim3(); } +__global__ void Dim3VectorKernel(dim3* vector, const uint32_t x) { *vector = dim3(x); } +__global__ void Dim3VectorKernel(dim3* vector, const uint32_t x, const uint32_t y) { + *vector = dim3(x, y); +} +__global__ void Dim3VectorKernel(dim3* vector, const uint32_t x, const uint32_t y, + const uint32_t z) { + *vector = dim3(x, y, z); +} + +/** + * Test Description + * ------------------------ + * - Creates a dim3 with an empty constructor: + * -# Expected result: dim3(1, 1, 1) + * - Calls dim3 from the device side + * Test source + * ------------------------ + * - unit/vector_types/dim3.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_dim3_Empty_Positive_Device") { + dim3 vector_h{0, 0, 0}; + dim3* vector_d; + HIP_CHECK(hipMalloc(&vector_d, sizeof(dim3))); + HIP_CHECK(hipMemcpy(vector_d, &vector_h, sizeof(dim3), hipMemcpyHostToDevice)); + Dim3VectorKernel<<<1, 1, 0, 0>>>(vector_d); + HIP_CHECK(hipMemcpy(&vector_h, vector_d, sizeof(dim3), hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(vector_d)); + + REQUIRE(vector_h.x == 1); + REQUIRE(vector_h.y == 1); + REQUIRE(vector_h.z == 1); +} + +/** + * Test Description + * ------------------------ + * - Creates a dim3 with an constructor with one parameter (X): + * -# Expected result: dim3(X, 1, 1) + * - Calls dim3 from the device side + * Test source + * ------------------------ + * - unit/vector_types/dim3.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_dim3_X_Positive_Device") { + dim3 vector_h{0, 0, 0}; + dim3* vector_d; + HIP_CHECK(hipMalloc(&vector_d, sizeof(dim3))); + HIP_CHECK(hipMemcpy(vector_d, &vector_h, sizeof(dim3), hipMemcpyHostToDevice)); + uint32_t value_x = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + Dim3VectorKernel<<<1, 1, 0, 0>>>(vector_d, value_x); + HIP_CHECK(hipMemcpy(&vector_h, vector_d, sizeof(dim3), hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(vector_d)); + + REQUIRE(vector_h.x == value_x); + REQUIRE(vector_h.y == 1); + REQUIRE(vector_h.z == 1); +} + +/** + * Test Description + * ------------------------ + * - Creates a dim3 with an constructor with two parameters (X, Y): + * -# Expected result: dim3(X, Y, 1) + * - Calls dim3 from the device side + * Test source + * ------------------------ + * - unit/vector_types/dim3.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_dim3_XY_Positive_Device") { + dim3 vector_h{0, 0, 0}; + dim3* vector_d; + HIP_CHECK(hipMalloc(&vector_d, sizeof(dim3))); + HIP_CHECK(hipMemcpy(vector_d, &vector_h, sizeof(dim3), hipMemcpyHostToDevice)); + uint32_t value_x = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + uint32_t value_y = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + Dim3VectorKernel<<<1, 1, 0, 0>>>(vector_d, value_x, value_y); + HIP_CHECK(hipMemcpy(&vector_h, vector_d, sizeof(dim3), hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(vector_d)); + + REQUIRE(vector_h.x == value_x); + REQUIRE(vector_h.y == value_y); + REQUIRE(vector_h.z == 1); +} + +/** + * Test Description + * ------------------------ + * - Creates a dim3 with an constructor with three parameters (X, Y, Z): + * -# Expected result: dim3(X, Y, Z) + * - Calls dim3 from the device side + * Test source + * ------------------------ + * - unit/vector_types/dim3.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_dim3_XYZ_Positive_Device") { + dim3 vector_h{0, 0, 0}; + dim3* vector_d; + HIP_CHECK(hipMalloc(&vector_d, sizeof(dim3))); + HIP_CHECK(hipMemcpy(vector_d, &vector_h, sizeof(dim3), hipMemcpyHostToDevice)); + uint32_t value_x = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + uint32_t value_y = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + + uint32_t value_z = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + Dim3VectorKernel<<<1, 1, 0, 0>>>(vector_d, value_x, value_y, value_z); + HIP_CHECK(hipMemcpy(&vector_h, vector_d, sizeof(dim3), hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(vector_d)); + + REQUIRE(vector_h.x == value_x); + REQUIRE(vector_h.y == value_y); + REQUIRE(vector_h.z == value_z); +} + +/** + * Test Description + * ------------------------ + * - Creates a dim3 with an empty constructor: + * -# Expected result: dim3(1, 1, 1) + * - Calls dim3 from the host side + * Test source + * ------------------------ + * - unit/vector_types/dim3.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_dim3_Empty_Positive_Host") { + dim3 vector = dim3(); + REQUIRE(vector.x == 1); + REQUIRE(vector.y == 1); + REQUIRE(vector.z == 1); +} + +/** + * Test Description + * ------------------------ + * - Creates a dim3 with an constructor with one parameter (X): + * -# Expected result: dim3(X, 1, 1) + * - Calls dim3 from the host side + * Test source + * ------------------------ + * - unit/vector_types/dim3.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_dim3_X_Positive_Host") { + uint32_t value_x = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + dim3 vector = dim3(value_x); + REQUIRE(vector.x == value_x); + REQUIRE(vector.y == 1); + REQUIRE(vector.z == 1); +} + +/** + * Test Description + * ------------------------ + * - Creates a dim3 with an constructor with two parameters (X, Y): + * -# Expected result: dim3(X, Y, 1) + * - Calls dim3 from the host side + * Test source + * ------------------------ + * - unit/vector_types/dim3.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_dim3_XY_Positive_Host") { + uint32_t value_x = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + uint32_t value_y = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + dim3 vector = dim3(value_x, value_y); + REQUIRE(vector.x == value_x); + REQUIRE(vector.y == value_y); + REQUIRE(vector.z == 1); +} + +/** + * Test Description + * ------------------------ + * - Creates a dim3 with an constructor with three parameters (X, Y, Z): + * -# Expected result: dim3(X, Y, Z) + * - Calls dim3 from the host side + * Test source + * ------------------------ + * - unit/vector_types/dim3.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_dim3_XYZ_Positive_Host") { + uint32_t value_x = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + uint32_t value_y = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + uint32_t value_z = + GENERATE(std::numeric_limits::min(), std::numeric_limits::max() / 2, + std::numeric_limits::max()); + dim3 vector = dim3(value_x, value_y, value_z); + REQUIRE(vector.x == value_x); + REQUIRE(vector.y == value_y); + REQUIRE(vector.z == value_z); +} From bd120fef5a516ab2bdee4e3de0138f6c53ab745c Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 18 Jul 2023 13:03:05 +0530 Subject: [PATCH 09/17] SWDEV-366390 - [catch2][dtest] Adding additional functional tests for hipGraphAddChildGraphNode() (#107) Change-Id: I132141db9d9fee9c8ed3a28f11b2c729042588f0 [ROCm/hip-tests commit: a23c6a140a2909d8405bfc93750182d1fcc2e903] --- .../config/config_amd_windows_MI2xx.json | 1 + .../config/config_amd_windows_common.json | 1 + .../unit/graph/hipGraphAddChildGraphNode.cc | 822 +++++++++++++++++- 3 files changed, 819 insertions(+), 5 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json index 3121abb125..7df6d38d78 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json @@ -93,6 +93,7 @@ "Unit_hipInit_Negative", "Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime", "Unit_hipStreamBeginCapture_captureComplexGraph", + "Unit_hipGraphAddChildGraphNode_MultGraphsAsSingleGraph" "Unit_hipMemGetAddressRange_Negative", "Unit_hipStreamValue_Wait64_Blocking_NoMask_Nor", "Unit_hipLaunchHostFunc_Graph", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json index 15e35fef94..372ae2bdca 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -95,6 +95,7 @@ "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread", "Note: intermittent Seg fault failure ", "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", + "Unit_hipGraphAddChildGraphNode_MultGraphsAsSingleGraph", "Unit_hipFuncSetCacheConfig_Positive_Basic", "Unit_hipFuncSetCacheConfig_Negative_Parameters", "Unit_hipFuncSetSharedMemConfig_Positive_Basic", diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc index a665180c2f..9c8bf4c098 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -22,10 +22,43 @@ Testcase Scenarios of hipGraphAddChildGraphNode API: Functional: 1. Create child graph as root node and execute the main graph. -2. Create multiple child graph nodes and check the behaviour -3. Clone the child graph node, Add new nodes and execute the cloned graph -4. Create child graph, add it to main graph and execute child graph -5. Pass original graph as child graph and execute the org graph +2. Create multiple child graph nodes and check the behaviour. +3. Clone the child graph node, Add new nodes and execute the cloned graph. +4. Create child graph, add it to main graph and execute child graph. +5. Pass original graph as child graph and execute the org graph. +6. This test case verifies nested graph functionality. Parent graph + containing child graph, which in turn, contains another child graph. + Execute the graph in loop taking random input data and Validate the + output in each iteration. +7. This test case verifies clones the nested graph created in scenario6. + Execute the cloned graph in loop taking random input data and Validate + the output in each iteration. +8. Verify if an empty graph can be added as child node. +9. Create the nested graph of scenario6 and update the property of add kernel + node (innermost graph) with subtract kernel functionality. Clone the graph. + Execute both the updated graph. +10. The updated nested graph in 9 is cloned and the cloned graph is then + executed and the result is validated. +11. Create the nested graph of 6 and update the block size and grid size + property of add kernel node. +12. Create the nested graph of 6 and delete the add kernel node + (innermost graph) and add a subtract kernel node. +13. The updated nested graph in 12 is cloned and the cloned graph is then + executed and the result is validated. +14. Create the nested graph of 6 and delete the add kernel node + (innermost graph), add a child graph that contains an event record node, + a subtract kernel node followed by another event record node. Clone the + graph. Execute both the original and cloned graph. +15. The updated nested graph in 14 is cloned and the cloned graph is then + executed and the result is validated. +16. Create one nested graph per GPU context. Execute all the created graphs + in their respective GPUs and validate the output. +17. Functional Test to use child node as barrier to wait for multiple nodes. + This test uses child nodes to resolve dependencies between graphs. 4 + graphs are created. Graph1 contains 3 independent memcpy h2d nodes, graph2 + contains 3 independent kernel nodes and graph3 contains 3 independent + memcpy d2h nodes. Graph1, graph2 and graph3 are added as child nodes in + graph4. Graph4 is validated for functionality. Negative: 1. Pass nullptr to graph node @@ -38,6 +71,7 @@ Negative: #include #include +#define TEST_LOOP_SIZE 50 /* This testcase verifies the negative scenarios of hipGraphAddChildGraphNode API @@ -422,3 +456,781 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_SingleChildNode") { HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(streamForGraph)); } + +// Kernel functions +static __global__ void ker_vec_mul(int *A, int *B, int *C) { + int i = threadIdx.x + blockDim.x * blockIdx.x; + C[i] = A[i]*B[i]; +} + +static __global__ void ker_vec_add(int *A, int *B) { + int i = threadIdx.x + blockDim.x * blockIdx.x; + A[i] = A[i] + B[i]; +} + +static __global__ void ker_vec_sub(int *A, int *B) { + int i = threadIdx.x + blockDim.x * blockIdx.x; + A[i] = A[i] - B[i]; +} + +static __global__ void ker_vec_sqr(int *A, int *B) { + int i = threadIdx.x + blockDim.x * blockIdx.x; + A[i] = B[i]*B[i]; +} + +enum class updateGraphNodeTests { + normalTest, + updateFunKerNodParamTest, + updateGrdBlkParamTest, + deleteAddNewKerNodTest, + addAnotherChildNodeTest +}; + +/** + Internal class for creating nested graphs. + */ +typedef class nestedGraph { + const int const_val1 = 11; + const int const_val2 = 7; + const int N = 1024; + size_t Nbytes; + const int threadsPerBlock = 256; + const int blocks = (N/threadsPerBlock); + const int threadsPerBlockUpd = 128; + const int blocksUpd = (N/threadsPerBlockUpd); + hipGraphNode_t memset_B1, memset_B2; + hipGraphNode_t memcpyH2D_A1, memcpyH2D_A2, memcpyD2H_A3; + hipGraphNode_t vec_mul1, vec_mul2, vec_add, vec_sqr, vec_sub; + hipGraphNode_t child_node1, child_node2, child_node3; + hipGraph_t graph[4]; // 4 level graph + hipKernelNodeParams kerNodeParams1{}, kerNodeParams2{}, + kerNodeParams3{}, kerNodeParams4{}; + int *A1_d, *A2_d, *A1_h, *A2_h, *A3_h; + int *B1_d, *B2_d, *C1_d, *C2_d; + hipMemsetParams memsetParams{}; + hipEvent_t eventstart, eventend; + hipGraphNode_t event_start, event_final; + + public: + // Create a nested Graph + nestedGraph() { + Nbytes = N * sizeof(int); + // Allocate device buffers + HIP_CHECK(hipMalloc(&A1_d, Nbytes)); + HIP_CHECK(hipMalloc(&A2_d, Nbytes)); + HIP_CHECK(hipMalloc(&B1_d, Nbytes)); + HIP_CHECK(hipMalloc(&B2_d, Nbytes)); + HIP_CHECK(hipMalloc(&C1_d, Nbytes)); + HIP_CHECK(hipMalloc(&C2_d, Nbytes)); + // Allocate host buffers + A1_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A1_h != nullptr); + A2_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A2_h != nullptr); + A3_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A3_h != nullptr); + // Create all the 3 level graphs + HIP_CHECK(hipGraphCreate(&graph[0], 0)); + HIP_CHECK(hipGraphCreate(&graph[1], 0)); + HIP_CHECK(hipGraphCreate(&graph[2], 0)); + HIP_CHECK(hipGraphCreate(&graph[3], 0)); + // Add the nodes to lowest level graph[2] + void* kernelArgs1[] = {&A1_d, &B1_d, &C1_d}; + kerNodeParams1.func = + reinterpret_cast(ker_vec_mul); + kerNodeParams1.gridDim = dim3(blocks); + kerNodeParams1.blockDim = dim3(threadsPerBlock); + kerNodeParams1.sharedMemBytes = 0; + kerNodeParams1.kernelParams = reinterpret_cast(kernelArgs1); + kerNodeParams1.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vec_mul1, graph[2], nullptr, 0, + &kerNodeParams1)); + void* kernelArgs2[] = {&A2_d, &B2_d, &C2_d}; + kerNodeParams2.func = + reinterpret_cast(ker_vec_mul); + kerNodeParams2.gridDim = dim3(blocks); + kerNodeParams2.blockDim = dim3(threadsPerBlock); + kerNodeParams2.sharedMemBytes = 0; + kerNodeParams2.kernelParams = reinterpret_cast(kernelArgs2); + kerNodeParams2.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vec_mul2, graph[2], nullptr, 0, + &kerNodeParams2)); + void* kernelArgs3[] = {&C1_d, &C2_d}; + kerNodeParams3.func = + reinterpret_cast(ker_vec_add); + kerNodeParams3.gridDim = dim3(blocks); + kerNodeParams3.blockDim = dim3(threadsPerBlock); + kerNodeParams3.sharedMemBytes = 0; + kerNodeParams3.kernelParams = reinterpret_cast(kernelArgs3); + kerNodeParams3.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vec_add, graph[2], nullptr, 0, + &kerNodeParams3)); + // Resolve Dependencies in graph[2] + HIP_CHECK(hipGraphAddDependencies(graph[2], &vec_mul1, &vec_add, 1)); + HIP_CHECK(hipGraphAddDependencies(graph[2], &vec_mul2, &vec_add, 1)); + // Add nodes to graph[1] + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B1_d); + memsetParams.value = const_val1; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(int); + memsetParams.width = N; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_B1, graph[1], nullptr, 0, + &memsetParams)); + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B2_d); + memsetParams.value = const_val2; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(int); + memsetParams.width = N; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_B2, graph[1], nullptr, 0, + &memsetParams)); + HIP_CHECK(hipGraphAddChildGraphNode(&child_node1, graph[1], + nullptr, 0, graph[2])); + void* kernelArgs4[] = {&C1_d, &C1_d}; + kerNodeParams3.func = + reinterpret_cast(ker_vec_sqr); + kerNodeParams3.gridDim = dim3(blocks); + kerNodeParams3.blockDim = dim3(threadsPerBlock); + kerNodeParams3.sharedMemBytes = 0; + kerNodeParams3.kernelParams = reinterpret_cast(kernelArgs4); + kerNodeParams3.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vec_sqr, graph[1], nullptr, 0, + &kerNodeParams3)); + HIP_CHECK(hipGraphAddDependencies(graph[1], &memset_B1, &child_node1, 1)); + HIP_CHECK(hipGraphAddDependencies(graph[1], &memset_B2, &child_node1, 1)); + HIP_CHECK(hipGraphAddDependencies(graph[1], &child_node1, &vec_sqr, 1)); + // Add nodes to graph[0] + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A1, graph[0], nullptr, + 0, A1_d, A1_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A2, graph[0], nullptr, + 0, A2_d, A2_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_A3, graph[0], nullptr, + 0, A3_h, C1_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddChildGraphNode(&child_node2, graph[0], + nullptr, 0, graph[1])); + HIP_CHECK(hipGraphAddDependencies(graph[0], &memcpyH2D_A1, + &child_node2, 1)); + HIP_CHECK(hipGraphAddDependencies(graph[0], &memcpyH2D_A2, + &child_node2, 1)); + HIP_CHECK(hipGraphAddDependencies(graph[0], &child_node2, + &memcpyD2H_A3, 1)); + } + // Fill Random Input Data + void fillRandInpData() { + unsigned int seed = time(nullptr); + for (int i = 0; i < N; i++) { + A1_h[i] = (HipTest::RAND_R(&seed) & 0xFF); + A2_h[i] = (HipTest::RAND_R(&seed) & 0xFF); + } + } + // Get the root graph + hipGraph_t* getRootGraph() { + return &graph[0]; + } + // Get the root graph + void updateInnermostNode(updateGraphNodeTests updatetype) { + hipGraph_t embGraph1, embGraph2; + // Get the embedded graph from child_node2 + HIP_CHECK(hipGraphChildGraphNodeGetGraph(child_node2, &embGraph2)); + size_t numNodes{}; + HIP_CHECK(hipGraphGetNodes(embGraph2, nullptr, &numNodes)); + hipGraphNode_t* nodes = + reinterpret_cast( + malloc(numNodes*sizeof(hipGraphNode_t))); + HIP_CHECK(hipGraphGetNodes(embGraph2, nodes, &numNodes)); + // Get the Graph node from the embedded graph + size_t nodeIdx = 0; + for (size_t idx = 0; idx < numNodes; idx++) { + hipGraphNodeType nodeType; + HIP_CHECK(hipGraphNodeGetType(nodes[idx], &nodeType)); + if (nodeType == hipGraphNodeTypeGraph) { + nodeIdx = idx; + break; + } + } + // Extract the embedded graph from the graph node + HIP_CHECK(hipGraphChildGraphNodeGetGraph(nodes[nodeIdx], &embGraph1)); + free(nodes); + numNodes = 0; + HIP_CHECK(hipGraphGetNodes(embGraph1, nullptr, &numNodes)); + nodes = reinterpret_cast( + malloc(numNodes*sizeof(hipGraphNode_t))); + // Get the kernel node from the extracted embedded graph + HIP_CHECK(hipGraphGetNodes(embGraph1, nodes, &numNodes)); + nodeIdx = 0; + hipKernelNodeParams nodeParam; + for (size_t idx = 0; idx < numNodes; idx++) { + hipGraphNodeType nodeType; + HIP_CHECK(hipGraphNodeGetType(nodes[idx], &nodeType)); + if (nodeType == hipGraphNodeTypeKernel) { + HIP_CHECK(hipGraphKernelNodeGetParams(nodes[idx], &nodeParam)); + if (nodeParam.func == reinterpret_cast(ker_vec_add)) { + nodeIdx = idx; + break; + } + } + } + if (updatetype == updateGraphNodeTests::updateFunKerNodParamTest) { + nodeParam.func = reinterpret_cast(ker_vec_sub); + HIP_CHECK(hipGraphKernelNodeSetParams(nodes[nodeIdx], &nodeParam)); + } else if (updatetype == updateGraphNodeTests::deleteAddNewKerNodTest) { + // delete the kernel add node + HIP_CHECK(hipGraphDestroyNode(nodes[nodeIdx])); + // add kernel subtract node to embGraph1 + void* kernelArgs[] = {&C1_d, &C2_d}; + kerNodeParams3.func = + reinterpret_cast(ker_vec_sub); + kerNodeParams3.gridDim = dim3(blocks); + kerNodeParams3.blockDim = dim3(threadsPerBlock); + kerNodeParams3.sharedMemBytes = 0; + kerNodeParams3.kernelParams = reinterpret_cast(kernelArgs); + kerNodeParams3.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vec_sub, embGraph1, nullptr, 0, + &kerNodeParams3)); + // Create new dependencies + for (size_t idx = 0; idx < numNodes; idx++) { + if (idx == nodeIdx) { + continue; + } + HIP_CHECK(hipGraphAddDependencies(embGraph1, &nodes[idx], + &vec_sub, 1)); + } + } else if (updatetype == updateGraphNodeTests::updateGrdBlkParamTest) { + nodeParam.blockDim = threadsPerBlockUpd; + nodeParam.gridDim = blocksUpd; + HIP_CHECK(hipGraphKernelNodeSetParams(nodes[nodeIdx], &nodeParam)); + } else if (updatetype == updateGraphNodeTests::addAnotherChildNodeTest) { + // delete the kernel add node + HIP_CHECK(hipGraphDestroyNode(nodes[nodeIdx])); + // add graph EventRecordNode -> Subtract Kernel -> EventRecordNode as + // child node + void* kernelArgs[] = {&C1_d, &C2_d}; + kerNodeParams3.func = + reinterpret_cast(ker_vec_sub); + kerNodeParams3.gridDim = dim3(blocks); + kerNodeParams3.blockDim = dim3(threadsPerBlock); + kerNodeParams3.sharedMemBytes = 0; + kerNodeParams3.kernelParams = reinterpret_cast(kernelArgs); + kerNodeParams3.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vec_sub, graph[3], nullptr, 0, + &kerNodeParams3)); + HIP_CHECK(hipEventCreate(&eventstart)); + HIP_CHECK(hipEventCreate(&eventend)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_start, graph[3], nullptr, + 0, eventstart)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_final, graph[3], nullptr, + 0, eventend)); + HIP_CHECK(hipGraphAddDependencies(graph[3], &event_start, + &vec_sub, 1)); + HIP_CHECK(hipGraphAddDependencies(graph[3], &vec_sub, + &event_final, 1)); + HIP_CHECK(hipGraphAddChildGraphNode(&child_node3, embGraph1, nullptr, + 0, graph[3])); + // Create new dependencies + for (size_t idx = 0; idx < numNodes; idx++) { + if (idx == nodeIdx) { + continue; + } + HIP_CHECK(hipGraphAddDependencies(embGraph1, &nodes[idx], + &child_node3, 1)); + } + } + free(nodes); + } + // Function to validate result + void validateOutData(updateGraphNodeTests updatetype) { + if ((updatetype == updateGraphNodeTests::normalTest) || + (updatetype == updateGraphNodeTests::updateGrdBlkParamTest)) { + for (int i = 0; i < N; i++) { + int result = (const_val1*A1_h[i] + const_val2*A2_h[i]); + result = result * result; + REQUIRE(result == A3_h[i]); + } + } else if ((updatetype == updateGraphNodeTests::deleteAddNewKerNodTest) + || (updatetype == updateGraphNodeTests::updateFunKerNodParamTest) + || (updatetype == updateGraphNodeTests::addAnotherChildNodeTest)) { + for (int i = 0; i < N; i++) { + int result = (const_val1*A1_h[i] - const_val2*A2_h[i]); + result = result * result; + REQUIRE(result == A3_h[i]); + } + } + } + // Destroy resources + ~nestedGraph() { + // Free all allocated buffers + HIP_CHECK(hipFree(C2_d)); + HIP_CHECK(hipFree(C1_d)); + HIP_CHECK(hipFree(B2_d)); + HIP_CHECK(hipFree(B1_d)); + HIP_CHECK(hipFree(A2_d)); + HIP_CHECK(hipFree(A1_d)); + free(A3_h); + free(A2_h); + free(A1_h); + HIP_CHECK(hipGraphDestroy(graph[3])); + HIP_CHECK(hipGraphDestroy(graph[2])); + HIP_CHECK(hipGraphDestroy(graph[1])); + HIP_CHECK(hipGraphDestroy(graph[0])); + } +} clNestedGraph; + +/** + Complex Scenario: This testcase verifies nested graph functionality. + Parent graph containing child graph, which in turn, contains another + child graph. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_Cmplx_NestedGraphs") { + hipGraph_t *graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + class nestedGraph nestedGraphObj; + graph = nestedGraphObj.getRootGraph(); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr, + nullptr, 0)); + for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) { + nestedGraphObj.fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + nestedGraphObj.validateOutData(updateGraphNodeTests::normalTest); + } + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); +} + +/** + Complex Scenario: This testcase verifies cloned nested graph functionality. + Clone the nested graph and execute the clone graph. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_CmplxClone_NestedGraphs") { + hipGraph_t *graph, clonedGraph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + class nestedGraph nestedGraphObj; + graph = nestedGraphObj.getRootGraph(); + HIP_CHECK(hipGraphClone(&clonedGraph, *graph)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedGraph, nullptr, + nullptr, 0)); + for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) { + nestedGraphObj.fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + nestedGraphObj.validateOutData(updateGraphNodeTests::normalTest); + } + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(clonedGraph)); +} + +/** + Scenario: Adding an empty graph to Child Graph Node. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_EmptyGraphAsChildNode") { + hipGraph_t graph, graphChild; + hipGraphNode_t child_node; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphCreate(&graphChild, 0)); + HIP_CHECK(hipGraphAddChildGraphNode(&child_node, graph, + nullptr, 0, graphChild)); + HIP_CHECK(hipGraphDestroy(graphChild)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + Complex Scenario: This testcase verifies the behavior of a nested graph + when one of the child graph node is updated. In this test the kernel node + function is updated to a different function. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_CmplxNstGrph_UpdKerFun") { + hipGraph_t *graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + class nestedGraph nestedGraphObj; + graph = nestedGraphObj.getRootGraph(); + nestedGraphObj.updateInnermostNode( + updateGraphNodeTests::updateFunKerNodParamTest); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr, + nullptr, 0)); + nestedGraphObj.fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + nestedGraphObj.validateOutData( + updateGraphNodeTests::updateFunKerNodParamTest); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); +} + +/** + Complex Scenario: This testcase verifies the behavior of a nested graph + when one of the child graph node is updated. In this test the kernel node + function is updated to a different function and the nested graph is cloned. + Execute the cloned graph and validate the results. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_CmplxNstGrph_UpdKerFun_Clone") { + hipGraph_t *graph, clonedGraph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + class nestedGraph nestedGraphObj; + graph = nestedGraphObj.getRootGraph(); + nestedGraphObj.updateInnermostNode( + updateGraphNodeTests::updateFunKerNodParamTest); + HIP_CHECK(hipGraphClone(&clonedGraph, *graph)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedGraph, nullptr, + nullptr, 0)); + nestedGraphObj.fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + nestedGraphObj.validateOutData( + updateGraphNodeTests::updateFunKerNodParamTest); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(clonedGraph)); +} + +/** + Complex Scenario: This testcase verifies the behavior of a nested graph + when one of the child graph node is updated. In this test the kernel node + parameters - blocksize and gridsize are updated. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_CmplxNstGrph_UpdKerDim") { + hipGraph_t *graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + class nestedGraph nestedGraphObj; + graph = nestedGraphObj.getRootGraph(); + nestedGraphObj.updateInnermostNode( + updateGraphNodeTests::updateGrdBlkParamTest); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr, + nullptr, 0)); + nestedGraphObj.fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + nestedGraphObj.validateOutData( + updateGraphNodeTests::updateGrdBlkParamTest); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); +} + +/** + Complex Scenario: This testcase verifies the behavior of a nested graph + when one of the nodes inside a child graph node is deleted and replaced with + a new node. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_CmplxNstGrph_DelAddNode") { + hipGraph_t *graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + class nestedGraph nestedGraphObj; + graph = nestedGraphObj.getRootGraph(); + nestedGraphObj.updateInnermostNode( + updateGraphNodeTests::deleteAddNewKerNodTest); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr, + nullptr, 0)); + nestedGraphObj.fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + nestedGraphObj.validateOutData( + updateGraphNodeTests::deleteAddNewKerNodTest); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); +} + +/** + Complex Scenario: This testcase verifies the behavior of a cloned nested + graph when one of the nodes inside a child graph node is deleted and + replaced with a new node. After modifying the original graph it is cloned + and the cloned graph is executed and validated. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_CmplxNstGrph_AddNode_Clone") { + hipGraph_t *graph, clonedGraph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + class nestedGraph nestedGraphObj; + graph = nestedGraphObj.getRootGraph(); + nestedGraphObj.updateInnermostNode( + updateGraphNodeTests::deleteAddNewKerNodTest); + HIP_CHECK(hipGraphClone(&clonedGraph, *graph)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedGraph, nullptr, + nullptr, 0)); + nestedGraphObj.fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + nestedGraphObj.validateOutData( + updateGraphNodeTests::deleteAddNewKerNodTest); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(clonedGraph)); +} + +/** + Complex Scenario: This testcase verifies the behavior of a nested graph + when one of the nodes inside a child graph node is deleted and replaced with + a new child graph node. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_CmplxNstGrph_AddChdNode") { + hipGraph_t *graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + class nestedGraph nestedGraphObj; + graph = nestedGraphObj.getRootGraph(); + nestedGraphObj.updateInnermostNode( + updateGraphNodeTests::deleteAddNewKerNodTest); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr, + nullptr, 0)); + nestedGraphObj.fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + nestedGraphObj.validateOutData( + updateGraphNodeTests::deleteAddNewKerNodTest); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); +} + +/** + Complex Scenario: This testcase verifies the behavior of a cloned nested graph + when one of the nodes inside a child graph node is deleted and replaced with + a new child graph node. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_CmplxNstGrph_AddChdNode_Clone") { + hipGraph_t *graph, clonedGraph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + class nestedGraph nestedGraphObj; + graph = nestedGraphObj.getRootGraph(); + nestedGraphObj.updateInnermostNode( + updateGraphNodeTests::deleteAddNewKerNodTest); + HIP_CHECK(hipGraphClone(&clonedGraph, *graph)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedGraph, nullptr, + nullptr, 0)); + nestedGraphObj.fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + nestedGraphObj.validateOutData( + updateGraphNodeTests::deleteAddNewKerNodTest); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(clonedGraph)); +} + +// Function to validate result +static void validateResults(int *A1_h, int *A2_h, size_t N) { + for (size_t i = 0; i < N; i++) { + int result = (A1_h[i]*A1_h[i]); + REQUIRE(result == A2_h[i]); + } +} + +/** + Functional Test to use child node as barrier to wait for multiple nodes. + This test uses child nodes to resolve dependencies between graphs. 4 + graphs are created. Graph1 contains 3 independent memcpy h2d nodes, graph2 + contains 3 independent kernel nodes and graph3 contains 3 independent + memcpy d2h nodes. Graph1, graph2 and graph3 are added as child nodes in + graph4. Graph4 is validated for functionality. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_MultGraphsAsSingleGraph") { + size_t size = 1024; + constexpr auto blocksPerCU = 6; + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, size); + hipGraph_t graph1, graph2, graph3, graph4; + std::vector nodeDependencies; + HIP_CHECK(hipGraphCreate(&graph1, 0)); + HIP_CHECK(hipGraphCreate(&graph2, 0)); + HIP_CHECK(hipGraphCreate(&graph3, 0)); + HIP_CHECK(hipGraphCreate(&graph4, 0)); + int *inputVec_d1{nullptr}, *inputVec_h1{nullptr}, *outputVec_h1{nullptr}, + *outputVec_d1{nullptr}; + int *inputVec_d2{nullptr}, *inputVec_h2{nullptr}, *outputVec_h2{nullptr}, + *outputVec_d2{nullptr}; + int *inputVec_d3{nullptr}, *inputVec_h3{nullptr}, *outputVec_h3{nullptr}, + *outputVec_d3{nullptr}; + // host and device allocation + HipTest::initArrays(&inputVec_d1, &outputVec_d1, nullptr, + &inputVec_h1, &outputVec_h1, nullptr, size, false); + HipTest::initArrays(&inputVec_d2, &outputVec_d2, nullptr, + &inputVec_h2, &outputVec_h2, nullptr, size, false); + HipTest::initArrays(&inputVec_d3, &outputVec_d3, nullptr, + &inputVec_h3, &outputVec_h3, nullptr, size, false); + // add nodes to graph + hipGraphNode_t memcpyH2D_1, memcpyH2D_2, memcpyH2D_3; + hipGraphNode_t vecSqr1, vecSqr2, vecSqr3; + hipGraphNode_t memcpyD2H_1, memcpyD2H_2, memcpyD2H_3; + hipGraphNode_t childGraphNode1, childGraphNode2, childGraphNode3; + // Create memcpy h2d nodes + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_1, graph1, nullptr, + 0, inputVec_d1, inputVec_h1, (sizeof(int)*size), hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_2, graph1, nullptr, + 0, inputVec_d2, inputVec_h2, (sizeof(int)*size), hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_3, graph1, nullptr, + 0, inputVec_d3, inputVec_h3, (sizeof(int)*size), hipMemcpyHostToDevice)); + // Create child node and add it to graph4 + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph4, nullptr, 0, + graph1)); + nodeDependencies.clear(); + nodeDependencies.push_back(childGraphNode1); + // Creating kernel nodes + hipKernelNodeParams kerNodeParams1{}, kerNodeParams2{}, kerNodeParams3{}; + void* kernelArgs1[] = {reinterpret_cast(&inputVec_d1), + reinterpret_cast(&outputVec_d1), + reinterpret_cast(&size)}; + kerNodeParams1.func = reinterpret_cast(HipTest::vector_square); + kerNodeParams1.gridDim = dim3(blocks); + kerNodeParams1.blockDim = dim3(threadsPerBlock); + kerNodeParams1.sharedMemBytes = 0; + kerNodeParams1.kernelParams = reinterpret_cast(kernelArgs1); + kerNodeParams1.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vecSqr1, graph2, nullptr, 0, + &kerNodeParams1)); + void* kernelArgs2[] = {reinterpret_cast(&inputVec_d2), + reinterpret_cast(&outputVec_d2), + reinterpret_cast(&size)}; + kerNodeParams2.func = reinterpret_cast(HipTest::vector_square); + kerNodeParams2.gridDim = dim3(blocks); + kerNodeParams2.blockDim = dim3(threadsPerBlock); + kerNodeParams2.sharedMemBytes = 0; + kerNodeParams2.kernelParams = reinterpret_cast(kernelArgs2); + kerNodeParams2.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vecSqr2, graph2, nullptr, 0, + &kerNodeParams2)); + void* kernelArgs3[] = {reinterpret_cast(&inputVec_d3), + reinterpret_cast(&outputVec_d3), + reinterpret_cast(&size)}; + kerNodeParams3.func = reinterpret_cast(HipTest::vector_square); + kerNodeParams3.gridDim = dim3(blocks); + kerNodeParams3.blockDim = dim3(threadsPerBlock); + kerNodeParams3.sharedMemBytes = 0; + kerNodeParams3.kernelParams = reinterpret_cast(kernelArgs3); + kerNodeParams3.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vecSqr3, graph2, nullptr, 0, + &kerNodeParams3)); + // Create child node and add it to graph4 + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode2, graph4, + nodeDependencies.data(), nodeDependencies.size(), graph2)); + nodeDependencies.clear(); + nodeDependencies.push_back(childGraphNode2); + // Create memcpy d2h nodes + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph3, nullptr, 0, + outputVec_h1, outputVec_d1, (sizeof(int)*size), hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph3, nullptr, 0, + outputVec_h2, outputVec_d2, (sizeof(int)*size), hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_3, graph3, nullptr, 0, + outputVec_h3, outputVec_d3, (sizeof(int)*size), hipMemcpyDeviceToHost)); + // Create child node and add it to graph4 + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode3, graph4, + nodeDependencies.data(), nodeDependencies.size(), graph3)); + nodeDependencies.clear(); + // Create executable graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec{nullptr}; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph4, nullptr, + nullptr, 0)); + // Execute graph + for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) { + // Inititalize random input data + unsigned int seed = time(nullptr); + for (size_t i = 0; i < size; i++) { + inputVec_h1[i] = (HipTest::RAND_R(&seed) & 0xFF); + inputVec_h2[i] = (HipTest::RAND_R(&seed) & 0xFF); + inputVec_h3[i] = (HipTest::RAND_R(&seed) & 0xFF); + } + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + validateResults(inputVec_h1, outputVec_h1, size); + validateResults(inputVec_h2, outputVec_h2, size); + validateResults(inputVec_h3, outputVec_h3, size); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + // Free + HipTest::freeArrays(inputVec_d1, outputVec_d1, nullptr, + inputVec_h1, outputVec_h1, nullptr, false); + HipTest::freeArrays(inputVec_d2, outputVec_d2, nullptr, + inputVec_h2, outputVec_h2, nullptr, false); + HipTest::freeArrays(inputVec_d3, outputVec_d3, nullptr, + inputVec_h3, outputVec_h3, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph4)); + HIP_CHECK(hipGraphDestroy(graph3)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipGraphDestroy(graph1)); +} + +/** + Complex Scenario: This testcase verifies the behavior of a nested graph + in multi GPU environment. Create one nested graph per GPU context. Execute + all the created graphs in their respective GPUs and validate the output. + */ +TEST_CASE("Unit_hipGraphAddChildGraphNode_CmplxNstGrph_MultGPU") { + int devcount = 0; + HIP_CHECK(hipGetDeviceCount(&devcount)); + // If only single GPU is detected then return + if (devcount < 2) { + SUCCEED("skipping the testcases as numDevices < 2"); + return; + } + hipGraph_t **graph = new hipGraph_t *[devcount](); + REQUIRE(graph != nullptr); + hipStream_t *streamForGraph = new hipStream_t[devcount]; + REQUIRE(streamForGraph != nullptr); + hipGraphExec_t *graphExec = new hipGraphExec_t[devcount]; + REQUIRE(graphExec != nullptr); + clNestedGraph** nestedGraphObj = new clNestedGraph *[devcount](); + REQUIRE(nestedGraphObj != nullptr); + // Create graph resources for each devices + for (int dev = 0; dev < devcount; dev++) { + HIP_CHECK(hipSetDevice(dev)); + nestedGraphObj[dev] = new clNestedGraph(); + REQUIRE(nestedGraphObj[dev] != nullptr); + graph[dev] = nestedGraphObj[dev]->getRootGraph(); + HIP_CHECK(hipStreamCreate(&streamForGraph[dev])); + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec[dev], *(graph[dev]), nullptr, + nullptr, 0)); + } + // Execute graph in each GPU + for (int dev = 0; dev < devcount; dev++) { + HIP_CHECK(hipSetDevice(dev)); + nestedGraphObj[dev]->fillRandInpData(); + HIP_CHECK(hipGraphLaunch(graphExec[dev], streamForGraph[dev])); + } + // Wait for each device to complete task and validate the results + for (int dev = 0; dev < devcount; dev++) { + HIP_CHECK(hipSetDevice(dev)); + HIP_CHECK(hipStreamSynchronize(streamForGraph[dev])); + nestedGraphObj[dev]->validateOutData( + updateGraphNodeTests::normalTest); + } + // Destroy graph resources + for (int dev = 0; dev < devcount; dev++) { + HIP_CHECK(hipStreamDestroy(streamForGraph[dev])); + HIP_CHECK(hipGraphExecDestroy(graphExec[dev])); + delete nestedGraphObj[dev]; + } + delete[] nestedGraphObj; + delete[] graphExec; + delete[] streamForGraph; + delete[] graph; +} From 2a01aa5b4e5648109e8f837c366b57206f5a8cf7 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 20 Jul 2023 09:55:37 +0530 Subject: [PATCH 10/17] Update performance_common.hh Replace std::reduce with std::accumulate due to RHEL build issues. [ROCm/hip-tests commit: 2c7f3749fb130759114b757dd0a8b23e9dac38da] --- projects/hip-tests/catch/include/performance_common.hh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip-tests/catch/include/performance_common.hh b/projects/hip-tests/catch/include/performance_common.hh index a40d6b3f5c..9b0c790444 100644 --- a/projects/hip-tests/catch/include/performance_common.hh +++ b/projects/hip-tests/catch/include/performance_common.hh @@ -155,7 +155,7 @@ template class Benchmark { time_ = .0; } - float sum = std::reduce(cbegin(samples), cend(samples)); + float sum = std::accumulate(cbegin(samples), cend(samples), .0); float mean = sum / samples.size(); float deviation = From 6fb65da9a13498791cb266ce7981b6d445d5046c Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 20 Jul 2023 10:15:03 +0530 Subject: [PATCH 11/17] SWDEV-403770 - Enable tests. (#330) Change-Id: I6e31c7ecea0f8343c37fd8a6249bf1cf1e455be6 [ROCm/hip-tests commit: faa2dd7cfbce2135ab32c825b9cd4ef92cc8f400] --- .../catch/hipTestMain/config/config_amd_linux_common.json | 3 --- 1 file changed, 3 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index 1d2c9e458b..2fec6cffa1 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -39,9 +39,6 @@ "Unit_hipDrvMemcpy3DAsync_Positive_Array", "Unit_hipMemRangeGetAttribute_Positive_AccessedBy_Basic", "Unit_hipMemRangeGetAttribute_Positive_AccessedBy_Partial_Range", - "Unit_hipStreamAttachMemAsync_Positive_Basic", - "Unit_hipStreamAttachMemAsync_Positive_AttachGlobal", - "Unit_hipStreamAttachMemAsync_Negative_Parameters", "Unit_hipMemGetAddressRange_Positive", "Unit_hipGraphAddMemcpyNode1D_Negative_Basic", "intermittent issue: corrupted double-linked list", From f72821a84d443f78a3b93cc9829f83862da8f78b Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 20 Jul 2023 10:16:56 +0530 Subject: [PATCH 12/17] SWDEV-408958 - Use LaunchDelayKernel and modify it to use same kernel based on real time clock for gfx10 and gfx11. (#370) Change-Id: Iea8a48e8cbfa1745c7d5535dc5820133a4104087 [ROCm/hip-tests commit: 04080c2e2eb88d0eb4e0c484d7b54961f658fca1] --- .../catch/include/hip_test_common.hh | 81 ------------------- projects/hip-tests/catch/include/utils.hh | 26 +++--- .../unit/event/Unit_hipEventElapsedTime.cc | 8 +- .../catch/unit/event/Unit_hipEventQuery.cc | 31 +++++-- .../catch/unit/event/hipEventDestroy.cc | 4 +- .../hip-tests/catch/unit/memory/hipFree.cc | 12 +-- .../catch/unit/memory/hipMemcpySync.cc | 4 +- .../catch/unit/memory/hipMemsetSync.cc | 3 +- .../unit/stream/hipStreamACb_MultiThread.cc | 32 +------- .../unit/stream/hipStreamCreateWithFlags.cc | 5 +- .../catch/unit/stream/hipStreamDestroy.cc | 3 +- .../catch/unit/stream/hipStreamQuery.cc | 9 +-- .../catch/unit/stream/hipStreamSynchronize.cc | 18 ++--- .../catch/unit/stream/hipStreamWaitEvent.cc | 52 +----------- 14 files changed, 79 insertions(+), 209 deletions(-) diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index d4c44a813f..7b4b08ee38 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -350,87 +350,6 @@ template <> struct MemTraits { } }; - -namespace { -static __global__ void waitKernel(size_t offset) { - auto start = clock(); - while ((clock() - start) < offset) { - } -} - -static __global__ void waitKernel_gfx11(size_t offset) { -#if HT_AMD - auto start = wall_clock64(); - while ((wall_clock64() - start) < offset) { - } -#endif -} - -// helper function used to set the device frequency variable -// estimates the number of clock ticks in 1 second -static size_t findTicksPerSecond() { - // first read the reported clockRate as a starting point - hipDeviceProp_t prop; - int device; - HIP_CHECK(hipGetDevice(&device)); - HIP_CHECK(hipGetDeviceProperties(&prop, device)); - size_t devFreq = static_cast(prop.clockRate); // in kHz - size_t clockTicksPerSecond = devFreq * 1000; - - // init - hipEvent_t start, stop; - HIP_CHECK(hipEventCreate(&start)); - HIP_CHECK(hipEventCreate(&stop)); - auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel; - // Warmup - hipLaunchKernelGGL(waitKernel_used, dim3(1), dim3(1), 0, 0, clockTicksPerSecond); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipDeviceSynchronize()); - - // try 10 times to find device frequency - // after 10 attempts the result is likely good enough so just accept it - for (int attempts = 10; attempts > 0; --attempts) { - HIP_CHECK(hipEventRecord(start)); - hipLaunchKernelGGL(waitKernel_used, dim3(1), dim3(1), 0, 0, clockTicksPerSecond); - HIP_CHECK(hipEventRecord(stop)); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipEventSynchronize(stop)); - - float executionTimeMs = 0; - HIP_CHECK(hipEventElapsedTime(&executionTimeMs, start, stop)); - - constexpr float tolerance = 20; - if (fabs(executionTimeMs - 1000) <= tolerance) { - // Timing is within accepted tolerance, break here - break; - } else { - clockTicksPerSecond = (clockTicksPerSecond * 1000) / executionTimeMs; - --attempts; - } - } - - // deinit - HIP_CHECK(hipEventDestroy(start)); - HIP_CHECK(hipEventDestroy(stop)); - return clockTicksPerSecond; -} -} // namespace - -// Launches a kernel which runs for specified amount of time -// Note: The current implementation uses HIP_CHECK which is not thread safe! -// Note: the function assumes execution on a single device and caches the number of clock ticks per -// second -static inline void runKernelForDuration(std::chrono::milliseconds duration, - hipStream_t stream = nullptr) { - // number of clocks the device is running at (device frequency) - // each translation unit will have a copy of ticksPerSecond but this function isn't designed for - // precision so that's acceptable. - static size_t ticksPerSecond = findTicksPerSecond(); - const auto millis = duration.count(); - auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel; - hipLaunchKernelGGL(waitKernel_used, dim3(1), dim3(1), 0, stream, ticksPerSecond * millis / 1000); -} - class BlockingContext { std::atomic_bool blocked{true}; hipStream_t stream; diff --git a/projects/hip-tests/catch/include/utils.hh b/projects/hip-tests/catch/include/utils.hh index 6dd8070b9c..5efd6a5125 100644 --- a/projects/hip-tests/catch/include/utils.hh +++ b/projects/hip-tests/catch/include/utils.hh @@ -122,9 +122,17 @@ template __global__ void VectorSet(T* const vec, const T value, siz // Will execute for atleast interval milliseconds static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) { while (interval--) { - uint64_t start = clock(); - while (clock() - start < ticks_per_ms) { + #if HT_AMD + uint64_t start = wall_clock64(); + while (wall_clock64() - start < ticks_per_ms) { + __builtin_amdgcn_s_sleep(10); } + #endif + #if HT_NVIDIA + uint64_t start = clock64(); + while (clock64() - start < ticks_per_ms) { + } + #endif } } @@ -140,14 +148,14 @@ __global__ void Iota(T* const out, size_t pitch, size_t w, size_t h, size_t d) { } } -inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) { +inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream = nullptr) { int ticks_per_ms = 0; - // Clock rate is in kHz => number of clock ticks in a millisecond - if (IsGfx11()) { - HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeWallClockRate, 0)); - } else { - HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0)); - } + #if HT_AMD + HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeWallClockRate, 0)); + #endif + #if HT_NVIDIA + HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0)); + #endif Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms); } diff --git a/projects/hip-tests/catch/unit/event/Unit_hipEventElapsedTime.cc b/projects/hip-tests/catch/unit/event/Unit_hipEventElapsedTime.cc index 93deb7f842..cf8d6fb919 100644 --- a/projects/hip-tests/catch/unit/event/Unit_hipEventElapsedTime.cc +++ b/projects/hip-tests/catch/unit/event/Unit_hipEventElapsedTime.cc @@ -23,7 +23,7 @@ THE SOFTWARE. #include #include #include - +#include /** * @addtogroup hipEventElapsedTime hipEventElapsedTime * @{ @@ -158,10 +158,7 @@ TEST_CASE("Unit_hipEventElapsedTime_NotReady_Negative") { // Record start event HIP_CHECK(hipEventRecord(start, nullptr)); - HipTest::BlockingContext b_context{nullptr}; - b_context.block_stream(); // blocked stream - REQUIRE(b_context.is_blocked()); - + LaunchDelayKernel(std::chrono::milliseconds(1000)); // Record stop event HIP_CHECK(hipEventRecord(stop, nullptr)); @@ -169,7 +166,6 @@ TEST_CASE("Unit_hipEventElapsedTime_NotReady_Negative") { float tElapsed = 1.0f; HIP_CHECK_ERROR(hipEventQuery(stop), hipErrorNotReady); HIP_ASSERT(hipEventElapsedTime(&tElapsed, start, stop) == hipErrorNotReady); - b_context.unblock_stream(); HIP_CHECK(hipStreamSynchronize(nullptr)); HIP_CHECK(hipEventDestroy(start)); diff --git a/projects/hip-tests/catch/unit/event/Unit_hipEventQuery.cc b/projects/hip-tests/catch/unit/event/Unit_hipEventQuery.cc index b9606fccf4..29f7b4d50c 100644 --- a/projects/hip-tests/catch/unit/event/Unit_hipEventQuery.cc +++ b/projects/hip-tests/catch/unit/event/Unit_hipEventQuery.cc @@ -18,7 +18,29 @@ THE SOFTWARE. */ #include +#include +/** + * @addtogroup hipEventQuery hipEventQuery + * @{ + * @ingroup EventTest + * `hipEventQuery(hipEvent_t event)` - + * Query the status of the specified event. + * ________________________ + * Test cases from other modules: + * - @ref Unit_hipEventIpc + */ +/** + * Test Description + * ------------------------ + * - Query events with a single and with multiple devices. + * Test source + * ------------------------ + * - unit/event/Unit_hipEventQuery.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipEventQuery_DifferentDevice") { hipEvent_t event1{}, event2{}; HIP_CHECK(hipEventCreate(&event1)); @@ -30,15 +52,10 @@ TEST_CASE("Unit_hipEventQuery_DifferentDevice") { HIP_CHECK(hipStreamCreate(&stream)); REQUIRE(stream != nullptr); - HipTest::BlockingContext b_context1{stream}; // og context - // Block stream { HIP_CHECK(hipSetDevice(0)); HIP_CHECK(hipEventRecord(event1, stream)); - - b_context1.block_stream(); // blocked stream - REQUIRE(b_context1.is_blocked()); - + LaunchDelayKernel(std::chrono::milliseconds(3000), stream); HIP_CHECK(hipEventRecord(event2, stream)); HIP_CHECK(hipEventSynchronize(event1)); @@ -58,8 +75,6 @@ TEST_CASE("Unit_hipEventQuery_DifferentDevice") { HIP_CHECK(hipEventQuery(event1)); HIP_CHECK_ERROR(hipEventQuery(event2), hipErrorNotReady); - b_context1.unblock_stream(); - HIP_CHECK(hipEventSynchronize(event2)); // Query, should be done now diff --git a/projects/hip-tests/catch/unit/event/hipEventDestroy.cc b/projects/hip-tests/catch/unit/event/hipEventDestroy.cc index d62921f01b..8836e3db7a 100644 --- a/projects/hip-tests/catch/unit/event/hipEventDestroy.cc +++ b/projects/hip-tests/catch/unit/event/hipEventDestroy.cc @@ -25,7 +25,7 @@ THE SOFTWARE. #include #include #include "hip/hip_runtime_api.h" - +#include /** * @addtogroup hipEventDestroy hipEventDestroy * @{ @@ -53,7 +53,7 @@ static inline void launchVectorAdd(float*& A_h, float*& B_h, float*& C_h, HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), A_h, 0)); HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&B_d), B_h, 0)); HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&C_d), C_h, 0)); - HipTest::runKernelForDuration(delay, stream); + LaunchDelayKernel(delay, stream); HipTest::vectorADD<<<1, 1, 0, stream>>>(A_d, B_d, C_d, vectorSize); } diff --git a/projects/hip-tests/catch/unit/memory/hipFree.cc b/projects/hip-tests/catch/unit/memory/hipFree.cc index 1db0ec9502..bc06a4c7a6 100644 --- a/projects/hip-tests/catch/unit/memory/hipFree.cc +++ b/projects/hip-tests/catch/unit/memory/hipFree.cc @@ -25,7 +25,7 @@ THE SOFTWARE. #include #include "hipArrayCommon.hh" #include "DriverContext.hh" - +#include /* * This testcase verifies [ hipFree || hipFreeArray || hipFreeType::ArrayDestroy || * hipFreeType::HostFree with hipHostMalloc ] @@ -54,7 +54,7 @@ TEST_CASE("Unit_hipFreeImplicitSyncDev") { size_t size_mult = GENERATE(1, 32, 64, 128, 256); HIP_CHECK(hipMalloc(&devPtr, sizeof(*devPtr) * size_mult)); - HipTest::runKernelForDuration(delay); + LaunchDelayKernel(delay); // make sure device is busy HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); HIP_CHECK(hipFree(devPtr)); @@ -67,7 +67,7 @@ TEST_CASE("Unit_hipFreeImplicitSyncHost") { HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(*hostPtr) * size_mult)); - HipTest::runKernelForDuration(delay); + LaunchDelayKernel(delay); // make sure device is busy HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); HIP_CHECK(hipHostFree(hostPtr)); @@ -88,7 +88,7 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo hipChannelFormatDesc desc = hipCreateChannelDesc(); HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayDefault)); - HipTest::runKernelForDuration(delay); + LaunchDelayKernel(delay); // make sure device is busy HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); HIP_CHECK(hipFreeArray(arrayPtr)); @@ -103,7 +103,7 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo cuDesc.Format = vec_info::format; cuDesc.NumChannels = vec_info::size; HIP_CHECK(hipArrayCreate(&cuArrayPtr, &cuDesc)); - HipTest::runKernelForDuration(delay); + LaunchDelayKernel(delay); // make sure device is busy HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); HIP_CHECK(hipArrayDestroy(cuArrayPtr)); @@ -120,7 +120,7 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo hipChannelFormatDesc desc = hipCreateChannelDesc(); HIP_CHECK(hipMallocArray(&arrayPtr, &desc, extent.width, extent.height, hipArrayDefault)); - HipTest::runKernelForDuration(delay); + LaunchDelayKernel(delay); // make sure device is busy HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); // Second free segfaults diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpySync.cc b/projects/hip-tests/catch/unit/memory/hipMemcpySync.cc index fbb165d892..bef3c953f5 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpySync.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpySync.cc @@ -21,7 +21,7 @@ THE SOFTWARE. */ #include "MemUtils.hh" - +#include /* * These testcases verify that synchronization behaviour for memcpy functions with respect to * the host. @@ -156,7 +156,7 @@ static void runMemcpyTests(hipStream_t stream, bool async, allocType type, memTy using namespace std::chrono_literals; const std::chrono::duration delay = 100ms; - HipTest::runKernelForDuration(delay, stream); + LaunchDelayKernel(delay, stream); memcpyCheck(type, memType, aPtr.first, data, fillerData, async, stream, fromHost); checkForSync(stream, async, type, fromHost); diff --git a/projects/hip-tests/catch/unit/memory/hipMemsetSync.cc b/projects/hip-tests/catch/unit/memory/hipMemsetSync.cc index f3de97e02f..973c7dc6e7 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemsetSync.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemsetSync.cc @@ -21,6 +21,7 @@ THE SOFTWARE. */ #include +#include /* * These testcases verify that synchronous memset functions are asynchronous with respect to the * host except when the target is pinned host memory or a Unified Memory region @@ -406,7 +407,7 @@ void runTests(allocType type, memSetType memsetType, MultiDData data, hipStream_ std::pair aPtr = initMemory(type, memsetType, data); using namespace std::chrono_literals; const std::chrono::duration delay = 100ms; - HipTest::runKernelForDuration(delay, stream); + LaunchDelayKernel(delay, stream); memsetCheck(aPtr.first, testValue, memsetType, data, async, stream); if (async || type == allocType::deviceMalloc) { diff --git a/projects/hip-tests/catch/unit/stream/hipStreamACb_MultiThread.cc b/projects/hip-tests/catch/unit/stream/hipStreamACb_MultiThread.cc index 2825914cba..9d1a780deb 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamACb_MultiThread.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamACb_MultiThread.cc @@ -25,6 +25,7 @@ multiple Threads. #include #include +#include static constexpr size_t N = 4096; static constexpr int numThreads = 1000; @@ -43,33 +44,6 @@ static __global__ void device_function(float* C_d, float* A_d, size_t Num) { for (size_t i = gputhread; i < Num; i += stride) { C_d[i] = A_d[i] * A_d[i]; } - - // Delay thread 1 only in the GPU - if (gputhread == 1) { - uint64_t wait_t = 3200000000, start = clock64(), cur; - do { - cur = clock64() - start; - } while (cur < wait_t); - } -} - -static __global__ void device_function_gfx11(float* C_d, float* A_d, size_t Num) { -#if HT_AMD - size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x); - size_t stride = blockDim.x * gridDim.x; - - for (size_t i = gputhread; i < Num; i += stride) { - C_d[i] = A_d[i] * A_d[i]; - } - - // Delay thread 1 only in the GPU - if (gputhread == 1) { - uint64_t wait_t = 3200000000, start = wall_clock64(), cur; - do { - cur = wall_clock64() - start; - } while (cur < wait_t); - } -#endif } static void HIPRT_CB Thread1_Callback(hipStream_t stream, hipError_t status, @@ -146,10 +120,10 @@ TEST_CASE("Unit_hipStreamAddCallback_MultipleThreads") { constexpr unsigned threadsPerBlock = 256; constexpr unsigned blocks = (N + 255)/threadsPerBlock; - auto device_function_used = IsGfx11() ? device_function_gfx11 : device_function; - hipLaunchKernelGGL((device_function_used), dim3(blocks), + hipLaunchKernelGGL((device_function), dim3(blocks), dim3(threadsPerBlock), 0, mystream, C_d, A_d, N); + LaunchDelayKernel(std::chrono::milliseconds(2000), mystream); HIP_CHECK(hipGetLastError()); HIP_CHECK( hipMemcpyAsync(C1_h, C_d, Nbytes, diff --git a/projects/hip-tests/catch/unit/stream/hipStreamCreateWithFlags.cc b/projects/hip-tests/catch/unit/stream/hipStreamCreateWithFlags.cc index ef38755028..5038037452 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamCreateWithFlags.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamCreateWithFlags.cc @@ -19,6 +19,7 @@ THE SOFTWARE. #include #include +#include namespace hipStreamCreateWithFlagsTests { @@ -69,11 +70,11 @@ TEST_CASE("Unit_hipStreamCreateWithFlags_DefaultStreamInteraction") { constexpr auto delay = std::chrono::milliseconds(500); SECTION("default stream waiting for created stream") { - HipTest::runKernelForDuration(delay, stream); + LaunchDelayKernel(delay, stream); REQUIRE(hipStreamQuery(defaultStream) == expectedError); } SECTION("created stream waiting for default stream") { - HipTest::runKernelForDuration(delay, defaultStream); + LaunchDelayKernel(delay, defaultStream); REQUIRE(hipStreamQuery(stream) == expectedError); } diff --git a/projects/hip-tests/catch/unit/stream/hipStreamDestroy.cc b/projects/hip-tests/catch/unit/stream/hipStreamDestroy.cc index 7c7658e64b..a2c0f287ec 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamDestroy.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamDestroy.cc @@ -18,6 +18,7 @@ THE SOFTWARE. */ #include #include +#include namespace hipStreamDestroyTests { @@ -80,7 +81,7 @@ TEST_CASE("Unit_hipStreamDestroy_WithPendingWork") { HIP_CHECK(hipMalloc(&deviceData, sizeof(int) * numDataPoints)); HIP_CHECK(hipMemset(deviceData, 0, sizeof(int) * numDataPoints)); - HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream); + LaunchDelayKernel(std::chrono::milliseconds(500), stream); setToOne<<<1, numDataPoints, 0, stream>>>(deviceData, numDataPoints); HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady); HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); diff --git a/projects/hip-tests/catch/unit/stream/hipStreamQuery.cc b/projects/hip-tests/catch/unit/stream/hipStreamQuery.cc index 8b99dc712c..6a1f306c6c 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamQuery.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamQuery.cc @@ -19,7 +19,7 @@ THE SOFTWARE. #include #include "streamCommon.hh" - +#include /** * @brief Check that querying a stream with no work returns hipSuccess * @@ -101,7 +101,7 @@ TEST_CASE("Unit_hipStreamQuery_SubmitWorkOnStreamAndQueryNullStream") { HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipStreamQuery(hip::nullStream)); - HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream); + LaunchDelayKernel(std::chrono::milliseconds(500), stream); HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); HIP_CHECK(hipDeviceSynchronize()); @@ -116,7 +116,7 @@ TEST_CASE("Unit_hipStreamQuery_SubmitWorkOnStreamAndQueryNullStream") { */ TEST_CASE("Unit_hipStreamQuery_NullStreamQuery") { HIP_CHECK(hipStreamQuery(hip::nullStream)); - HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::nullStream); + LaunchDelayKernel(std::chrono::milliseconds(500), hip::nullStream); HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); HIP_CHECK(hipStreamSynchronize(hip::nullStream)); @@ -130,8 +130,7 @@ TEST_CASE("Unit_hipStreamQuery_WithPendingWork") { hipStream_t waitingStream{nullptr}; HIP_CHECK(hipStreamCreate(&waitingStream)); - HipTest::runKernelForDuration(std::chrono::milliseconds(500), waitingStream); - + LaunchDelayKernel(std::chrono::milliseconds(500), waitingStream); HIP_CHECK_ERROR(hipStreamQuery(waitingStream), hipErrorNotReady); HIP_CHECK(hipStreamSynchronize(waitingStream)); HIP_CHECK(hipStreamQuery(waitingStream)); diff --git a/projects/hip-tests/catch/unit/stream/hipStreamSynchronize.cc b/projects/hip-tests/catch/unit/stream/hipStreamSynchronize.cc index a31546a52a..d1faaa4f04 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamSynchronize.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamSynchronize.cc @@ -19,7 +19,7 @@ THE SOFTWARE. #include #include "streamCommon.hh" - +#include namespace hipStreamSynchronizeTest { /** @@ -62,7 +62,7 @@ TEST_CASE("Unit_hipStreamSynchronize_FinishWork") { HIP_CHECK(hipStreamCreate(&stream)); } - HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream); + LaunchDelayKernel(std::chrono::milliseconds(500), stream); HIP_CHECK(hipStreamSynchronize(stream)); HIP_CHECK(hipStreamQuery(stream)); @@ -86,15 +86,15 @@ TEST_CASE("Unit_hipStreamSynchronize_NullStreamSynchronization") { } for (int i = 0; i < totalStreams; ++i) { - HipTest::runKernelForDuration(std::chrono::milliseconds(1000), streams[i]); + LaunchDelayKernel(std::chrono::milliseconds(1000), streams[i]); } + HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); + for (int i = 0; i < totalStreams; ++i) { HIP_CHECK_ERROR(hipStreamQuery(streams[i]), hipErrorNotReady); } - HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); - HIP_CHECK(hipStreamSynchronize(hip::nullStream)); HIP_CHECK(hipStreamQuery(hip::nullStream)); @@ -123,8 +123,8 @@ TEST_CASE("Unit_hipStreamSynchronize_SynchronizeStreamAndQueryNullStream") { HIP_CHECK(hipStreamCreate(&stream1)); HIP_CHECK(hipStreamCreate(&stream2)); - HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream1); - HipTest::runKernelForDuration(std::chrono::milliseconds(2000), stream2); + LaunchDelayKernel(std::chrono::milliseconds(500), stream1); + LaunchDelayKernel(std::chrono::milliseconds(2000), stream2); SECTION("Do not use NullStream") {} SECTION("Submit Kernel to NullStream") { @@ -157,10 +157,10 @@ TEST_CASE("Unit_hipStreamSynchronize_SynchronizeStreamAndQueryNullStream") { * */ TEST_CASE("Unit_hipStreamSynchronize_NullStreamAndStreamPerThread") { - HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::streamPerThread); + LaunchDelayKernel(std::chrono::milliseconds(500), hip::streamPerThread); HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); HIP_CHECK_ERROR(hipStreamQuery(hip::streamPerThread), hipErrorNotReady); - HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::nullStream); + LaunchDelayKernel(std::chrono::milliseconds(500), hip::nullStream); HIP_CHECK(hipStreamSynchronize(hip::nullStream)) HIP_CHECK_ERROR(hipStreamQuery(hip::streamPerThread), hipSuccess); HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipSuccess); diff --git a/projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc b/projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc index 7965bc19ef..9877d548c1 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc @@ -25,7 +25,7 @@ Unit_hipStreamWaitEvent_DifferentStreams - Test waiting for an event on a differ */ #include - +#include TEST_CASE("Unit_hipStreamWaitEvent_Negative") { enum class StreamTestType { NullStream = 0, StreamPerThread, CreatedStream }; @@ -79,35 +79,6 @@ TEST_CASE("Unit_hipStreamWaitEvent_UninitializedStream_Negative") { } #endif -// Since we can not use atomic*_system on every gpu, we will use wait based on clock rate. -// This wont be accurate since clock rate of a GPU varies depending on many variables including -// thermals, load, utilization -__global__ void waitKernel(int clockRate, int seconds) { - auto start = clock(); - auto ms = seconds * 1000; - long long waitTill = clockRate * (long long)ms; - while (1) { - auto end = clock(); - if ((end - start) > waitTill) { - return; - } - } -} - -__global__ void waitKernel_gfx11(int clockRate, int seconds) { -#if HT_AMD - auto start = wall_clock64(); - auto ms = seconds * 1000; - long long waitTill = clockRate * (long long)ms; - while (1) { - auto end = wall_clock64(); - if ((end - start) > waitTill) { - return; - } - } -#endif -} - TEST_CASE("Unit_hipStreamWaitEvent_Default") { hipStream_t stream{nullptr}; hipEvent_t waitEvent{nullptr}; @@ -118,15 +89,7 @@ TEST_CASE("Unit_hipStreamWaitEvent_Default") { REQUIRE(stream != nullptr); REQUIRE(waitEvent != nullptr); - int deviceId {}; - HIP_CHECK(hipGetDevice(&deviceId)); - - hipDeviceProp_t prop{}; - HIP_CHECK(hipGetDeviceProperties(&prop, deviceId)); - auto clockRate = prop.clockRate; - - auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel; - waitKernel_used<<<1, 1, 0, stream>>>(clockRate, 2); // Wait for 2 seconds + LaunchDelayKernel(std::chrono::milliseconds(2000), stream); HIP_CHECK(hipEventRecord(waitEvent, stream)); @@ -154,15 +117,8 @@ TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") { REQUIRE(streamBlockedOnStreamA != nullptr); REQUIRE(waitEvent != nullptr); - int deviceId {}; - HIP_CHECK(hipGetDevice(&deviceId)); + LaunchDelayKernel(std::chrono::milliseconds(3000), blockedStreamA); - hipDeviceProp_t prop{}; - HIP_CHECK(hipGetDeviceProperties(&prop, deviceId)); - auto clockRate = prop.clockRate; - auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel; - waitKernel_used<<<1, 1, 0, blockedStreamA>>>(clockRate, - 3); // wait for 3 seconds HIP_CHECK(hipEventRecord(waitEvent, blockedStreamA)); // Make sure stream is waiting for data to be set @@ -170,7 +126,7 @@ TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") { HIP_CHECK(hipStreamWaitEvent(streamBlockedOnStreamA, waitEvent, 0)); - waitKernel_used<<<1, 1, 0, streamBlockedOnStreamA>>>(clockRate, 2); // Wait for 2 seconds + LaunchDelayKernel(std::chrono::milliseconds(2000), streamBlockedOnStreamA); HIP_CHECK(hipStreamSynchronize(unblockingStream)); From fa900628f1f41484c7a2ab552f522563c4213bcd Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 20 Jul 2023 10:17:26 +0530 Subject: [PATCH 13/17] SWDEV-404873 - Add double quotes that handles space in path for windows OS. (#371) Change-Id: I4752072914aded5e39105913f5bd79fcaf4b499b [ROCm/hip-tests commit: 573aefb65255e734a406daf030ece19913adc09a] --- projects/hip-tests/catch/include/hip_test_process.hh | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/projects/hip-tests/catch/include/hip_test_process.hh b/projects/hip-tests/catch/include/hip_test_process.hh index e65be29d7e..ee5facd704 100644 --- a/projects/hip-tests/catch/include/hip_test_process.hh +++ b/projects/hip-tests/catch/include/hip_test_process.hh @@ -85,6 +85,10 @@ class SpawnProc { INFO("Testing that capture file does not exist already: " << tmpFileName); REQUIRE(!fs::exists(tmpFileName)); } + if (TestContext::get().isWindows()) { + exeName = (exeName.find(" ", 0) == std::string::npos) ? exeName : ("\"" + exeName + "\""); + tmpFileName = (tmpFileName.find(" ", 0) == std::string::npos) ? tmpFileName : ("\"" + tmpFileName + "\""); + } } int run(std::string commandLineArgs = "") { @@ -100,6 +104,9 @@ class SpawnProc { execCmd += " > "; execCmd += tmpFileName; } + if (TestContext::get().isWindows()) { + execCmd = (execCmd.find(" ", 0) == std::string::npos) ? execCmd : ("\"" + execCmd + "\""); + } auto res = std::system(execCmd.c_str()); if (captureOutput) { From 969c118fa5a968f11e65cd57595a5ae62ebcdd24 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 20 Jul 2023 10:18:30 +0530 Subject: [PATCH 14/17] SWDEV-409588 - hipMemsetDSync tests should not use the size as 1 while using D32/D16 APIs (#372) Change-Id: I32c9bed860ddf4fe0d7bba21dce9bd720168c396 [ROCm/hip-tests commit: 74c143392b78282f8891e82bfe7037fcb6727417] --- .../config/config_amd_linux_common.json | 1 - .../catch/unit/memory/hipMemsetSync.cc | 24 ++++++++++--------- 2 files changed, 13 insertions(+), 12 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index 2fec6cffa1..48cbccbb52 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -90,7 +90,6 @@ "Unit_hipMemGetInfo_ParaNonDiv", "Unit_hipMemGetInfo_ParaMultiSmall", "Unit_hipMemGetInfo_Negative", - "Unit_hipMemsetDSync - int8_t", "Unit_hipGraphClone_Test_hipGraphExecMemcpyNodeSetParams", "Unit_hipGraphClone_Test_hipGraphMemcpyNodeSetParams1D_and_exec", "Unit_hipStreamValue_Wait64_Blocking_NoMask_And", diff --git a/projects/hip-tests/catch/unit/memory/hipMemsetSync.cc b/projects/hip-tests/catch/unit/memory/hipMemsetSync.cc index 973c7dc6e7..950833f2b7 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemsetSync.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemsetSync.cc @@ -52,8 +52,8 @@ struct MultiDData { // set of helper functions to tidy the nested switch statements template -static std::pair deviceMallocHelper(memSetType memType, size_t dataW, size_t dataH, size_t dataD, - size_t& dataPitch) { +static std::pair deviceMallocHelper(memSetType memType, size_t dataW, size_t dataH, + size_t dataD, size_t& dataPitch) { size_t elementSize = sizeof(T); size_t sizeInBytes = elementSize * dataW * dataH * dataD; T* aPtr{}; @@ -89,7 +89,8 @@ static std::pair deviceMallocHelper(memSetType memType, size_t dataW, siz } template -static std::pair hostMallocHelper(size_t dataW, size_t dataH, size_t dataD, size_t& dataPitch) { +static std::pair hostMallocHelper(size_t dataW, size_t dataH, size_t dataD, + size_t& dataPitch) { size_t elementSize = sizeof(T); size_t sizeInBytes = elementSize * dataW * dataH * dataD; T* aPtr; @@ -101,7 +102,8 @@ static std::pair hostMallocHelper(size_t dataW, size_t dataH, size_t dat } template -static std::pair hostRegisteredHelper(size_t dataW, size_t dataH, size_t dataD, size_t& dataPitch) { +static std::pair hostRegisteredHelper(size_t dataW, size_t dataH, size_t dataD, + size_t& dataPitch) { size_t elementSize = sizeof(T); size_t sizeInBytes = elementSize * dataW * dataH * dataD; T* aPtr = new T[dataW * dataH * dataD]; @@ -448,7 +450,7 @@ TEST_CASE("Unit_hipMemsetSync") { allocType::devRegistered); memSetType memset_type = memSetType::hipMemset; MultiDData data; - data.width = GENERATE(1, 1024); + data.width = GENERATE(512, 1024); doMemsetTest(type, memset_type, data); } @@ -461,7 +463,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemsetDSync", "", int8_t, int16_t, uint32_t) { allocType::hostMalloc, allocType::devRegistered); memSetType memset_type; MultiDData data; - data.width = GENERATE(1, 1024); + data.width = GENERATE(512, 1024); if (std::is_same::value) { memset_type = memSetType::hipMemsetD8; @@ -483,8 +485,8 @@ TEST_CASE("Unit_hipMemset2DSync") { allocType::hostRegisted, allocType::devRegistered); memSetType memset_type = memSetType::hipMemset2D; MultiDData data; - data.width = GENERATE(1, 1024); - data.height = GENERATE(1, 1024); + data.width = GENERATE(512, 1024); + data.height = GENERATE(512, 1024); doMemsetTest(mallocType, memset_type, data); } @@ -498,9 +500,9 @@ TEST_CASE("Unit_hipMemset3DSync") { allocType::hostRegisted, allocType::devRegistered); memSetType memset_type = memSetType::hipMemset3D; MultiDData data; - data.width = GENERATE(1, 256); - data.height = GENERATE(1, 256); - data.depth = GENERATE(1, 256); + data.width = GENERATE(128, 256); + data.height = GENERATE(128, 256); + data.depth = GENERATE(128, 256); doMemsetTest(mallocType, memset_type, data); } From 0b6f88a5d3a4f86074df47a714b0ab5a44842436 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 20 Jul 2023 10:20:05 +0530 Subject: [PATCH 15/17] SWDEV-409113,SWDEV-409114,SWDEV-409112,SWDEV-409111,SWDEV-409110,SWDEV-409109,SWDEV-409096,SWDEV-409105 - Remove strong assumptions on hipMemGetInfo reported memory. (#373) Change-Id: I6751e20690b93a3db35d98e45d295e5465387c5a [ROCm/hip-tests commit: b8fb6f88b9ee464fbc3adc86cc5a5e2445477ad5] --- .../config/config_amd_linux_common.json | 9 - .../catch/unit/memory/hipArrayCommon.hh | 6 - .../catch/unit/memory/hipArrayCreate.cc | 17 +- .../catch/unit/memory/hipMalloc3D.cc | 16 +- .../catch/unit/memory/hipMemGetInfo.cc | 556 ------------------ 5 files changed, 3 insertions(+), 601 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index 48cbccbb52..b8bf15e523 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -81,15 +81,6 @@ "=== Below tests fail in stress test on 23/06/23 ===", "Unit_hipIpcMemAccess_ParameterValidation", "Unit_hipMemcpy2DFromArrayAsync_Positive_Synchronization_Behavior", - "Unit_hipMalloc3D_SmallandBigChunks", - "Unit_hipMalloc3D_MultiThread", - "Unit_hipArrayCreate_DiffSizes", - "Unit_hipArrayCreate_MultiThread", - "hipMemGetInfo_DifferentMallocSmall", - "Unit_hipMemGetInfo_ParaSmall", - "Unit_hipMemGetInfo_ParaNonDiv", - "Unit_hipMemGetInfo_ParaMultiSmall", - "Unit_hipMemGetInfo_Negative", "Unit_hipGraphClone_Test_hipGraphExecMemcpyNodeSetParams", "Unit_hipGraphClone_Test_hipGraphMemcpyNodeSetParams1D_and_exec", "Unit_hipStreamValue_Wait64_Blocking_NoMask_And", diff --git a/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh b/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh index b0beeb3126..b19a32a206 100644 --- a/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh +++ b/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh @@ -67,12 +67,6 @@ template void checkDataIsAscending(const std::vector& hostData) REQUIRE(allMatch); } -inline size_t getFreeMem() { - size_t free = 0, total = 0; - HIP_CHECK(hipMemGetInfo(&free, &total)); - return free; -} - struct Sizes { int max1D; std::array max2D; diff --git a/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc b/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc index 70a8636922..a22386070b 100644 --- a/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc +++ b/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc @@ -42,12 +42,9 @@ static constexpr auto ARRAY_LOOP{100}; * bigger chunks of data. * Two scenarios are verified in this API * 1. SmallArray: Allocates NUM_W*NUM_H in a loop and - * releases the memory and verifies the meminfo. + * releases the memory. * 2. BigArray: Allocates BIGNUM_W*BIGNUM_H in a loop and - * releases the memory and verifies the meminfo - * - * In both cases, the memory info before allocation and - * after releasing the memory should be the same. + * releases the memory. * */ @@ -71,9 +68,6 @@ static void ArrayCreate_DiffSizes(int gpu) { for (int i = 0; i < ARRAY_LOOP; i++) { HIP_CHECK_THREAD(hipArrayDestroy(array[i])); } - - HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); - REQUIRE_THREAD(pavail == avail); } } @@ -94,7 +88,6 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { devCnt = HipTest::getDeviceCount(); - const size_t pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { threadlist.push_back(std::thread(ArrayCreate_DiffSizes, i)); } @@ -103,12 +96,6 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { t.join(); } HIP_CHECK_THREAD_FINALIZE(); - const size_t avail = getFreeMem(); - - if (pavail != avail) { - WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); - REQUIRE(false); - } } diff --git a/projects/hip-tests/catch/unit/memory/hipMalloc3D.cc b/projects/hip-tests/catch/unit/memory/hipMalloc3D.cc index 386eff0a51..7e0e249bee 100644 --- a/projects/hip-tests/catch/unit/memory/hipMalloc3D.cc +++ b/projects/hip-tests/catch/unit/memory/hipMalloc3D.cc @@ -31,9 +31,7 @@ static constexpr auto CHUNK_LOOP{100}; static constexpr auto BIG_SIZE{100}; /* This API verifies hipMalloc3D API by allocating memory in smaller chunks for -CHUNK_LOOP iterations and checks for the memory leaks by get the memory -info before and after the hipMalloc3D API and the difference should -match with the allocated memory +CHUNK_LOOP iterations */ static void MemoryAlloc3DDiffSizes(int gpu) { HIPCHECK(hipSetDevice(gpu)); @@ -53,10 +51,6 @@ static void MemoryAlloc3DDiffSizes(int gpu) { for (int i = 0; i < CHUNK_LOOP; i++) { HIPCHECK(hipFree(devPitchedPtr[i].ptr)); } - HIPCHECK(hipMemGetInfo(&avail, &tot)); - if ((pavail != avail)) { - HIPASSERT(false); - } } } @@ -95,8 +89,6 @@ TEST_CASE("Unit_hipMalloc3D_MultiThread") { devCnt = HipTest::getDeviceCount(); - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); for (int i = 0; i < devCnt; i++) { threadlist.push_back(std::thread(Malloc3DThreadFunc, i)); } @@ -104,10 +96,4 @@ TEST_CASE("Unit_hipMalloc3D_MultiThread") { for (auto &t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - - if (pavail != avail) { - WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); - REQUIRE(false); - } } diff --git a/projects/hip-tests/catch/unit/memory/hipMemGetInfo.cc b/projects/hip-tests/catch/unit/memory/hipMemGetInfo.cc index 80feaeb239..0aa34fbfaf 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemGetInfo.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemGetInfo.cc @@ -17,563 +17,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ - #include -#include -#include - -/* - * This testcase verifies hipMemGetInfo API - * 1. Different memory chunk allocation - * 1.1. hipMalloc - smallest memory chunck that can be allocated is 1024 - * 1.2. hipMallocArray - * 1.3. hipMalloc3D - * 1.3. hipMalloc3DArray - * 2. Allocation using different threads - * 3. Negative: Invalid args - * - */ - -struct MinAlloc { - private: - int value; - MinAlloc() { - size_t freeMemInit; - size_t totalMemInit; - - unsigned int* A_mem{nullptr}; - size_t mallocSize{1}; - - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - // allocate 1 byte - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), mallocSize)); - - size_t freeMemRet; - size_t totalMemRet; - // actual allocation should be bigger to reflect the minimum allocation on device - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - REQUIRE(freeMemInit >= freeMemRet); - HIP_CHECK(hipFree(A_mem)); - - // store the size of minimum allocation - value = (freeMemInit - freeMemRet); - } - - public: - static int Get() { - static MinAlloc instance; - return instance.value; - } -}; - -// if the memory being allocated is not divisible by the minimum allocation add an extra minimum -// allocation AddedAllocation = InitialAllocation + (MinAllocation - divisionRemainer) -void fixAllocSize(size_t& allocation) { - REQUIRE(MinAlloc::Get() >= 0); - if (allocation % MinAlloc::Get() != 0) { - auto adjustment = allocation % MinAlloc::Get(); // FIXME This does mod by zero - adjustment = MinAlloc::Get() - adjustment; - allocation = allocation + adjustment; - } -} - -// Print information about memory -#define MEMINFO(totalMem, freeMemInit, freeMemRet, usedMem) \ - INFO("Total memory: \t\t\t" << totalMem << "\n" \ - << "Memory used: \t\t\t\t" << freeMemInit - freeMemRet << "\n" \ - << "Free memory after alloc: \t\t" << freeMemRet << "\n" \ - << "Free memory initally: \t\t" << freeMemInit << "\n" \ - << "Memory assumed to be used: \t\t" << usedMem); - - -TEST_CASE("Unit_hipMemGetInfo_DifferentMallocSmall") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - - unsigned int* A_mem{nullptr}; - size_t freeMemRet; - size_t totalMemRet; - // allocate smaller chunk than minimum - size_t Malloc1Size = 2; - - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); - - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size); - - auto assumedFreeMem = freeMemInit - Malloc1Size; - // Free memory should be less than assumed for - // single allocation smaller than min allocation chunk - REQUIRE(freeMemRet < assumedFreeMem); - // confirms that allocated memory is at least equal to smallest allocation - assumedFreeMem = freeMemInit - MinAlloc::Get(); - REQUIRE(freeMemRet <= assumedFreeMem); - - HIP_CHECK(hipFree(A_mem)); - - // allocate smallest chunk of memory - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MinAlloc::Get())); - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - - MEMINFO(totalMemRet, freeMemInit, freeMemRet, MinAlloc::Get()); - - assumedFreeMem = freeMemInit - MinAlloc::Get(); - // confirms that allocated memory is at least equal to smallest allocation - REQUIRE(freeMemRet <= assumedFreeMem); - - HIP_CHECK(hipFree(A_mem)); -} - -#if 0 // FIXME_jatinx Disabled for now because the formula to calulcate memget info is incorrect - // To be enabled after correct formula is found. - -TEST_CASE("Unit_hipMemGetInfo_DifferentMallocLarge") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - - unsigned int* A_mem{nullptr}; - unsigned int* B_mem{nullptr}; - - size_t freeMemRet; - size_t totalMemRet; - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t prop; - HIP_CHECK(hipGetDeviceProperties(&prop, device)); - auto totalMemory = prop.totalGlobalMem; - - - // allocate half of free mem - auto Malloc1Size = freeMemInit >> 1; - // if the allocation is not divisible by the MinAllocation - // take into account and add padding - fixAllocSize(Malloc1Size); - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); - - // allocate an extra quarter of free mem - auto Malloc2Size = Malloc1Size >> 1; - fixAllocSize(Malloc2Size); - HIP_CHECK(hipMalloc(reinterpret_cast(&B_mem), Malloc2Size)); - - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - - MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size + Malloc2Size); - // check if device property total memory is the same as - // total memory returned from hipMemGetInfo - REQUIRE(totalMemory == totalMemRet); - auto allocSize = Malloc1Size + Malloc2Size; - auto assumedFreeMem = freeMemInit - allocSize; - - REQUIRE(freeMemRet <= assumedFreeMem); - HIP_CHECK(hipFree(A_mem)); - HIP_CHECK(hipFree(B_mem)); -} - - -TEST_CASE("Unit_hipMemGetInfo_DifferentMallocMultiSmall") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - - unsigned int* A_mem{nullptr}; - unsigned int* B_mem{nullptr}; - size_t freeMemRet; - size_t totalMemRet; - - // Allocate memory that is a quarter of the min allocation - // Expected behaviour is to reuse the min allocation memory - size_t MallocSize = MinAlloc::Get() >> 2; - - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); - HIP_CHECK(hipMalloc(reinterpret_cast(&B_mem), MallocSize)); - - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - MEMINFO(totalMemRet, freeMemInit, freeMemRet, MallocSize * 2); - - - auto assumedFreeMem = freeMemInit - (MallocSize * 2); - // freeMemRet should be FreeMem - (1 * MinAlloc) - // instead of FreeMem - (MinAlloc * 2) - // since MinAlloc > MallocSize*2 - REQUIRE(freeMemRet < assumedFreeMem); - fixAllocSize(MallocSize); - assumedFreeMem = freeMemInit - (MallocSize * 2); - // Ensure memory allocated is less than 2 * minimum allocation - REQUIRE(freeMemRet > assumedFreeMem); - - // confirms that allocated memory is at least equal to Min Allocation - assumedFreeMem = freeMemInit - MinAlloc::Get(); - REQUIRE(freeMemRet <= assumedFreeMem); - HIP_CHECK(hipFree(A_mem)); - HIP_CHECK(hipFree(B_mem)); -} - -TEST_CASE("Unit_hipMemGetInfo_DifferentMallocNotDiv") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - - unsigned int* A_mem{nullptr}; - size_t freeMemRet; - size_t totalMemRet; - // Allocate memory that is just a bit larger than the min allocation - // Expected behaviour is to allocate 2x min allocation size - size_t MallocSize = MinAlloc::Get() + 1; - - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); - - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - MEMINFO(totalMemRet, freeMemInit, freeMemRet, MallocSize); - - - auto freeMemExpected = freeMemInit - MallocSize; - // Free Memory after allocation should be less than - // expected free memory - REQUIRE(freeMemRet < freeMemExpected); - // confirms that allocated memory is at least 2 x Min Allocaton - fixAllocSize(MallocSize); - freeMemExpected = freeMemInit - MallocSize; - REQUIRE(freeMemRet <= freeMemExpected); - HIP_CHECK(hipFree(A_mem)); -} - - -TEMPLATE_TEST_CASE("Unit_hipMemGetInfo_MallocArray", "", int, int4, char) { - // get initial mem data - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - - // create and allocate an Array - hipArray_t arrayPtr{}; - - auto bytesPerItem = sizeof(TestType); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - hipExtent extent{}; - extent.width = GENERATE(32, 128, 256, 512, 1024); - - extent.height = GENERATE(0, 32, 128, 256, 512, 1024); - - HIP_CHECK(hipMallocArray(&arrayPtr, &desc, extent.width, extent.height, hipArrayDefault)); - - // check if memory is correct - size_t freeMemRet; - size_t totalMemRet; - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - - // calculate used memory, take into account 1D array (height = 0) - size_t usedMem = bytesPerItem * extent.width * (extent.height != 0 ? extent.height : 1); - - // ensure we allocate at least the min allocation for the array - fixAllocSize(usedMem); - MEMINFO(totalMemRet, freeMemInit, freeMemRet, usedMem); - - size_t assumedFreeMem = freeMemInit - usedMem; - - REQUIRE(freeMemRet <= assumedFreeMem); - - HIP_CHECK(hipFreeArray(arrayPtr)); -} - -TEST_CASE("Unit_hipMemGetInfo_Malloc3D") { - // Get initial memory - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - - // Allocate 3D object - hipExtent extent{}; - // extent is given in bytes for with - extent.width = GENERATE(32, 128, 256, 512); - extent.height = GENERATE(32, 128, 256, 512); - extent.depth = GENERATE(32, 128, 256, 512); - hipPitchedPtr A_mem{}; - HIP_CHECK(hipMalloc3D(&A_mem, extent)); - - // Get memory after allocation - size_t freeMemRet; - size_t totalMemRet; - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - - // Verify result - size_t mallocSize = A_mem.pitch * extent.height * extent.depth; - fixAllocSize(mallocSize); - - size_t assumedFreeMem = freeMemInit - mallocSize; - MEMINFO(totalMemRet, freeMemInit, freeMemRet, mallocSize); - - REQUIRE(freeMemRet <= assumedFreeMem); - - HIP_CHECK(hipFree(A_mem.ptr)); -} - -TEMPLATE_TEST_CASE("Unit_hipMemGetInfo_Malloc3DArray", "", char, int, int4) { - // Get initial memory - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - // Allocate 3D object - hipArray_t arrayPtr{}; - size_t sizeInBytes = (size_t)sizeof(TestType); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - - int device; - HIP_CHECK(hipGetDevice(&device)); - int allignSize{0}; - hipDeviceGetAttribute(&allignSize, hipDeviceAttributeTextureAlignment, device); - -#if HT_NVIDIA - auto flag = GENERATE(hipArrayDefault, hipArrayLayered, hipArrayCubemap, - hipArrayLayered | hipArrayCubemap); -#else - // hipArrayCubemap not supported on AMD - auto flag = GENERATE(hipArrayDefault, hipArrayLayered); -#endif - - hipExtent extent{}; - extent.width = GENERATE(32, 128, 256, 512); - extent.height = GENERATE(0, 32, 128, 256, 512); - if (flag == hipArrayCubemap) { - // width must be equal to height, and depth must be six. - extent.height = extent.width; - extent.depth = 6; - } else if (flag == hipArrayLayered | hipArrayCubemap) { - // width must be equal to height, and depth must be a multiple six. - extent.height = extent.width; - extent.depth = 6 * GENERATE(4, 8, 16, 32); - } else if (extent.height == 0 && flag != hipArrayLayered) { - // if height = 0 the depth must be 0 unless using hipArrayLayered flag - extent.depth = 0; - } else { - extent.depth = GENERATE(32, 128, 256, 512); - } - - - // Get memory after allocation - auto h = extent.height == 0 ? 1 : extent.height; - auto d = extent.depth == 0 ? 1 : extent.depth; - auto w = extent.width * sizeInBytes; - size_t mallocSize = w * h * d; - - HIP_CHECK(hipMalloc3DArray(&arrayPtr, &desc, extent, flag)); - - // Verify result - size_t freeMemRet; - size_t totalMemRet; - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - - // Sometimes hipMemGetInfo reports that no new memory has be allocated for testcase - // take this into account - if (freeMemInit == freeMemRet) { - // no new memory allocation has occured verify that memory trying - // to be allocated is less than a min allocation block - MEMINFO(totalMemRet, freeMemInit, freeMemRet, mallocSize); - REQUIRE(mallocSize <= static_cast(MinAlloc::Get())); - - } else { - // account for min allocation - fixAllocSize(mallocSize); - - MEMINFO(totalMemRet, freeMemInit, freeMemRet, mallocSize); - size_t assumedFreeMem = freeMemInit - mallocSize; - REQUIRE(freeMemRet <= assumedFreeMem); - } - HIP_CHECK(hipFreeArray(arrayPtr)); -} - - -TEST_CASE("Unit_hipMemGetInfo_ParaLarge") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - unsigned int* A_mem{nullptr}; - unsigned int* B_mem{nullptr}; - - // allocate half of free mem - auto Malloc1Size = freeMemInit >> 1; - // if the allocation is not divisible by the MinAllocation - // take into account and add padding - fixAllocSize(Malloc1Size); - std::thread t1( - [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); }); - - // allocate an extra quarter of free mem - auto Malloc2Size = Malloc1Size >> 1; - fixAllocSize(Malloc2Size); - std::thread t2( - [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&B_mem), Malloc2Size)); }); - - t1.join(); - t2.join(); - HIP_CHECK_THREAD_FINALIZE(); - - size_t freeMemRet; - size_t totalMemRet; - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - - MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size + Malloc2Size); - auto allocSize = Malloc1Size + Malloc2Size; - REQUIRE(freeMemRet <= freeMemInit - allocSize); - - HIP_CHECK(hipFree(A_mem)); - HIP_CHECK(hipFree(B_mem)); -} - -#endif - -TEST_CASE("Unit_hipMemGetInfo_ParaSmall") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - unsigned int* A_mem{nullptr}; - // allocate smaller chunk than minimum - size_t Malloc1Size = 2; - - std::thread t1( - [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)) }); - t1.join(); - HIP_CHECK_THREAD_FINALIZE(); - size_t freeMemRet; - size_t totalMemRet; - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size); - - - auto assumedFreeMem = freeMemInit - Malloc1Size; - // Free memory should be less than assumed for - // single allocation smaller than min allocation chunk - REQUIRE(freeMemRet < assumedFreeMem); - // confirms that allocated memory is at least equal to smallest allocation allowed - assumedFreeMem = freeMemInit - MinAlloc::Get(); - REQUIRE(freeMemRet <= assumedFreeMem); - - HIP_CHECK(hipFree(A_mem)); - - // allocate smallest chunck of memory - std::thread t2( - [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), MinAlloc::Get())); }); - t2.join(); - HIP_CHECK_THREAD_FINALIZE(); - - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - - MEMINFO(totalMemRet, freeMemInit, freeMemRet, MinAlloc::Get()); - - assumedFreeMem = freeMemInit - MinAlloc::Get(); - REQUIRE(freeMemRet <= assumedFreeMem); - - HIP_CHECK(hipFree(A_mem)); -} - -TEST_CASE("Unit_hipMemGetInfo_ParaNonDiv") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - unsigned int* A_mem{nullptr}; - - // Allocate memory that is just 1 byte larger than the min allocation - // Expected behaviour is to allocate 2x min allocation size - size_t Malloc1Size = MinAlloc::Get() + 1; - - std::thread t1( - [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); }); - t1.join(); - HIP_CHECK_THREAD_FINALIZE(); - - size_t freeMemRet; - size_t totalMemRet; - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size); - - - auto allocSize = freeMemInit - Malloc1Size; - // should not be equal - REQUIRE(freeMemRet != allocSize); - // confirms that allocated memory is equal to 2 x Min Allocaton - allocSize = MinAlloc::Get() * 2; - auto assumedAllocSize = freeMemInit - allocSize; - REQUIRE(freeMemRet <= assumedAllocSize); - HIP_CHECK(hipFree(A_mem)); -} - -TEST_CASE("Unit_hipMemGetInfo_ParaMultiSmall") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - unsigned int* A_mem{nullptr}; - unsigned int* B_mem{nullptr}; - - // Allocate memory that is a quarter of the min allocation - // Expected behaviour is to reuse the min allocation memory - size_t MallocSize = MinAlloc::Get() >> 2; - - std::thread t1( - [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); }); - std::thread t2( - [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&B_mem), MallocSize)); }); - - t1.join(); - t2.join(); - HIP_CHECK_THREAD_FINALIZE(); - - size_t freeMemRet; - size_t totalMemRet; - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - MEMINFO(totalMemRet, freeMemInit, freeMemRet, MallocSize * 2); - - auto assumedFreeMem = freeMemInit - MallocSize * 2; - // freeMemRet should be less than assumedFreeMem - REQUIRE(freeMemRet < assumedFreeMem); - // confirms that allocated memory is equal to Min Allocation - assumedFreeMem = freeMemInit - MinAlloc::Get(); - REQUIRE(freeMemRet <= assumedFreeMem); - HIP_CHECK(hipFree(A_mem)); - HIP_CHECK(hipFree(B_mem)); -} - - -TEST_CASE("Unit_hipMemGetInfo_Negative") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - - unsigned int* A_mem{nullptr}; - auto MallocSize = MinAlloc::Get(); - - SECTION("Zero allocation") { - size_t freeMemRet; - size_t totalMemRet; - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), 0)); - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - - REQUIRE(freeMemRet == freeMemInit); - } - SECTION("Nullptr as first param passed to hipMemGetInfo") { - size_t* freeMemRet = nullptr; - size_t totalMemRet; - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); - // Segfaults on AMD and returns hipSuccess on Nvidia - HIP_CHECK(hipMemGetInfo(freeMemRet, &totalMemRet)); - } - SECTION("Nullptr as second param passed to hipMemGetInfo") { - size_t freeMemRet; - size_t* totalMemRet = nullptr; - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); - // Segfaults on AMD and returns hipSuccess on Nvidia - HIP_CHECK(hipMemGetInfo(&freeMemRet, totalMemRet)); - } - SECTION("Nullptr as both params passed to hipMemGetInfo") { - size_t* freeMemRet = nullptr; - size_t* totalMemRet = nullptr; - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); - // Segfaults on AMD and returns hipSuccess on Nvidia - HIP_CHECK(hipMemGetInfo(freeMemRet, totalMemRet)); - } - - HIP_CHECK(hipFree(A_mem)); -} TEST_CASE("Unit_hipMemGetInfo_FreeLessThanTotal") { unsigned int *A_mem{nullptr}; From f8e1cba1ad259e1262f50e39b13c8976f5ce0ec9 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 20 Jul 2023 10:22:17 +0530 Subject: [PATCH 16/17] SWDEV-402381 - Add hipCheckErrors for HIP API calls in samples (#375) Change-Id: I335d7e780362fc59fd2d90939b4c8b8a7231ffc7 [ROCm/hip-tests commit: 7cc53f992ffc070b7dea7dae82f99db161f59098] --- .../0_Intro/bit_extract/CMakeLists.txt | 2 + .../samples/0_Intro/bit_extract/Makefile | 3 +- .../0_Intro/bit_extract/bit_extract.cpp | 29 ++++------- .../samples/0_Intro/module_api/CMakeLists.txt | 2 + .../samples/0_Intro/module_api/Makefile | 9 ++-- .../0_Intro/module_api/defaultDriver.cpp | 27 +++++----- .../0_Intro/module_api/launchKernelHcc.cpp | 33 ++++++------- .../samples/0_Intro/module_api/runKernel.cpp | 31 +++++------- .../0_Intro/module_api_global/CMakeLists.txt | 2 + .../0_Intro/module_api_global/Makefile | 3 +- .../0_Intro/module_api_global/runKernel.cpp | 24 ++++----- .../1_Utils/hipDispatchLatency/CMakeLists.txt | 2 + .../1_Utils/hipDispatchLatency/Makefile | 3 +- .../hipDispatchEnqueueRateMT.cpp | 27 ++++------ .../hipDispatchLatency/hipDispatchLatency.cpp | 37 +++++++------- .../samples/1_Utils/hipInfo/CMakeLists.txt | 2 + .../samples/1_Utils/hipInfo/Makefile | 3 +- .../samples/1_Utils/hipInfo/hipInfo.cpp | 29 +++-------- .../0_MatrixTranspose/CMakeLists.txt | 2 + .../2_Cookbook/0_MatrixTranspose/Makefile | 3 +- .../0_MatrixTranspose/MatrixTranspose.cpp | 15 +++--- .../2_Cookbook/10_inline_asm/CMakeLists.txt | 2 + .../samples/2_Cookbook/10_inline_asm/Makefile | 4 +- .../2_Cookbook/10_inline_asm/inline_asm.cpp | 43 ++++++++-------- .../11_texture_driver/CMakeLists.txt | 2 + .../2_Cookbook/11_texture_driver/Makefile | 3 +- .../11_texture_driver/texture2dDrv.cpp | 43 +++++++--------- .../CMakeLists.txt | 2 + .../MatrixTranspose.cpp | 15 +++--- .../2_Cookbook/13_occupancy/CMakeLists.txt | 2 + .../samples/2_Cookbook/13_occupancy/Makefile | 4 +- .../2_Cookbook/13_occupancy/occupancy.cpp | 49 +++++++++---------- .../2_Cookbook/14_gpu_arch/CMakeLists.txt | 2 + .../samples/2_Cookbook/14_gpu_arch/Makefile | 4 +- .../2_Cookbook/14_gpu_arch/gpuarch.cpp | 14 ++---- .../16_assembly_to_executable/Makefile | 5 +- .../16_assembly_to_executable/square.cpp | 28 ++++------- .../17_llvm_ir_to_executable/Makefile | 5 +- .../17_llvm_ir_to_executable/square.cpp | 28 ++++------- .../18_cmake_hip_device/CMakeLists.txt | 4 ++ .../2_Cookbook/18_cmake_hip_device/square.cpp | 29 ++++------- .../2_Cookbook/19_cmake_lang/CMakeLists.txt | 3 ++ .../19_cmake_lang/MatrixTranspose.cpp | 16 +++--- .../2_Cookbook/1_hipEvent/CMakeLists.txt | 2 + .../samples/2_Cookbook/1_hipEvent/Makefile | 3 +- .../2_Cookbook/1_hipEvent/hipEvent.cpp | 43 ++++++++-------- .../21_cmake_hip_cxx_clang/CMakeLists.txt | 2 + .../21_cmake_hip_cxx_clang/square.cpp | 27 ++++------ .../22_cmake_hip_lang/CMakeLists.txt | 2 + .../2_Cookbook/22_cmake_hip_lang/square.hip | 27 ++++------ .../2_Cookbook/23_cmake_hiprtc/CMakeLists.txt | 2 + .../2_Cookbook/23_cmake_hiprtc/saxpy.cpp | 31 ++++++------ .../2_Cookbook/3_shared_memory/CMakeLists.txt | 2 + .../2_Cookbook/3_shared_memory/Makefile | 3 +- .../3_shared_memory/sharedMemory.cpp | 16 +++--- .../samples/2_Cookbook/4_shfl/CMakeLists.txt | 2 + .../samples/2_Cookbook/4_shfl/Makefile | 3 +- .../samples/2_Cookbook/4_shfl/shfl.cpp | 16 +++--- .../samples/2_Cookbook/5_2dshfl/2dshfl.cpp | 15 +++--- .../2_Cookbook/5_2dshfl/CMakeLists.txt | 2 + .../samples/2_Cookbook/5_2dshfl/Makefile | 3 +- .../6_dynamic_shared/CMakeLists.txt | 4 ++ .../2_Cookbook/6_dynamic_shared/Makefile | 3 +- .../6_dynamic_shared/dynamic_shared.cpp | 15 +++--- .../2_Cookbook/7_streams/CMakeLists.txt | 2 + .../samples/2_Cookbook/7_streams/Makefile | 3 +- .../samples/2_Cookbook/7_streams/stream.cpp | 25 +++++----- .../2_Cookbook/9_unroll/CMakeLists.txt | 2 + .../samples/2_Cookbook/9_unroll/Makefile | 3 +- .../samples/2_Cookbook/9_unroll/unroll.cpp | 17 ++++--- .../hip-tests/samples/common/hip_helper.h | 38 ++++++++++++++ 71 files changed, 460 insertions(+), 448 deletions(-) create mode 100644 projects/hip-tests/samples/common/hip_helper.h diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/CMakeLists.txt b/projects/hip-tests/samples/0_Intro/bit_extract/CMakeLists.txt index d51e4d974a..4986dc1b1f 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/CMakeLists.txt +++ b/projects/hip-tests/samples/0_Intro/bit_extract/CMakeLists.txt @@ -48,5 +48,7 @@ set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) # Create the excutable add_executable(bit_extract bit_extract.cpp) +target_include_directories(bit_extract PRIVATE ../../common) + # Link with HIP target_link_libraries(bit_extract hip::host) diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile index 6f4d824ba8..939aca00a5 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile +++ b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile @@ -29,6 +29,7 @@ ifeq (,$(HIP_PATH)) endif HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) HIPCC=$(HIP_PATH)/bin/hipcc +INCLUDES := -I../../common # Show how to use PLATFORM to specify different options for each compiler: ifeq (${HIP_PLATFORM}, nvcc) @@ -38,7 +39,7 @@ endif EXE=bit_extract $(EXE): bit_extract.cpp - $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + $(HIPCC) $(HIPCC_FLAGS) $(INCLUDES) $< -o $@ all: $(EXE) diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp b/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp index cf0440dd57..3f7636bc08 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp +++ b/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp @@ -23,16 +23,7 @@ THE SOFTWARE. #include #include #include "hip/hip_runtime.h" - -#define CHECK(cmd) \ - { \ - hipError_t error = cmd; \ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \ - __FILE__, __LINE__); \ - exit(EXIT_FAILURE); \ - } \ - } +#include "hip_helper.h" __global__ void bit_extract_kernel(uint32_t* C_d, const uint32_t* A_d, size_t N) { size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); @@ -69,28 +60,28 @@ int main(int argc, char* argv[]) { #endif int deviceId; - CHECK(hipGetDevice(&deviceId)); + checkHipErrors(hipGetDevice(&deviceId)); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, deviceId)); + checkHipErrors(hipGetDeviceProperties(&props, deviceId)); printf("info: running on device #%d %s\n", deviceId, props.name); printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); A_h = (uint32_t*)malloc(Nbytes); - CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess); + checkHipErrors(A_h == 0 ? hipErrorOutOfMemory : hipSuccess); C_h = (uint32_t*)malloc(Nbytes); - CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess); + checkHipErrors(C_h == 0 ? hipErrorOutOfMemory : hipSuccess); for (size_t i = 0; i < N; i++) { A_h[i] = i; } printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); - CHECK(hipMalloc(&A_d, Nbytes)); - CHECK(hipMalloc(&C_d, Nbytes)); + checkHipErrors(hipMalloc(&A_d, Nbytes)); + checkHipErrors(hipMalloc(&C_d, Nbytes)); printf("info: copy Host2Device\n"); - CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + checkHipErrors(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); printf("info: launch 'bit_extract_kernel' \n"); const unsigned blocks = 512; @@ -98,7 +89,7 @@ int main(int argc, char* argv[]) { hipLaunchKernelGGL(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); printf("info: copy Device2Host\n"); - CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + checkHipErrors(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); printf("info: check result\n"); for (size_t i = 0; i < N; i++) { @@ -106,7 +97,7 @@ int main(int argc, char* argv[]) { if (C_h[i] != Agold) { fprintf(stderr, "mismatch detected.\n"); printf("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]); - CHECK(hipErrorUnknown); + checkHipErrors(hipErrorUnknown); } } printf("PASSED!\n"); diff --git a/projects/hip-tests/samples/0_Intro/module_api/CMakeLists.txt b/projects/hip-tests/samples/0_Intro/module_api/CMakeLists.txt index cefe6e2c79..7106df44c0 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/CMakeLists.txt +++ b/projects/hip-tests/samples/0_Intro/module_api/CMakeLists.txt @@ -22,6 +22,8 @@ project(module_api) cmake_minimum_required(VERSION 3.10) +include_directories(../../common) + if (NOT DEFINED ROCM_PATH ) set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." ) endif () diff --git a/projects/hip-tests/samples/0_Intro/module_api/Makefile b/projects/hip-tests/samples/0_Intro/module_api/Makefile index 118d16a7ef..ee3e68e067 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/Makefile +++ b/projects/hip-tests/samples/0_Intro/module_api/Makefile @@ -27,20 +27,21 @@ ifeq (,$(HIP_PATH)) endif HIPCC=$(HIP_PATH)/bin/hipcc HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) +INCLUDES := -I../../common all: vcpy_kernel.code runKernel.hip.out launchKernelHcc.hip.out defaultDriver.hip.out runKernel.hip.out: runKernel.cpp - $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + $(HIPCC) $(HIPCC_FLAGS) $(INCLUDES) $< -o $@ launchKernelHcc.hip.out: launchKernelHcc.cpp - $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + $(HIPCC) $(HIPCC_FLAGS) $(INCLUDES) $< -o $@ defaultDriver.hip.out: defaultDriver.cpp - $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + $(HIPCC) $(HIPCC_FLAGS) $(INCLUDES) $< -o $@ vcpy_kernel.code: vcpy_kernel.cpp - $(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@ + $(HIPCC) --genco $(GENCO_FLAGS) $(INCLUDES) $^ -o $@ clean: rm -f *.code *.out diff --git a/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp b/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp index 9443842026..fc759f3f1c 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. #include #include #include +#include "hip_helper.h" #define LEN 64 #define SIZE LEN << 2 @@ -45,25 +46,25 @@ int main() { hipInit(0); hipDevice_t device; hipCtx_t context; - hipDeviceGet(&device, 0); - hipCtxCreate(&context, 0, device); + checkHipErrors(hipDeviceGet(&device, 0)); + checkHipErrors(hipCtxCreate(&context, 0, device)); - hipMalloc((void**)&Ad, SIZE); - hipMalloc((void**)&Bd, SIZE); + checkHipErrors(hipMalloc((void**)&Ad, SIZE)); + checkHipErrors(hipMalloc((void**)&Bd, SIZE)); - hipMemcpyHtoD(Ad, A, SIZE); - hipMemcpyHtoD(Bd, B, SIZE); + checkHipErrors(hipMemcpyHtoD(Ad, A, SIZE)); + checkHipErrors(hipMemcpyHtoD(Bd, B, SIZE)); hipModule_t Module; hipFunction_t Function; - hipModuleLoad(&Module, fileName); - hipModuleGetFunction(&Function, Module, kernel_name); + checkHipErrors(hipModuleLoad(&Module, fileName)); + checkHipErrors(hipModuleGetFunction(&Function, Module, kernel_name)); void* args[2] = {&Ad, &Bd}; - hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr); + checkHipErrors(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr)); - hipMemcpyDtoH(B, Bd, SIZE); + checkHipErrors(hipMemcpyDtoH(B, Bd, SIZE)); int mismatchCount = 0; for (uint32_t i = 0; i < LEN; i++) { if (A[i] != B[i]) { @@ -78,10 +79,10 @@ int main() { std::cout << "FAILED!\n"; }; - hipFree(Ad); - hipFree(Bd); + checkHipErrors(hipFree(Ad)); + checkHipErrors(hipFree(Bd)); delete[] A; delete[] B; - hipCtxDestroy(context); + checkHipErrors(hipCtxDestroy(context)); return 0; } diff --git a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp index 464f6d8851..3c44a375c7 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -25,6 +25,7 @@ THE SOFTWARE. #include #include #include +#include "hip_helper.h" #ifdef __HIP_PLATFORM_AMD__ #include @@ -36,12 +37,6 @@ THE SOFTWARE. #define fileName "vcpy_kernel.code" #define kernel_name "hello_world" -#define HIP_CHECK(status) \ - if (status != hipSuccess) { \ - std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \ - exit(0); \ - } - int main() { float *A, *B; hipDeviceptr_t Ad, Bd; @@ -56,18 +51,18 @@ int main() { hipInit(0); hipDevice_t device; hipCtx_t context; - hipDeviceGet(&device, 0); - hipCtxCreate(&context, 0, device); + checkHipErrors(hipDeviceGet(&device, 0)); + checkHipErrors(hipCtxCreate(&context, 0, device)); - hipMalloc((void**)&Ad, SIZE); - hipMalloc((void**)&Bd, SIZE); + checkHipErrors(hipMalloc((void**)&Ad, SIZE)); + checkHipErrors(hipMalloc((void**)&Bd, SIZE)); - hipMemcpyHtoD(Ad, A, SIZE); - hipMemcpyHtoD(Bd, B, SIZE); + checkHipErrors(hipMemcpyHtoD(Ad, A, SIZE)); + checkHipErrors(hipMemcpyHtoD(Bd, B, SIZE)); hipModule_t Module; hipFunction_t Function; - HIP_CHECK(hipModuleLoad(&Module, fileName)); - HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + checkHipErrors(hipModuleLoad(&Module, fileName)); + checkHipErrors(hipModuleGetFunction(&Function, Module, kernel_name)); struct { void* _Ad; @@ -83,10 +78,10 @@ int main() { void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; - HIP_CHECK( + checkHipErrors( hipExtModuleLaunchKernel(Function, LEN, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config, 0)); - hipMemcpyDtoH(B, Bd, SIZE); + checkHipErrors(hipMemcpyDtoH(B, Bd, SIZE)); int mismatchCount = 0; for (uint32_t i = 0; i < LEN; i++) { @@ -102,10 +97,10 @@ int main() { std::cout << "FAILED!\n"; }; - hipFree(Ad); - hipFree(Bd); + checkHipErrors(hipFree(Ad)); + checkHipErrors(hipFree(Bd)); delete[] A; delete[] B; - hipCtxDestroy(context); + checkHipErrors(hipCtxDestroy(context)); return 0; } diff --git a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp index c2de0c6c0d..f4af0332fa 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp @@ -26,6 +26,7 @@ THE SOFTWARE. #include #include #include +#include "hip_helper.h" #define LEN 64 #define SIZE LEN << 2 @@ -33,12 +34,6 @@ THE SOFTWARE. #define fileName "vcpy_kernel.code" #define kernel_name "hello_world" -#define HIP_CHECK(status) \ - if (status != hipSuccess) { \ - std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \ - exit(0); \ - } - int main() { float *A, *B; hipDeviceptr_t Ad, Bd; @@ -53,18 +48,18 @@ int main() { hipInit(0); hipDevice_t device; hipCtx_t context; - hipDeviceGet(&device, 0); - hipCtxCreate(&context, 0, device); + checkHipErrors(hipDeviceGet(&device, 0)); + checkHipErrors(hipCtxCreate(&context, 0, device)); - hipMalloc((void**)&Ad, SIZE); - hipMalloc((void**)&Bd, SIZE); + checkHipErrors(hipMalloc((void**)&Ad, SIZE)); + checkHipErrors(hipMalloc((void**)&Bd, SIZE)); - hipMemcpyHtoD(Ad, A, SIZE); - hipMemcpyHtoD(Bd, B, SIZE); + checkHipErrors(hipMemcpyHtoD(Ad, A, SIZE)); + checkHipErrors(hipMemcpyHtoD(Bd, B, SIZE)); hipModule_t Module; hipFunction_t Function; - HIP_CHECK(hipModuleLoad(&Module, fileName)); - HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + checkHipErrors(hipModuleLoad(&Module, fileName)); + checkHipErrors(hipModuleGetFunction(&Function, Module, kernel_name)); struct { void* _Ad; @@ -79,9 +74,9 @@ int main() { void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; - HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config)); + checkHipErrors(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config)); - hipMemcpyDtoH(B, Bd, SIZE); + checkHipErrors(hipMemcpyDtoH(B, Bd, SIZE)); int mismatchCount = 0; for (uint32_t i = 0; i < LEN; i++) { @@ -97,10 +92,10 @@ int main() { std::cout << "FAILED!\n"; }; - hipFree(Ad); + checkHipErrors(hipFree(Ad)); hipFree(Bd); delete[] A; delete[] B; - hipCtxDestroy(context); + checkHipErrors(hipCtxDestroy(context)); return 0; } diff --git a/projects/hip-tests/samples/0_Intro/module_api_global/CMakeLists.txt b/projects/hip-tests/samples/0_Intro/module_api_global/CMakeLists.txt index c3b147c60c..186a99d164 100644 --- a/projects/hip-tests/samples/0_Intro/module_api_global/CMakeLists.txt +++ b/projects/hip-tests/samples/0_Intro/module_api_global/CMakeLists.txt @@ -50,5 +50,7 @@ add_custom_target( add_dependencies(runKernel.hip.out codeobj) +target_include_directories(runKernel.hip.out PRIVATE ../../common) + # Link with HIP target_link_libraries(runKernel.hip.out hip::host) \ No newline at end of file diff --git a/projects/hip-tests/samples/0_Intro/module_api_global/Makefile b/projects/hip-tests/samples/0_Intro/module_api_global/Makefile index fdbc77f65f..732eec63aa 100644 --- a/projects/hip-tests/samples/0_Intro/module_api_global/Makefile +++ b/projects/hip-tests/samples/0_Intro/module_api_global/Makefile @@ -27,11 +27,12 @@ ifeq (,$(HIP_PATH)) endif HIPCC=$(HIP_PATH)/bin/hipcc HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) +INCLUDES := -I../../common all: vcpy_kernel.code runKernel.hip.out runKernel.hip.out: runKernel.cpp - $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + $(HIPCC) $(HIPCC_FLAGS) $(INCLUDES) $< -o $@ vcpy_kernel.code: vcpy_kernel.cpp $(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@ diff --git a/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp index 23f4ec2750..40711dd4b2 100644 --- a/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp @@ -31,7 +31,7 @@ THE SOFTWARE. #define SIZE LEN * sizeof(float) #define fileName "vcpy_kernel.code" -#define HIP_CHECK(cmd) \ +#define checkHipErrors(cmd) \ { \ hipError_t status = cmd; \ if (status != hipSuccess) { \ @@ -64,23 +64,23 @@ int main() { hipMemcpyHtoD(hipDeviceptr_t(Ad), A, SIZE); hipMemcpyHtoD((hipDeviceptr_t)(Bd), B, SIZE); hipModule_t Module; - HIP_CHECK(hipModuleLoad(&Module, fileName)); + checkHipErrors(hipModuleLoad(&Module, fileName)); float myDeviceGlobal_h = 42.0; float* deviceGlobal; size_t deviceGlobalSize; - HIP_CHECK(hipModuleGetGlobal((void**)&deviceGlobal, &deviceGlobalSize, Module, "myDeviceGlobal")); - HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h, deviceGlobalSize)); + checkHipErrors(hipModuleGetGlobal((void**)&deviceGlobal, &deviceGlobalSize, Module, "myDeviceGlobal")); + checkHipErrors(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h, deviceGlobalSize)); #define ARRAY_SIZE 16 float myDeviceGlobalArray_h[ARRAY_SIZE]; float *myDeviceGlobalArray; size_t myDeviceGlobalArraySize; - HIP_CHECK(hipModuleGetGlobal((void**)&myDeviceGlobalArray, &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray")); + checkHipErrors(hipModuleGetGlobal((void**)&myDeviceGlobalArray, &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray")); for (int i = 0; i < ARRAY_SIZE; i++) { myDeviceGlobalArray_h[i] = i * 1000.0f; - HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray), &myDeviceGlobalArray_h, myDeviceGlobalArraySize)); + checkHipErrors(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray), &myDeviceGlobalArray_h, myDeviceGlobalArraySize)); } struct { @@ -98,8 +98,8 @@ int main() { { hipFunction_t Function; - HIP_CHECK(hipModuleGetFunction(&Function, Module, "hello_world")); - HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config)); + checkHipErrors(hipModuleGetFunction(&Function, Module, "hello_world")); + checkHipErrors(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config)); hipMemcpyDtoH(B, Bd, SIZE); @@ -123,13 +123,13 @@ int main() { { hipFunction_t Function; - HIP_CHECK(hipModuleGetFunction(&Function, Module, "test_globals")); + checkHipErrors(hipModuleGetFunction(&Function, Module, "test_globals")); int val =-1; - HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,Function)); + checkHipErrors(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,Function)); printf("Shared Size Bytes = %d\n",val); - HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_NUM_REGS, Function)); + checkHipErrors(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_NUM_REGS, Function)); printf("Num Regs = %d\n",val); - HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config)); + checkHipErrors(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config)); hipMemcpyDtoH(B, Bd, SIZE); diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/CMakeLists.txt b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/CMakeLists.txt index d0a453f1d5..33caab73b7 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/CMakeLists.txt +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/CMakeLists.txt @@ -22,6 +22,8 @@ project(hipDispatchLatency) cmake_minimum_required(VERSION 3.10) +include_directories(../../common) + if (NOT DEFINED ROCM_PATH ) set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." ) endif () diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/Makefile b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/Makefile index 988827071c..6643e3f034 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/Makefile +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/Makefile @@ -26,8 +26,9 @@ ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc -std=c++11 +INCLUDES := -I../../common -CXXFLAGS = -O3 +CXXFLAGS = -O3 $(INCLUDES) all: test_kernel.code hipDispatchLatency.out hipDispatchEnqueueRateMT.out diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp index 9bc3258b4d..aced502d8f 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp @@ -22,6 +22,7 @@ THE SOFTWARE. #ifdef __HIP_PLATFORM_AMD__ #include "hip/hip_ext.h" #endif +#include "hip_helper.h" #include #include #include @@ -41,16 +42,6 @@ THE SOFTWARE. #define failed(...) \ abort(); -#define HIPCHECK(error) \ - { \ - hipError_t localError = error; \ - if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ - printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(localError), \ - localError, #error, __FILE__, __LINE__); \ - failed("API returned error code."); \ - } \ - } - __global__ void EmptyKernel() {} @@ -87,12 +78,12 @@ void hipModuleLaunchKernel_enqueue_rate(const std::vector& buffer, std::at { //resources necessary for this thread hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); + checkHipErrors(hipStreamCreate(&stream)); hipModule_t module; hipFunction_t function; - HIPCHECK(hipModuleLoadData(&module, &buffer[0])); - HIPCHECK(hipModuleGetFunction(&function, module, "test")); + checkHipErrors(hipModuleLoadData(&module, &buffer[0])); + checkHipErrors(hipModuleGetFunction(&function, module, "test")); void* kernel_params = nullptr; std::array results; @@ -103,13 +94,13 @@ void hipModuleLaunchKernel_enqueue_rate(const std::vector& buffer, std::at for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { auto start = std::chrono::high_resolution_clock::now(); - HIPCHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, stream, &kernel_params, nullptr)); + checkHipErrors(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, stream, &kernel_params, nullptr)); auto stop = std::chrono::high_resolution_clock::now(); results[i] = std::chrono::duration(stop - start).count(); } - HIPCHECK(hipModuleUnload(module)); + checkHipErrors(hipModuleUnload(module)); print_timing("Thread ID : " + std::to_string(tid) + " , " + "hipModuleLaunchKernel enqueue rate", results); - HIPCHECK(hipStreamDestroy(stream)); + checkHipErrors(hipStreamDestroy(stream)); } // Measure time taken to enqueue a kernel on the GPU using hipLaunchKernelGGL @@ -117,7 +108,7 @@ void hipLaunchKernelGGL_enqueue_rate(const std::vector& buffer, std::atomi { //resources necessary for this thread hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); + checkHipErrors(hipStreamCreate(&stream)); std::array results; //synchronize all threads, before running @@ -131,7 +122,7 @@ void hipLaunchKernelGGL_enqueue_rate(const std::vector& buffer, std::atomi results[i] = std::chrono::duration(stop - start).count(); } print_timing("Thread ID : " + std::to_string(tid) + " , " + "hipLaunchKernelGGL enqueue rate", results); - HIPCHECK(hipStreamDestroy(stream)); + checkHipErrors(hipStreamDestroy(stream)); } // Simple thread pool diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index 5b96f3a199..525dafb45c 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -21,6 +21,7 @@ THE SOFTWARE. #ifdef __HIP_PLATFORM_AMD__ #include "hip/hip_ext.h" #endif +#include "hip_helper.h" #include #include #include @@ -66,19 +67,19 @@ void print_timing(std::string test, const std::array &re int main() { hipStream_t stream0 = 0; hipDevice_t device; - hipDeviceGet(&device, 0); + checkHipErrors(hipDeviceGet(&device, 0)); hipCtx_t context; - hipCtxCreate(&context, 0, device); + checkHipErrors(hipCtxCreate(&context, 0, device)); hipModule_t module; hipFunction_t function; - hipModuleLoad(&module, FILE_NAME); - hipModuleGetFunction(&function, module, KERNEL_NAME); + checkHipErrors(hipModuleLoad(&module, FILE_NAME)); + checkHipErrors(hipModuleGetFunction(&function, module, KERNEL_NAME)); void* params = nullptr; std::array results; hipEvent_t start, stop; - hipEventCreate(&start); - hipEventCreate(&stop); + checkHipErrors(hipEventCreate(&start)); + checkHipErrors(hipEventCreate(&stop)); /************************************************************************************/ /* HIP kernel launch enqueue rate: */ @@ -88,7 +89,7 @@ int main() { // Timing hipModuleLaunchKernel for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { auto start = std::chrono::high_resolution_clock::now(); - hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, ¶ms, nullptr); + checkHipErrors(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, ¶ms, nullptr)); auto stop = std::chrono::high_resolution_clock::now(); results[i] = std::chrono::duration(stop - start).count(); } @@ -110,11 +111,11 @@ int main() { //Timing around the dispatch for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { - hipEventRecord(start, 0); + checkHipErrors(hipEventRecord(start, 0)); hipLaunchKernelGGL((EmptyKernel), dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream0); - hipEventRecord(stop, 0); - hipEventSynchronize(stop); - hipEventElapsedTime(&results[i], start, stop); + checkHipErrors(hipEventRecord(stop, 0)); + checkHipErrors(hipEventSynchronize(stop)); + checkHipErrors(hipEventElapsedTime(&results[i], start, stop)); } print_timing("Timing around single dispatch latency", results); @@ -124,18 +125,18 @@ int main() { /*********************************************************************************/ for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { - hipEventRecord(start, 0); + checkHipErrors(hipEventRecord(start, 0)); for (int j = 0; j < BATCH_SIZE; j++) { hipLaunchKernelGGL((EmptyKernel), dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream0); } - hipEventRecord(stop, 0); - hipEventSynchronize(stop); - hipEventElapsedTime(&results[i], start, stop); + checkHipErrors(hipEventRecord(stop, 0)); + checkHipErrors(hipEventSynchronize(stop)); + checkHipErrors(hipEventElapsedTime(&results[i], start, stop)); } print_timing("Batch dispatch latency", results, BATCH_SIZE); - hipEventDestroy(start); - hipEventDestroy(stop); - hipCtxDestroy(context); + checkHipErrors(hipEventDestroy(start)); + checkHipErrors(hipEventDestroy(stop)); + checkHipErrors(hipCtxDestroy(context)); } diff --git a/projects/hip-tests/samples/1_Utils/hipInfo/CMakeLists.txt b/projects/hip-tests/samples/1_Utils/hipInfo/CMakeLists.txt index 6192c8ecff..60cf123282 100644 --- a/projects/hip-tests/samples/1_Utils/hipInfo/CMakeLists.txt +++ b/projects/hip-tests/samples/1_Utils/hipInfo/CMakeLists.txt @@ -57,6 +57,8 @@ add_executable(hipInfo hipInfo.cpp) # Link with HIP target_link_libraries(hipInfo hip::host) +target_include_directories(hipInfo PRIVATE ../../common) + # Used only when make install is called # when hipInfo is built as part of compute project # hipInfo.exe will be installed to install/hip/bin path diff --git a/projects/hip-tests/samples/1_Utils/hipInfo/Makefile b/projects/hip-tests/samples/1_Utils/hipInfo/Makefile index c6c343dbd1..c5bc1a9c74 100644 --- a/projects/hip-tests/samples/1_Utils/hipInfo/Makefile +++ b/projects/hip-tests/samples/1_Utils/hipInfo/Makefile @@ -26,13 +26,14 @@ ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc +INCLUDES := -I../../common EXE=hipInfo all: install $(EXE): hipInfo.cpp - $(HIPCC) hipInfo.cpp -o $@ + $(HIPCC) hipInfo.cpp $(INCLUDES) -o $@ install: $(EXE) cp $(EXE) $(HIP_PATH)/bin diff --git a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp index 28128b6d33..56978818d9 100644 --- a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp @@ -23,6 +23,7 @@ THE SOFTWARE. #include #include #include "hip/hip_runtime.h" +#include "hip_helper.h" #define KNRM "\x1B[0m" #define KRED "\x1B[31m" @@ -33,20 +34,6 @@ THE SOFTWARE. #define KCYN "\x1B[36m" #define KWHT "\x1B[37m" -#define failed(...) \ - printf("%serror: ", KRED); \ - printf(__VA_ARGS__); \ - printf("\n"); \ - printf("error: TEST FAILED\n%s", KNRM); \ - exit(EXIT_FAILURE); - -#define HIPCHECK(error) \ - if (error != hipSuccess) { \ - printf("%serror: '%s'(%d) at %s:%d%s\n", KRED, hipGetErrorString(error), error, __FILE__, \ - __LINE__, KNRM); \ - failed("API returned error code."); \ - } - void printCompilerInfo() { #ifdef __NVCC__ printf("compiler: nvcc\n"); @@ -76,7 +63,7 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "device#" << deviceId << endl; hipDeviceProp_t props = {0}; - HIPCHECK(hipGetDeviceProperties(&props, deviceId)); + checkHipErrors(hipGetDeviceProperties(&props, deviceId)); cout << setw(w1) << "Name: " << props.name << endl; cout << setw(w1) << "pciBusID: " << props.pciBusID << endl; @@ -149,11 +136,11 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "gcnArchName: " << props.gcnArchName << endl; #endif int deviceCnt; - hipGetDeviceCount(&deviceCnt); + checkHipErrors(hipGetDeviceCount(&deviceCnt)); cout << setw(w1) << "peers: "; for (int i = 0; i < deviceCnt; i++) { int isPeer; - hipDeviceCanAccessPeer(&isPeer, i, deviceId); + checkHipErrors(hipDeviceCanAccessPeer(&isPeer, i, deviceId)); if (isPeer) { cout << "device#" << i << " "; } @@ -162,7 +149,7 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "non-peers: "; for (int i = 0; i < deviceCnt; i++) { int isPeer; - hipDeviceCanAccessPeer(&isPeer, i, deviceId); + checkHipErrors(hipDeviceCanAccessPeer(&isPeer, i, deviceId)); if (!isPeer) { cout << "device#" << i << " "; } @@ -185,7 +172,7 @@ void printDeviceProp(int deviceId) { size_t free, total; - hipMemGetInfo(&free, &total); + checkHipErrors(hipMemGetInfo(&free, &total)); cout << fixed << setprecision(2); cout << setw(w1) << "memInfo.total: " << bytesToGB(total) << " GB" << endl; @@ -202,10 +189,10 @@ int main(int argc, char* argv[]) { int deviceCnt; - HIPCHECK(hipGetDeviceCount(&deviceCnt)); + checkHipErrors(hipGetDeviceCount(&deviceCnt)); for (int i = 0; i < deviceCnt; i++) { - hipSetDevice(i); + checkHipErrors(hipSetDevice(i)); printDeviceProp(i); } diff --git a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/CMakeLists.txt index ae4735c1ae..015e6a047e 100644 --- a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/CMakeLists.txt @@ -40,5 +40,7 @@ set(CMAKE_BUILD_TYPE Release) # Create the excutable add_executable(MatrixTranspose MatrixTranspose.cpp) +target_include_directories(MatrixTranspose PRIVATE ../../common) + # Link with HIP target_link_libraries(MatrixTranspose hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Makefile b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Makefile index 6d5b787510..308efa6dc6 100644 --- a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Makefile @@ -30,6 +30,7 @@ HIPCC=$(HIP_PATH)/bin/hipcc TARGET=hcc +INCLUDES := -I../../common SOURCES = MatrixTranspose.cpp OBJECTS = $(SOURCES:.cpp=.o) @@ -40,7 +41,7 @@ EXECUTABLE=./MatrixTranspose all: $(EXECUTABLE) test -CXXFLAGS =-g +CXXFLAGS =-g $(INCLUDES) CXX=$(HIPCC) diff --git a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp index 8444cff851..205f3b6966 100644 --- a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp +++ b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. // hip header file #include "hip/hip_runtime.h" +#include "hip_helper.h" #define WIDTH 1024 @@ -61,7 +62,7 @@ int main() { float* gpuTransposeMatrix; hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + checkHipErrors(hipGetDeviceProperties(&devProp, 0)); std::cout << "Device name " << devProp.name << std::endl; @@ -78,11 +79,11 @@ int main() { } // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + checkHipErrors(hipMalloc((void**)&gpuMatrix, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float))); // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + checkHipErrors(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); // Lauching kernel from host hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), @@ -90,7 +91,7 @@ int main() { gpuMatrix, WIDTH); // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + checkHipErrors(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost)); // CPU MatrixTranspose computation matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); @@ -110,8 +111,8 @@ int main() { } // free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + checkHipErrors(hipFree(gpuMatrix)); + checkHipErrors(hipFree(gpuTransposeMatrix)); // free the resources on host side free(Matrix); diff --git a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/CMakeLists.txt index ac4e586772..f20827df3f 100644 --- a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/CMakeLists.txt @@ -40,5 +40,7 @@ set(CMAKE_BUILD_TYPE Release) # Create the excutable add_executable(inline_asm inline_asm.cpp) +target_include_directories(inline_asm PRIVATE ../../common) + # Link with HIP target_link_libraries(inline_asm hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Makefile b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Makefile index 58f013a2ba..31928eb2dc 100644 --- a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Makefile @@ -32,7 +32,7 @@ TARGET=hcc SOURCES = inline_asm.cpp OBJECTS = $(SOURCES:.cpp=.o) - +INCLUDES := -I../../common EXECUTABLE=./inline_asm .PHONY: test @@ -40,7 +40,7 @@ EXECUTABLE=./inline_asm all: $(EXECUTABLE) test -CXXFLAGS =-g +CXXFLAGS =-g $(INCLUDES) CXX=$(HIPCC) diff --git a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp index 8145b3c86e..39bc86f1b6 100644 --- a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp +++ b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. // hip header file #include "hip/hip_runtime.h" +#include "hip_helper.h" #define WIDTH 1024 @@ -59,13 +60,13 @@ int main() { float* gpuTransposeMatrix; hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + checkHipErrors(hipGetDeviceProperties(&devProp, 0)); std::cout << "Device name " << devProp.name << std::endl; hipEvent_t start, stop; - hipEventCreate(&start); - hipEventCreate(&stop); + checkHipErrors(hipEventCreate(&start)); + checkHipErrors(hipEventCreate(&stop)); float eventMs = 1.0f; int i; @@ -81,25 +82,25 @@ int main() { } // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + checkHipErrors(hipMalloc((void**)&gpuMatrix, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float))); // Record the start event - hipEventRecord(start, NULL); + checkHipErrors(hipEventRecord(start, NULL)); // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + checkHipErrors(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); // Record the stop event - hipEventRecord(stop, NULL); - hipEventSynchronize(stop); + checkHipErrors(hipEventRecord(stop, NULL)); + checkHipErrors(hipEventSynchronize(stop)); - hipEventElapsedTime(&eventMs, start, stop); + checkHipErrors(hipEventElapsedTime(&eventMs, start, stop)); printf("hipMemcpyHostToDevice time taken = %6.3fms\n", eventMs); // Record the start event - hipEventRecord(start, NULL); + checkHipErrors(hipEventRecord(start, NULL)); // Lauching kernel from host hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), @@ -107,24 +108,24 @@ int main() { gpuMatrix, WIDTH); // Record the stop event - hipEventRecord(stop, NULL); - hipEventSynchronize(stop); + checkHipErrors(hipEventRecord(stop, NULL)); + checkHipErrors(hipEventSynchronize(stop)); - hipEventElapsedTime(&eventMs, start, stop); + checkHipErrors(hipEventElapsedTime(&eventMs, start, stop)); printf("kernel Execution time = %6.3fms\n", eventMs); // Record the start event - hipEventRecord(start, NULL); + checkHipErrors(hipEventRecord(start, NULL)); // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + checkHipErrors(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost)); // Record the stop event - hipEventRecord(stop, NULL); - hipEventSynchronize(stop); + checkHipErrors(hipEventRecord(stop, NULL)); + checkHipErrors(hipEventSynchronize(stop)); - hipEventElapsedTime(&eventMs, start, stop); + checkHipErrors(hipEventElapsedTime(&eventMs, start, stop)); printf("hipMemcpyDeviceToHost time taken = %6.3fms\n", eventMs); @@ -147,8 +148,8 @@ int main() { } // free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + checkHipErrors(hipFree(gpuMatrix)); + checkHipErrors(hipFree(gpuTransposeMatrix)); // free the resources on host side free(Matrix); diff --git a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/CMakeLists.txt index f93b2e791e..55aa40b2a6 100644 --- a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/CMakeLists.txt @@ -50,5 +50,7 @@ add_custom_target( add_dependencies(texture2dDrv codeobj) +target_include_directories(texture2dDrv PRIVATE ../../common) + # Link with HIP target_link_libraries(texture2dDrv hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/Makefile b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/Makefile index f149005aa5..25737ea2e2 100644 --- a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/Makefile @@ -27,11 +27,12 @@ ifeq (,$(HIP_PATH)) endif HIPCC=$(HIP_PATH)/bin/hipcc HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) +INCLUDES := -I../../common all: tex2dKernel.code texture2dDrv.out texture2dDrv.out: texture2dDrv.cpp - $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + $(HIPCC) $(HIPCC_FLAGS) $(INCLUDES) $< -o $@ tex2dKernel.code: tex2dKernel.cpp $(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@ diff --git a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index f03b3873bf..e30c1ba911 100644 --- a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -24,21 +24,12 @@ THE SOFTWARE. #include #include #include +#include "hip_helper.h" #define fileName "tex2dKernel.code" bool testResult = true; -#define HIP_CHECK(cmd) \ - { \ - hipError_t status = cmd; \ - if (status != hipSuccess) { \ - std::cout << "error: #" << status << " (" << hipGetErrorString(status) \ - << ") at line:" << __LINE__ << ": " << #cmd << std::endl; \ - abort(); \ - } \ - } - template::value>::type *t = nullptr> static inline hipArray_Format getArrayFormat() { @@ -154,11 +145,11 @@ bool runTest(hipModule_t &module, const char *refName, const char *funcName) { hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); hipArray_t array; - HIP_CHECK(hipMallocArray(&array, &channelDesc, width, height)); + checkHipErrors(hipMallocArray(&array, &channelDesc, width, height)); const size_t spitch = width * sizeof(T); - HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, hData, spitch, width * sizeof(T), + checkHipErrors(hipMemcpy2DToArray(array, 0, 0, hData, spitch, width * sizeof(T), height, hipMemcpyHostToDevice)); hipResourceDesc resDesc; @@ -175,10 +166,10 @@ bool runTest(hipModule_t &module, const char *refName, const char *funcName) { texDesc.normalizedCoords = 0; hipTextureObject_t texObj; - HIP_CHECK(hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr)); + checkHipErrors(hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr)); T *dData = NULL; - HIP_CHECK(hipMalloc((void** )&dData, size)); + checkHipErrors(hipMalloc((void** )&dData, size)); struct { void *_Ad; @@ -197,18 +188,18 @@ bool runTest(hipModule_t &module, const char *refName, const char *funcName) { HIP_LAUNCH_PARAM_BUFFER_SIZE, &sizeTemp, HIP_LAUNCH_PARAM_END }; hipFunction_t Function; - HIP_CHECK(hipModuleGetFunction(&Function, module, funcName)); + checkHipErrors(hipModuleGetFunction(&Function, module, funcName)); int temp1 = width / 16; int temp2 = height / 16; - HIP_CHECK( + checkHipErrors( hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void** )&config)); - HIP_CHECK(hipDeviceSynchronize()); + checkHipErrors(hipDeviceSynchronize()); T *hOutputData = (T*) malloc(size); memset(hOutputData, 0, size); - HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + checkHipErrors(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); for (int i = 0; i < height; i++) { for (int j = 0; j < width; j++) { @@ -219,9 +210,9 @@ bool runTest(hipModule_t &module, const char *refName, const char *funcName) { } } } - HIP_CHECK(hipDestroyTextureObject(texObj)); - HIP_CHECK(hipFree(dData)); - HIP_CHECK(hipFreeArray(array)); + checkHipErrors(hipDestroyTextureObject(texObj)); + checkHipErrors(hipFree(dData)); + checkHipErrors(hipFreeArray(array)); free(hOutputData); free(hData); printf("%s test %s ...\n", funcName, testResult ? "PASSED" : "FAILED"); @@ -231,7 +222,7 @@ bool runTest(hipModule_t &module, const char *refName, const char *funcName) { inline bool isImageSupported() { int imageSupport = 1; #ifdef __HIP_PLATFORM_AMD__ - HIP_CHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, + checkHipErrors(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, 0)); #endif return imageSupport != 0; @@ -242,10 +233,10 @@ int main(int argc, char** argv) { printf("Texture is not support on the device. Skipped.\n"); return 0; } - HIP_CHECK(hipInit(0)); - HIP_CHECK(hipSetDevice(0)); + checkHipErrors(hipInit(0)); + checkHipErrors(hipSetDevice(0)); hipModule_t module; - HIP_CHECK(hipModuleLoad(&module, fileName)); + checkHipErrors(hipModuleLoad(&module, fileName)); testResult = testResult && runTest(module, "texChar", "tex2dKernelChar"); testResult = testResult && runTest(module, "texShort", "tex2dKernelShort"); testResult = testResult && runTest(module, "texInt", "tex2dKernelInt"); @@ -255,7 +246,7 @@ int main(int argc, char** argv) { testResult = testResult && runTest(module, "texInt4", "tex2dKernelInt4"); testResult = testResult && runTest(module, "texFloat4", "tex2dKernelFloat4"); - HIP_CHECK(hipModuleUnload(module)); + checkHipErrors(hipModuleUnload(module)); printf("texture2dDrv %s ...\n", testResult ? "PASSED" : "FAILED"); return testResult ? EXIT_SUCCESS : EXIT_FAILURE; } diff --git a/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt index 7045a891e1..9bc99e36ae 100644 --- a/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt @@ -51,6 +51,8 @@ set(MY_NVCC_OPTIONS) set_source_files_properties(${MY_SOURCE_FILES} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} CLANG_OPTIONS ${MY_CLANG_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS}) +target_include_directories(${MY_TARGET_NAME} PRIVATE ../../common) + # Search for rocm in common locations list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH}) find_package(hip QUIET CONFIG) diff --git a/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/MatrixTranspose.cpp b/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/MatrixTranspose.cpp index 8444cff851..205f3b6966 100644 --- a/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/MatrixTranspose.cpp +++ b/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/MatrixTranspose.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. // hip header file #include "hip/hip_runtime.h" +#include "hip_helper.h" #define WIDTH 1024 @@ -61,7 +62,7 @@ int main() { float* gpuTransposeMatrix; hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + checkHipErrors(hipGetDeviceProperties(&devProp, 0)); std::cout << "Device name " << devProp.name << std::endl; @@ -78,11 +79,11 @@ int main() { } // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + checkHipErrors(hipMalloc((void**)&gpuMatrix, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float))); // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + checkHipErrors(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); // Lauching kernel from host hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), @@ -90,7 +91,7 @@ int main() { gpuMatrix, WIDTH); // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + checkHipErrors(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost)); // CPU MatrixTranspose computation matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); @@ -110,8 +111,8 @@ int main() { } // free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + checkHipErrors(hipFree(gpuMatrix)); + checkHipErrors(hipFree(gpuTransposeMatrix)); // free the resources on host side free(Matrix); diff --git a/projects/hip-tests/samples/2_Cookbook/13_occupancy/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/13_occupancy/CMakeLists.txt index 481cf58b26..44d0303068 100644 --- a/projects/hip-tests/samples/2_Cookbook/13_occupancy/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/13_occupancy/CMakeLists.txt @@ -40,5 +40,7 @@ set(CMAKE_BUILD_TYPE Release) # Create the excutable add_executable(occupancy occupancy.cpp) +target_include_directories(occupancy PRIVATE ../../common) + # Link with HIP target_link_libraries(occupancy hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/13_occupancy/Makefile b/projects/hip-tests/samples/2_Cookbook/13_occupancy/Makefile index 73f0753afb..dd4037418f 100644 --- a/projects/hip-tests/samples/2_Cookbook/13_occupancy/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/13_occupancy/Makefile @@ -26,7 +26,7 @@ ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc - +INCLUDES := -I../../common EXE=./occupancy .PHONY: test @@ -34,7 +34,7 @@ EXE=./occupancy all: test $(EXE): occupancy.cpp - $(HIPCC) $^ -o $@ + $(HIPCC) $(INCLUDES) $^ -o $@ test: $(EXE) $(EXE) diff --git a/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp b/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp index 4f51b61c6e..497b670f04 100644 --- a/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp +++ b/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp @@ -19,14 +19,9 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include +#include "hip_helper.h" #define NUM 1000000 -#define HIP_CHECK(status) \ - if (status != hipSuccess) { \ - std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \ - exit(0); \ - } - // Device (Kernel) function __global__ void multiply(float* C, float* A, float* B, int N){ @@ -47,11 +42,11 @@ void multiplyCPU(float* C, float* A, float* B, int N){ void launchKernel(float* C, float* A, float* B, bool manual){ hipDeviceProp_t devProp; - HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + checkHipErrors(hipGetDeviceProperties(&devProp, 0)); hipEvent_t start, stop; - HIP_CHECK(hipEventCreate(&start)); - HIP_CHECK(hipEventCreate(&stop)); + checkHipErrors(hipEventCreate(&start)); + checkHipErrors(hipEventCreate(&stop)); float eventMs = 1.0f; const unsigned threadsperblock = 32; const unsigned blocks = (NUM/threadsperblock)+1; @@ -66,28 +61,28 @@ void launchKernel(float* C, float* A, float* B, bool manual){ std::cout << std::endl << "Manual Configuration with block size " << blockSize << std::endl; } else{ - HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&mingridSize, &blockSize, multiply, 0, 0)); + checkHipErrors(hipOccupancyMaxPotentialBlockSize(&mingridSize, &blockSize, multiply, 0, 0)); std::cout << std::endl << "Automatic Configuation based on hipOccupancyMaxPotentialBlockSize " << std::endl; std::cout << "Suggested blocksize is " << blockSize << ", Minimum gridsize is " << mingridSize << std::endl; gridSize = (NUM/blockSize)+1; } // Record the start event - HIP_CHECK(hipEventRecord(start, NULL)); + checkHipErrors(hipEventRecord(start, NULL)); // Launching the Kernel from Host hipLaunchKernelGGL(multiply, dim3(gridSize), dim3(blockSize), 0, 0, C, A, B, NUM); // Record the stop event - HIP_CHECK(hipEventRecord(stop, NULL)); - HIP_CHECK(hipEventSynchronize(stop)); + checkHipErrors(hipEventRecord(stop, NULL)); + checkHipErrors(hipEventSynchronize(stop)); - HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop)); + checkHipErrors(hipEventElapsedTime(&eventMs, start, stop)); printf("kernel Execution time = %6.3fms\n", eventMs); //Calculate Occupancy int numBlock = 0; - HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, multiply, blockSize, 0)); + checkHipErrors(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, multiply, blockSize, 0)); if(devProp.maxThreadsPerMultiProcessor){ std::cout << "Theoretical Occupancy is " << (double)numBlock* blockSize/devProp.maxThreadsPerMultiProcessor * 100 << "%" << std::endl; @@ -113,14 +108,14 @@ int main() { } // allocate the memory on the device side - HIP_CHECK(hipMalloc((void**)&Ad, NUM * sizeof(float))); - HIP_CHECK(hipMalloc((void**)&Bd, NUM * sizeof(float))); - HIP_CHECK(hipMalloc((void**)&C0d, NUM * sizeof(float))); - HIP_CHECK(hipMalloc((void**)&C1d, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&Ad, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&Bd, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&C0d, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&C1d, NUM * sizeof(float))); // Memory transfer from host to device - HIP_CHECK(hipMemcpy(Ad,A,NUM * sizeof(float), hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy(Bd,B,NUM * sizeof(float), hipMemcpyHostToDevice)); + checkHipErrors(hipMemcpy(Ad,A,NUM * sizeof(float), hipMemcpyHostToDevice)); + checkHipErrors(hipMemcpy(Bd,B,NUM * sizeof(float), hipMemcpyHostToDevice)); //Kernel launch with manual/default block size launchKernel(C0d, Ad, Bd, 1); @@ -129,8 +124,8 @@ int main() { launchKernel(C1d, Ad, Bd, 0); // Memory transfer from device to host - HIP_CHECK(hipMemcpy(C0,C0d, NUM * sizeof(float), hipMemcpyDeviceToHost)); - HIP_CHECK(hipMemcpy(C1,C1d, NUM * sizeof(float), hipMemcpyDeviceToHost)); + checkHipErrors(hipMemcpy(C0,C0d, NUM * sizeof(float), hipMemcpyDeviceToHost)); + checkHipErrors(hipMemcpy(C1,C1d, NUM * sizeof(float), hipMemcpyDeviceToHost)); // CPU computation multiplyCPU(cpuC, A, B, NUM); @@ -163,10 +158,10 @@ int main() { printf("\nAutomatic Test PASSED!\n"); } - HIP_CHECK(hipFree(Ad)); - HIP_CHECK(hipFree(Bd)); - HIP_CHECK(hipFree(C0d)); - HIP_CHECK(hipFree(C1d)); + checkHipErrors(hipFree(Ad)); + checkHipErrors(hipFree(Bd)); + checkHipErrors(hipFree(C0d)); + checkHipErrors(hipFree(C1d)); free(A); free(B); diff --git a/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/CMakeLists.txt index 084f273096..91354b11cc 100644 --- a/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/CMakeLists.txt @@ -40,5 +40,7 @@ set(CMAKE_BUILD_TYPE Release) # Create the excutable add_executable(gpuarch gpuarch.cpp) +target_include_directories(gpuarch PRIVATE ../../common) + # Link with HIP target_link_libraries(gpuarch hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/Makefile b/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/Makefile index c730c10a06..a3b1115780 100644 --- a/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/Makefile @@ -26,7 +26,7 @@ ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc - +INCLUDES := -I../../common EXE=./gpuarch .PHONY: test @@ -34,7 +34,7 @@ EXE=./gpuarch all: test $(EXE): gpuarch.cpp - $(HIPCC) $^ -o $@ + $(HIPCC) $(INCLUDES) $^ -o $@ test: $(EXE) $(EXE) diff --git a/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp b/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp index f1b521fcd1..35cd93a2e9 100644 --- a/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp +++ b/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp @@ -25,12 +25,6 @@ THE SOFTWARE. #define SIZE (BLOCKS_PER_GRID * THREADS_PER_BLOCK) #define NOT_SUPPORTED -99 // dummy number indicates unsupported operation -#define HIP_STATUS_CHECK(status) \ - if (status != hipSuccess) { \ - std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \ - exit(0); \ - } - // Using __gfx*__ macro one can have GPU architecture specific code flow // For example: If below kernel runs on gfx908 it will increment 'in' by 'value' and store into // 'out' @@ -57,8 +51,8 @@ int main() { int32_t* hInput = static_cast(malloc(NBytes)); int32_t* hOutput = static_cast(malloc(NBytes)); - HIP_STATUS_CHECK(hipMalloc(&dInput, NBytes)); - HIP_STATUS_CHECK(hipMalloc(&dOutput, NBytes)); + checkHipErrors(hipMalloc(&dInput, NBytes)); + checkHipErrors(hipMalloc(&dOutput, NBytes)); // Initialize host input/output buffers for (int i = 0; i < SIZE; ++i) { @@ -67,14 +61,14 @@ int main() { } // Initialize device input buffer - HIP_STATUS_CHECK(hipMemcpy(dInput, hInput, NBytes, hipMemcpyHostToDevice)); + checkHipErrors(hipMemcpy(dInput, hInput, NBytes, hipMemcpyHostToDevice)); // Launch kernel hipLaunchKernelGGL(incrementKernel, dim3(BLOCKS_PER_GRID), dim3(THREADS_PER_BLOCK), 0, 0, dInput, dOutput, incrementValue, SIZE); // Copy result back to host buffer - HIP_STATUS_CHECK(hipMemcpy(hOutput, dOutput, NBytes, hipMemcpyDeviceToHost)); + checkHipErrors(hipMemcpy(hOutput, dOutput, NBytes, hipMemcpyDeviceToHost)); bool flag = true; // verify data diff --git a/projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/Makefile b/projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/Makefile index 56917be6c3..e0615b84dd 100644 --- a/projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/Makefile @@ -30,6 +30,7 @@ HIPCC=$(HIP_PATH)/bin/hipcc CLANG=$(HIP_PATH)/llvm/bin/clang LLVM_MC=$(HIP_PATH)/llvm/bin/llvm-mc CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/llvm/bin/clang-offload-bundler +INCLUDES := -I../../common SRCS=square.cpp @@ -57,8 +58,8 @@ GPU_ARCH9=gfx1103 all: src_to_asm asm_to_exec src_to_asm: - $(HIPCC) -c -S --cuda-host-only -target x86_64-linux-gnu -o $(SQ_HOST_ASM) $(SRCS) - $(HIPCC) -c -S --cuda-device-only --offload-arch=$(GPU_ARCH1) --offload-arch=$(GPU_ARCH2) --offload-arch=$(GPU_ARCH3) --offload-arch=$(GPU_ARCH4) --offload-arch=$(GPU_ARCH5) --offload-arch=$(GPU_ARCH6) --offload-arch=$(GPU_ARCH7) --offload-arch=$(GPU_ARCH8) --offload-arch=$(GPU_ARCH9) $(SRCS) + $(HIPCC) -c -S $(INCLUDES) --cuda-host-only -target x86_64-linux-gnu -o $(SQ_HOST_ASM) $(SRCS) + $(HIPCC) -c -S $(INCLUDES) --cuda-device-only --offload-arch=$(GPU_ARCH1) --offload-arch=$(GPU_ARCH2) --offload-arch=$(GPU_ARCH3) --offload-arch=$(GPU_ARCH4) --offload-arch=$(GPU_ARCH5) --offload-arch=$(GPU_ARCH6) --offload-arch=$(GPU_ARCH7) --offload-arch=$(GPU_ARCH8) --offload-arch=$(GPU_ARCH9) $(SRCS) # You may modify the .s assembly files before the next step # By default, their names will be: diff --git a/projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/square.cpp b/projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/square.cpp index a04bf625ca..296be3cbbe 100644 --- a/projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/square.cpp +++ b/projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/square.cpp @@ -19,15 +19,7 @@ THE SOFTWARE. #include #include - -#define CHECK(cmd) \ -{\ - hipError_t error = cmd;\ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} +#include "hip_helper.h" /* This kernel is a placeholder for the kernel in assembly generated by this * sample. It will be replaced by the kernel in assembly. @@ -55,14 +47,14 @@ int main(int argc, char *argv[]) size_t Nbytes = N * sizeof(float); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); + checkHipErrors(hipGetDeviceProperties(&props, 0/*deviceID*/)); printf ("info: running on device %s\n", props.name); printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); A_h = (float*)malloc(Nbytes); - CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); + checkHipErrors(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); C_h = (float*)malloc(Nbytes); - CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); + checkHipErrors(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); // Fill with Phi + i for (size_t i=0; i>> (C_d, A_d, N); printf ("info: copy Device2Host\n"); - CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + checkHipErrors ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); - printf ("info: check result\n"); + printf ("info: checkHipErrors result\n"); for (size_t i=0; i #include - -#define CHECK(cmd) \ -{\ - hipError_t error = cmd;\ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} +#include "hip_helper.h" /* This kernel is a placeholder for the kernel in LLVM IR generated by this * sample. It will be replaced by the kernel in LLVM IR. @@ -55,14 +47,14 @@ int main(int argc, char *argv[]) size_t Nbytes = N * sizeof(float); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); + checkHipErrors(hipGetDeviceProperties(&props, 0/*deviceID*/)); printf ("info: running on device %s\n", props.name); printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); A_h = (float*)malloc(Nbytes); - CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); + checkHipErrors(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); C_h = (float*)malloc(Nbytes); - CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); + checkHipErrors(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); // Fill with Phi + i for (size_t i=0; i>> (C_d, A_d, N); printf ("info: copy Device2Host\n"); - CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + checkHipErrors ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); - printf ("info: check result\n"); + printf ("info: checkHipErrors result\n"); for (size_t i=0; i #include - -#define CHECK(cmd) \ -{\ - hipError_t error = cmd;\ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} - +#include "hip_helper.h" /* * Square each element in the array A and write to array C. @@ -57,14 +48,14 @@ int main(int argc, char *argv[]) size_t Nbytes = N * sizeof(float); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); + checkHipErrors(hipGetDeviceProperties(&props, 0/*deviceID*/)); printf ("info: running on device %s\n", props.name); printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); A_h = (float*)malloc(Nbytes); - CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess ); + checkHipErrors(A_h == 0 ? hipErrorOutOfMemory : hipSuccess ); C_h = (float*)malloc(Nbytes); - CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess ); + checkHipErrors(C_h == 0 ? hipErrorOutOfMemory : hipSuccess ); // Fill with Phi + i for (size_t i=0; i #include - -#define CHECK(cmd) \ -{\ - hipError_t error = cmd;\ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} - +#include "hip_helper.h" /* * Square each element in the array A and write to array C. @@ -57,14 +48,14 @@ int main(int argc, char *argv[]) size_t Nbytes = N * sizeof(float); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); + checkHipErrors(hipGetDeviceProperties(&props, 0/*deviceID*/)); printf ("info: running on device %s\n", props.name); printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); A_h = (float*)malloc(Nbytes); - CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess ); + checkHipErrors(A_h == 0 ? hipErrorOutOfMemory : hipSuccess ); C_h = (float*)malloc(Nbytes); - CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess ); + checkHipErrors(C_h == 0 ? hipErrorOutOfMemory : hipSuccess ); // Fill with Phi + i for (size_t i=0; i #include - -#define CHECK(cmd) \ -{\ - hipError_t error = cmd;\ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} - +#include "hip_helper.h" /* * Square each element in the array A and write to array C. @@ -57,14 +48,14 @@ int main(int argc, char *argv[]) size_t Nbytes = N * sizeof(float); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); + checkHipErrors(hipGetDeviceProperties(&props, 0/*deviceID*/)); printf ("info: running on device %s\n", props.name); printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); A_h = (float*)malloc(Nbytes); - CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess ); + checkHipErrors(A_h == 0 ? hipErrorOutOfMemory : hipSuccess ); C_h = (float*)malloc(Nbytes); - CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess ); + checkHipErrors(C_h == 0 ? hipErrorOutOfMemory : hipSuccess ); // Fill with Phi + i for (size_t i=0; i #include +#include #include #include @@ -69,7 +70,7 @@ int main() hipDeviceProp_t props; int device = 0; - hipGetDeviceProperties(&props, device); + checkHipErrors(hipGetDeviceProperties(&props, device)); const char* options[] = {}; @@ -100,8 +101,8 @@ int main() hipModule_t module; hipFunction_t kernel; - hipModuleLoadData(&module, code.data()); - hipModuleGetFunction(&kernel, module, "saxpy"); + checkHipErrors(hipModuleLoadData(&module, code.data())); + checkHipErrors(hipModuleGetFunction(&kernel, module, "saxpy")); size_t n = NUM_THREADS * NUM_BLOCKS; size_t bufferSize = n * sizeof(float); @@ -117,11 +118,11 @@ int main() } hipDeviceptr_t dX, dY, dOut; - hipMalloc((void **)&dX, bufferSize); - hipMalloc((void **)&dY, bufferSize); - hipMalloc((void **)&dOut, bufferSize); - hipMemcpyHtoD(dX, hX.get(), bufferSize); - hipMemcpyHtoD(dY, hY.get(), bufferSize); + checkHipErrors(hipMalloc((void **)&dX, bufferSize)); + checkHipErrors(hipMalloc((void **)&dY, bufferSize)); + checkHipErrors(hipMalloc((void **)&dOut, bufferSize)); + checkHipErrors(hipMemcpyHtoD(dX, hX.get(), bufferSize)); + checkHipErrors(hipMemcpyHtoD(dY, hY.get(), bufferSize)); struct { float a_; @@ -136,9 +137,9 @@ int main() HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; - hipModuleLaunchKernel(kernel, NUM_BLOCKS, 1, 1, NUM_THREADS, 1, 1, - 0, nullptr, nullptr, config); - hipMemcpyDtoH(hOut.get(), dOut, bufferSize); + checkHipErrors(hipModuleLaunchKernel(kernel, NUM_BLOCKS, 1, 1, NUM_THREADS, 1, 1, + 0, nullptr, nullptr, config)); + checkHipErrors(hipMemcpyDtoH(hOut.get(), dOut, bufferSize)); for (size_t i = 0; i < n; ++i) { if (fabs(a * hX[i] + hY[i] - hOut[i]) > fabs(hOut[i])* 1e-6) { @@ -146,11 +147,11 @@ int main() } } - hipFree((void *)dX); - hipFree((void *)dY); - hipFree((void *)dOut); + checkHipErrors(hipFree((void *)dX)); + checkHipErrors(hipFree((void *)dY)); + checkHipErrors(hipFree((void *)dOut)); - hipModuleUnload(module); + checkHipErrors(hipModuleUnload(module)); cout << "SAXPY test completed" << endl; } diff --git a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/CMakeLists.txt index f2f5196373..bab322511d 100644 --- a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/CMakeLists.txt @@ -40,5 +40,7 @@ set(CMAKE_BUILD_TYPE Release) # Create the excutable add_executable(sharedMemory sharedMemory.cpp) +target_include_directories(sharedMemory PRIVATE ../../common) + # Link with HIP target_link_libraries(sharedMemory hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Makefile b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Makefile index bbd7daace3..c8571cf21f 100644 --- a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Makefile @@ -32,6 +32,7 @@ TARGET=hcc SOURCES = sharedMemory.cpp OBJECTS = $(SOURCES:.cpp=.o) +INCLUDES := -I../../common EXECUTABLE=./sharedMemory @@ -40,7 +41,7 @@ EXECUTABLE=./sharedMemory all: $(EXECUTABLE) test -CXXFLAGS =-g +CXXFLAGS =-g $(INCLUDES) CXX=$(HIPCC) diff --git a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp index 8bd489dbf3..b648288a8f 100644 --- a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp +++ b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp @@ -24,7 +24,7 @@ THE SOFTWARE. // hip header file #include "hip/hip_runtime.h" - +#include "hip_helper.h" #define WIDTH 64 @@ -66,7 +66,7 @@ int main() { float* gpuTransposeMatrix; hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + checkHipErrors(hipGetDeviceProperties(&devProp, 0)); std::cout << "Device name " << devProp.name << std::endl; @@ -83,11 +83,11 @@ int main() { } // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + checkHipErrors(hipMalloc((void**)&gpuMatrix, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float))); // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + checkHipErrors(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); // Lauching kernel from host hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), @@ -95,7 +95,7 @@ int main() { gpuMatrix, WIDTH); // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + checkHipErrors(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost)); // CPU MatrixTranspose computation matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); @@ -116,8 +116,8 @@ int main() { } // free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + checkHipErrors(hipFree(gpuMatrix)); + checkHipErrors(hipFree(gpuTransposeMatrix)); // free the resources on host side free(Matrix); diff --git a/projects/hip-tests/samples/2_Cookbook/4_shfl/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/4_shfl/CMakeLists.txt index 394d5852fb..d1e8a4843f 100644 --- a/projects/hip-tests/samples/2_Cookbook/4_shfl/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/4_shfl/CMakeLists.txt @@ -40,5 +40,7 @@ set(CMAKE_BUILD_TYPE Release) # Create the excutable add_executable(shfl shfl.cpp) +target_include_directories(shfl PRIVATE ../../common) + # Link with HIP target_link_libraries(shfl hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile b/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile index de94a3e546..6305ad10e0 100644 --- a/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile @@ -36,6 +36,7 @@ TARGET=hcc SOURCES = shfl.cpp OBJECTS = $(SOURCES:.cpp=.o) +INCLUDES := -I../../common EXECUTABLE=./shfl @@ -44,7 +45,7 @@ EXECUTABLE=./shfl all: $(EXECUTABLE) test -CXXFLAGS =-g +CXXFLAGS =-g $(INCLUDES) CXX=$(HIPCC) diff --git a/projects/hip-tests/samples/2_Cookbook/4_shfl/shfl.cpp b/projects/hip-tests/samples/2_Cookbook/4_shfl/shfl.cpp index de1ff7a950..6ef968bab3 100644 --- a/projects/hip-tests/samples/2_Cookbook/4_shfl/shfl.cpp +++ b/projects/hip-tests/samples/2_Cookbook/4_shfl/shfl.cpp @@ -24,7 +24,7 @@ THE SOFTWARE. // hip header file #include "hip/hip_runtime.h" - +#include "hip_helper.h" #define WIDTH 4 @@ -63,7 +63,7 @@ int main() { float* gpuTransposeMatrix; hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + checkHipErrors(hipGetDeviceProperties(&devProp, 0)); std::cout << "Device name " << devProp.name << std::endl; @@ -80,18 +80,18 @@ int main() { } // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + checkHipErrors(hipMalloc((void**)&gpuMatrix, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float))); // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + checkHipErrors(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); // Lauching kernel from host hipLaunchKernelGGL(matrixTranspose, dim3(1), dim3(THREADS_PER_BLOCK_X * THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, gpuMatrix, WIDTH); // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + checkHipErrors(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost)); // CPU MatrixTranspose computation matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); @@ -112,8 +112,8 @@ int main() { } // free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + checkHipErrors(hipFree(gpuMatrix)); + checkHipErrors(hipFree(gpuTransposeMatrix)); // free the resources on host side free(Matrix); diff --git a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/2dshfl.cpp b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/2dshfl.cpp index 269ad58383..d65af3ace2 100644 --- a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/2dshfl.cpp +++ b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/2dshfl.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. // hip header file #include "hip/hip_runtime.h" +#include "hip_helper.h" #define WIDTH 4 @@ -61,7 +62,7 @@ int main() { float* gpuTransposeMatrix; hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + checkHipErrors(hipGetDeviceProperties(&devProp, 0)); std::cout << "Device name " << devProp.name << std::endl; @@ -78,18 +79,18 @@ int main() { } // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + checkHipErrors(hipMalloc((void**)&gpuMatrix, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float))); // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + checkHipErrors(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); // Lauching kernel from host hipLaunchKernelGGL(matrixTranspose, dim3(1), dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, gpuMatrix, WIDTH); // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + checkHipErrors(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost)); // CPU MatrixTranspose computation matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); @@ -110,8 +111,8 @@ int main() { } // free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + checkHipErrors(hipFree(gpuMatrix)); + checkHipErrors(hipFree(gpuTransposeMatrix)); // free the resources on host side free(Matrix); diff --git a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/CMakeLists.txt index d0ab52859d..80fa6e08b3 100644 --- a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/CMakeLists.txt @@ -39,5 +39,7 @@ set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) # Create the excutable add_executable(2dshfl 2dshfl.cpp) +target_include_directories(2dshfl PRIVATE ../../common) + # Link with HIP target_link_libraries(2dshfl hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile index 91afcfc53a..116d38057a 100644 --- a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile @@ -36,6 +36,7 @@ TARGET=hcc SOURCES = 2dshfl.cpp OBJECTS = $(SOURCES:.cpp=.o) +INCLUDES := -I../../common EXECUTABLE=./2dshfl @@ -44,7 +45,7 @@ EXECUTABLE=./2dshfl all: $(EXECUTABLE) test -CXXFLAGS =-g +CXXFLAGS =-g $(INCLUDES) CXX=$(HIPCC) diff --git a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/CMakeLists.txt index 73c4fe621b..cb90fb6b74 100644 --- a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/CMakeLists.txt @@ -22,6 +22,8 @@ project(dynamic_shared) cmake_minimum_required(VERSION 3.10) +include_directories(../../common) + if (NOT DEFINED ROCM_PATH ) set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." ) endif () @@ -39,5 +41,7 @@ set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) # Create the excutable add_executable(dynamic_shared dynamic_shared.cpp) +target_include_directories(dynamic_shared PRIVATE ../../common) + # Link with HIP target_link_libraries(dynamic_shared hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Makefile b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Makefile index d95ca76085..8db78af246 100644 --- a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Makefile @@ -32,6 +32,7 @@ TARGET=hcc SOURCES = dynamic_shared.cpp OBJECTS = $(SOURCES:.cpp=.o) +INCLUDES := -I../../common EXECUTABLE=./dynamic_shared @@ -40,7 +41,7 @@ EXECUTABLE=./dynamic_shared all: $(EXECUTABLE) test -CXXFLAGS =-g +CXXFLAGS =-g $(INCLUDES) CXX=$(HIPCC) diff --git a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp index 531d94c5be..e538d3e074 100644 --- a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp +++ b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. // hip header file #include "hip/hip_runtime.h" +#include "hip_helper.h" #define WIDTH 16 @@ -65,7 +66,7 @@ int main() { float* gpuTransposeMatrix; hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + checkHipErrors(hipGetDeviceProperties(&devProp, 0)); std::cout << "Device name " << devProp.name << std::endl; @@ -82,11 +83,11 @@ int main() { } // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + checkHipErrors(hipMalloc((void**)&gpuMatrix, NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float))); // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + checkHipErrors(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); // Lauching kernel from host hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), @@ -94,7 +95,7 @@ int main() { 0, gpuTransposeMatrix, gpuMatrix, WIDTH); // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + checkHipErrors(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost)); // CPU MatrixTranspose computation matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); @@ -115,8 +116,8 @@ int main() { } // free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + checkHipErrors(hipFree(gpuMatrix)); + checkHipErrors(hipFree(gpuTransposeMatrix)); // free the resources on host side free(Matrix); diff --git a/projects/hip-tests/samples/2_Cookbook/7_streams/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/7_streams/CMakeLists.txt index 2d95541905..e1133da770 100644 --- a/projects/hip-tests/samples/2_Cookbook/7_streams/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/7_streams/CMakeLists.txt @@ -39,5 +39,7 @@ set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) # Create the excutable add_executable(stream stream.cpp) +target_include_directories(stream PRIVATE ../../common) + # Link with HIP target_link_libraries(stream hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/7_streams/Makefile b/projects/hip-tests/samples/2_Cookbook/7_streams/Makefile index 70dcd4c879..e55d9c8191 100644 --- a/projects/hip-tests/samples/2_Cookbook/7_streams/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/7_streams/Makefile @@ -32,6 +32,7 @@ TARGET=hcc SOURCES = stream.cpp OBJECTS = $(SOURCES:.cpp=.o) +INCLUDES := -I../../common EXECUTABLE=./stream @@ -40,7 +41,7 @@ EXECUTABLE=./stream all: $(EXECUTABLE) test -CXXFLAGS =-g +CXXFLAGS =-g $(INCLUDES) CXX=$(HIPCC) diff --git a/projects/hip-tests/samples/2_Cookbook/7_streams/stream.cpp b/projects/hip-tests/samples/2_Cookbook/7_streams/stream.cpp index 06da516444..89bafb2a31 100644 --- a/projects/hip-tests/samples/2_Cookbook/7_streams/stream.cpp +++ b/projects/hip-tests/samples/2_Cookbook/7_streams/stream.cpp @@ -22,6 +22,7 @@ THE SOFTWARE. #include #include +#include "hip_helper.h" #define WIDTH 32 @@ -66,11 +67,11 @@ void MultipleStream(float** data, float* randArray, float** gpuTransposeMatrix, const int num_streams = 2; hipStream_t streams[num_streams]; - for (int i = 0; i < num_streams; i++) hipStreamCreate(&streams[i]); + for (int i = 0; i < num_streams; i++) checkHipErrors(hipStreamCreate(&streams[i])); for (int i = 0; i < num_streams; i++) { - hipMalloc((void**)&data[i], NUM * sizeof(float)); - hipMemcpyAsync(data[i], randArray, NUM * sizeof(float), hipMemcpyHostToDevice, streams[i]); + checkHipErrors(hipMalloc((void**)&data[i], NUM * sizeof(float))); + checkHipErrors(hipMemcpyAsync(data[i], randArray, NUM * sizeof(float), hipMemcpyHostToDevice, streams[i])); } hipLaunchKernelGGL(matrixTranspose_static_shared, @@ -84,12 +85,12 @@ void MultipleStream(float** data, float* randArray, float** gpuTransposeMatrix, streams[1], gpuTransposeMatrix[1], data[1], width); for (int i = 0; i < num_streams; i++) - hipMemcpyAsync(TransposeMatrix[i], gpuTransposeMatrix[i], NUM * sizeof(float), - hipMemcpyDeviceToHost, streams[i]); + checkHipErrors(hipMemcpyAsync(TransposeMatrix[i], gpuTransposeMatrix[i], NUM * sizeof(float), + hipMemcpyDeviceToHost, streams[i])); } int main() { - hipSetDevice(0); + checkHipErrors(hipSetDevice(0)); float *data[2], *TransposeMatrix[2], *gpuTransposeMatrix[2], *randArray; @@ -100,8 +101,8 @@ int main() { TransposeMatrix[0] = (float*)malloc(NUM * sizeof(float)); TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix[0], NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float)); + checkHipErrors(hipMalloc((void**)&gpuTransposeMatrix[0], NUM * sizeof(float))); + checkHipErrors(hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float))); for (int i = 0; i < NUM; i++) { randArray[i] = (float)i * 1.0f; @@ -109,7 +110,7 @@ int main() { MultipleStream(data, randArray, gpuTransposeMatrix, TransposeMatrix, width); - hipDeviceSynchronize(); + checkHipErrors(hipDeviceSynchronize()); // verify the results int errors = 0; @@ -128,11 +129,11 @@ int main() { free(randArray); for (int i = 0; i < 2; i++) { - hipFree(data[i]); - hipFree(gpuTransposeMatrix[i]); + checkHipErrors(hipFree(data[i])); + checkHipErrors(hipFree(gpuTransposeMatrix[i])); free(TransposeMatrix[i]); } - hipDeviceReset(); + checkHipErrors(hipDeviceReset()); return 0; } diff --git a/projects/hip-tests/samples/2_Cookbook/9_unroll/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/9_unroll/CMakeLists.txt index 258138f2b9..070f66fe23 100644 --- a/projects/hip-tests/samples/2_Cookbook/9_unroll/CMakeLists.txt +++ b/projects/hip-tests/samples/2_Cookbook/9_unroll/CMakeLists.txt @@ -39,5 +39,7 @@ set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) # Create the excutable add_executable(unroll unroll.cpp) +target_include_directories(unroll PRIVATE ../../common) + # Link with HIP target_link_libraries(unroll hip::host) diff --git a/projects/hip-tests/samples/2_Cookbook/9_unroll/Makefile b/projects/hip-tests/samples/2_Cookbook/9_unroll/Makefile index 657f879ac5..49bf51ece1 100644 --- a/projects/hip-tests/samples/2_Cookbook/9_unroll/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/9_unroll/Makefile @@ -36,6 +36,7 @@ TARGET=hcc SOURCES = unroll.cpp OBJECTS = $(SOURCES:.cpp=.o) +INCLUDES := -I../../common EXECUTABLE=./unroll @@ -44,7 +45,7 @@ EXECUTABLE=./unroll all: $(EXECUTABLE) test -CXXFLAGS =-g +CXXFLAGS =-g $(INCLUDES) CXX=$(HIPCC) diff --git a/projects/hip-tests/samples/2_Cookbook/9_unroll/unroll.cpp b/projects/hip-tests/samples/2_Cookbook/9_unroll/unroll.cpp index 18f910a5dd..230ff60112 100644 --- a/projects/hip-tests/samples/2_Cookbook/9_unroll/unroll.cpp +++ b/projects/hip-tests/samples/2_Cookbook/9_unroll/unroll.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. // hip header file #include "hip/hip_runtime.h" +#include "hip_helper.h" #define LENGTH 4 @@ -59,7 +60,7 @@ int main() { int* gpuSumMatrix; hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + checkHipErrors(hipGetDeviceProperties(&devProp, 0)); std::cout << "Device name " << devProp.name << std::endl; @@ -76,19 +77,19 @@ int main() { } // Allocated Device Memory - hipMalloc((void**)&gpuMatrix, SIZE * sizeof(int)); - hipMalloc((void**)&gpuSumMatrix, LENGTH * sizeof(int)); + checkHipErrors(hipMalloc((void**)&gpuMatrix, SIZE * sizeof(int))); + checkHipErrors(hipMalloc((void**)&gpuSumMatrix, LENGTH * sizeof(int))); // Memory Copy to Device - hipMemcpy(gpuMatrix, Matrix, SIZE * sizeof(int), hipMemcpyHostToDevice); - hipMemcpy(gpuSumMatrix, cpuSumMatrix, LENGTH * sizeof(float), hipMemcpyHostToDevice); + checkHipErrors(hipMemcpy(gpuMatrix, Matrix, SIZE * sizeof(int), hipMemcpyHostToDevice)); + checkHipErrors(hipMemcpy(gpuSumMatrix, cpuSumMatrix, LENGTH * sizeof(float), hipMemcpyHostToDevice)); // Launch device kernels hipLaunchKernelGGL(gpuMatrixRowSum, dim3(BLOCKS_PER_GRID), dim3(THREADS_PER_BLOCK), 0, 0, gpuMatrix, gpuSumMatrix, LENGTH); // Memory copy back to device - hipMemcpy(sumMatrix, gpuSumMatrix, LENGTH * sizeof(int), hipMemcpyDeviceToHost); + checkHipErrors(hipMemcpy(sumMatrix, gpuSumMatrix, LENGTH * sizeof(int), hipMemcpyDeviceToHost)); // Cpu implementation matrixRowSum(Matrix, cpuSumMatrix, LENGTH); @@ -110,8 +111,8 @@ int main() { } // GPU Free - hipFree(gpuMatrix); - hipFree(gpuSumMatrix); + checkHipErrors(hipFree(gpuMatrix)); + checkHipErrors(hipFree(gpuSumMatrix)); // CPU Free free(Matrix); diff --git a/projects/hip-tests/samples/common/hip_helper.h b/projects/hip-tests/samples/common/hip_helper.h new file mode 100644 index 0000000000..4b25346c92 --- /dev/null +++ b/projects/hip-tests/samples/common/hip_helper.h @@ -0,0 +1,38 @@ +/* +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" + +#ifndef checkHipErrors +#define checkHipErrors(err) __checkHipErrors(err, __FILE__, __LINE__) + +inline void __checkHipErrors(hipError_t err, const char *file, const int line) { + if (HIP_SUCCESS != err) { + const char *errorStr = hipGetErrorString(err); + fprintf(stderr, + "checkHipErrors() HIP API error = %04d \"%s\" from file <%s>, " + "line %i.\n", + err, errorStr, file, line); + exit(EXIT_FAILURE); + } +} +#endif From 27a7591e94cb9593156f0b684def742921c8e7ad Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 26 Jul 2023 16:54:35 +0530 Subject: [PATCH 17/17] Temporarily disable unstable tests (#376) * Update config_amd_linux_MI2xx.json * Update config_amd_linux_common.json [ROCm/hip-tests commit: f7295e3ea3351ba25bda7afa0204efca0ddf748e] --- .../catch/hipTestMain/config/config_amd_linux_MI2xx.json | 3 ++- .../catch/hipTestMain/config/config_amd_linux_common.json | 5 ++++- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_MI2xx.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_MI2xx.json index e8a37c2830..f1a914fa6c 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_MI2xx.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_MI2xx.json @@ -16,7 +16,8 @@ "Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores", "Unit_hipSignalExternalSemaphoresAsync_Vulkan_Negative_Parameters", "Unit_hipImportExternalSemaphore_Vulkan_Negative_Parameters", - "Unit_hipDestroyExternalSemaphore_Vulkan_Negative_Parameters" + "Unit_hipDestroyExternalSemaphore_Vulkan_Negative_Parameters", + "Unit_Grid_Group_Sync_Positive_Basic" ] } diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index b8bf15e523..e6b83d620a 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -110,6 +110,9 @@ "=== Below tests fail in stress test on 13/07/23 ===", "Unit_deviceAllocation_Malloc_ComplexDataType", "Unit_deviceAllocation_New_ComplexDataType", - "Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_1" + "Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_1", + "=== Below tests fail in stress test on 24/07/23 ===", + "Unit_hipStreamCreateWithPriority_ValidateWithEvents", + "Unit_hipEventIpc" ] }