AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge 'develop' into 'amd-staging'

Просмотр исходного кода 
Change-Id: Ice59310b128db9e211da42b3a2f2dd108e076122


[ROCm/hip-tests commit: 777c0c4288]
Этот коммит содержится в:
Jenkins
2023-07-26 23:10:14 +00:00
родитель c5c3441e02 27a7591e94
Коммит 7b987fb698
111 изменённых файлов: 2795 добавлений и 1350 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/hip-tests/catch/CMakeLists.txt
+1
Просмотреть файл
@@ -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)
projects/hip-tests/catch/hipTestMain/config/config_amd_linux_MI2xx.json
+2 -1
Просмотреть файл
@@ -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"
]
}
projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json
+8 -14
Просмотреть файл
@@ -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",
@@ -63,6 +60,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",
@@ -77,19 +76,11 @@
"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",
"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_hipMemsetDSync - int8_t",
"Unit_hipGraphClone_Test_hipGraphExecMemcpyNodeSetParams",
"Unit_hipGraphClone_Test_hipGraphMemcpyNodeSetParams1D_and_exec",
"Unit_hipStreamValue_Wait64_Blocking_NoMask_And",
@@ -119,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"
]
}
projects/hip-tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json
+1
Просмотреть файл
@@ -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",
projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json
+6
Просмотреть файл
@@ -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",
@@ -180,10 +181,15 @@
"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",
"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 ===",
projects/hip-tests/catch/hipTestMain/main.cc
+30 -2
Просмотреть файл
@@ -1,7 +1,10 @@
#define CATCH_CONFIG_RUNNER
#include <cmd_options.hh>
#include <hip_test_common.hh>
#include <iostream>
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;
}
projects/hip-tests/catch/include/cmd_options.hh
+33
Просмотреть файл
@@ -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;
projects/hip-tests/catch/include/cpu_grid.h
+154
Просмотреть файл
@@ -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 <optional>
#include <hip_test_common.hh>
#include <hip/hip_runtime_api.h>
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<unsigned int> 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<dim3> 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<dim3> 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<int>(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<int>(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<int>(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<int>(i * sm), 1, 1); },
values(multipliers)),
map([sm = props.multiProcessorCount](
double i) { return dim3(1, static_cast<int>(i * sm), 1); },
values(multipliers)),
map([sm = props.multiProcessorCount](
double i) { return dim3(1, 1, static_cast<int>(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<int>(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<int>(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<int>(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<int>(i * sm), 1, 1); },
values(multipliers)),
map([sm = props.multiProcessorCount](
double i) { return dim3(1, static_cast<int>(i * sm), 1); },
values(multipliers)),
map([sm = props.multiProcessorCount](
double i) { return dim3(1, 1, static_cast<int>(i * sm)); },
values(multipliers)),
dim3(5, 5, 5));
}
projects/hip-tests/catch/include/hip_test_common.hh
-81
Просмотреть файл
@@ -350,87 +350,6 @@ template <> struct MemTraits<MemcpyAsync> {
}
};
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<size_t>(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;
projects/hip-tests/catch/include/hip_test_defgroups.hh
+20
Просмотреть файл
@@ -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
* @{
@@ -128,3 +135,16 @@ THE SOFTWARE.
* This section describes the various kernel functions invocation.
* @}
*/
/**
* @defgroup DeviceLanguageTest Device Language
* @{
* This section describes tests for the Device Language API.
* @}
*/
/**
* @defgroup DeviceLanguageTest Device Language
* @{
* This section describes tests for the Device Language API.
*/
projects/hip-tests/catch/include/hip_test_process.hh
+7
Просмотреть файл
@@ -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) {
projects/hip-tests/catch/include/performance_common.hh
+247
Просмотреть файл
@@ -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 <algorithm>
#include <chrono>
#include <memory>
#include <numeric>
#include <type_traits>
#include <vector>
#include <cmd_options.hh>
#include <hip_test_common.hh>
#include <resource_guards.hh>
#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<float, std::milli> ms = stop_ - start_;
Record(ms.count());
}
private:
std::chrono::time_point<std::chrono::steady_clock> start_;
std::chrono::time_point<std::chrono::steady_clock> stop_;
};
template <typename Derived> 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<float(float)>;
void RegisterModifier(const ModifierSignature& modifier) { modifier_ = modifier; }
template <typename... Args> std::tuple<float, float, float, float> Run(Args&&... args) {
AddSectionName(std::to_string(iterations_));
AddSectionName(std::to_string(warmups_));
auto& derived = static_cast<Derived&>(*this);
current_ = kWarmup;
for (size_t i = 0u; i < warmups_; ++i) {
PrintProgress("warmup", static_cast<int>(100.f * (i + 1) / warmups_));
derived(args...);
}
time_ = .0;
std::vector<float> samples;
samples.reserve(iterations_);
for (current_ = 0; current_ < iterations_; ++current_) {
PrintProgress("measurement", static_cast<int>(100.f * (current_ + 1) / iterations_));
derived(args...);
if (modifier_) time_ = modifier_(time_);
samples.push_back(time_);
time_ = .0;
}
float sum = std::accumulate(cbegin(samples), cend(samples), .0);
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 <bool event_based>
using TimerType = std::conditional_t<event_based, EventTimer, CpuTimer>;
template <bool event_based = false>
std::unique_ptr<TimerType<event_based>> GetTimer(hipStream_t stream = nullptr) {
return std::make_unique<TimerType<event_based>>(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<TIMER_TYPE>(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";
}
}
projects/hip-tests/catch/include/utils.hh
+32 -9
Просмотреть файл
@@ -20,6 +20,7 @@ THE SOFTWARE.
#pragma once
#include <chrono>
#include <optional>
#include <hip_test_common.hh>
#include <hip/hip_runtime_api.h>
@@ -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 <typename T, typename F>
static inline void ArrayAllOf(const T* arr, uint32_t count, F value_gen) {
for (auto i = 0u; i < count; ++i) {
const std::optional<T> 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 <typename T, typename F>
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) {
@@ -107,9 +122,17 @@ template <typename T> __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
}
}
@@ -125,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);
}
projects/hip-tests/catch/performance/CMakeLists.txt
+21
Просмотреть файл
@@ -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)
projects/hip-tests/catch/performance/example/CMakeLists.txt
+28
Просмотреть файл
@@ -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)
projects/hip-tests/catch/performance/example/example.cc
+59
Просмотреть файл
@@ -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 <hip_test_common.hh>
#include <performance_common.hh>
#include <resource_guards.hh>
class ExampleBenchmark : public Benchmark<ExampleBenchmark> {
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<void> dst(LinearAllocs::hipMalloc, 4_MB);
benchmark.Run(dst.ptr());
}
projects/hip-tests/catch/unit/CMakeLists.txt
+2 -1
Просмотреть файл
@@ -44,7 +44,8 @@ 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()
add_subdirectory(vector_types)
projects/hip-tests/catch/unit/clock/hipClockCheck.cc
+142 -67
Просмотреть файл
@@ -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 <hip_test_checkers.hh>
#include <hip/hip_ext.h>
#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));
}
projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt
+1
Просмотреть файл
@@ -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
projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh
+71
Просмотреть файл
@@ -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 <hip_test_common.hh>
#include <hip/hip_cooperative_groups.h>
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 <class T> 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;
}
projects/hip-tests/catch/unit/cooperativeGrps/grid_group.cc
+285
Просмотреть файл
@@ -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 <cpu_grid.h>
#include <resource_guards.hh>
#include <utils.hh>
/**
* @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<unsigned int> uint_arr_dev(LinearAllocs::hipMalloc,
grid.thread_count_ * sizeof(unsigned int));
LinearAllocGuard<unsigned int> 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<unsigned int> uint_arr_dev(LinearAllocs::hipMalloc,
grid.thread_count_ * sizeof(unsigned int));
LinearAllocGuard<unsigned int> 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<unsigned int> uint_arr_dev(LinearAllocs::hipMalloc,
array_len * sizeof(unsigned int));
LinearAllocGuard<unsigned int> uint_arr(LinearAllocs::hipHostMalloc,
array_len * sizeof(unsigned int));
LinearAllocGuard<unsigned int> 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<void*>(&atomic_val_ptr);
params[1] = reinterpret_cast<void*>(&uint_arr_dev_ptr);
params[2] = reinterpret_cast<void*>(&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);
}
}
projects/hip-tests/catch/unit/event/Unit_hipEventElapsedTime.cc
+2 -6
Просмотреть файл
@@ -23,7 +23,7 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <iostream>
#include <utils.hh>
/**
* @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));
projects/hip-tests/catch/unit/event/Unit_hipEventQuery.cc
+23 -8
Просмотреть файл
@@ -18,7 +18,29 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <utils.hh>
/**
* @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
projects/hip-tests/catch/unit/event/hipEventDestroy.cc
+2 -2
Просмотреть файл
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include <hip_test_kernels.hh>
#include <hip_test_common.hh>
#include "hip/hip_runtime_api.h"
#include <utils.hh>
/**
* @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<void**>(&A_d), A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&B_d), B_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&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);
}
projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc
+817 -5
Просмотреть файл
@@ -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 <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#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<int*>(malloc(Nbytes));
REQUIRE(A1_h != nullptr);
A2_h = reinterpret_cast<int*>(malloc(Nbytes));
REQUIRE(A2_h != nullptr);
A3_h = reinterpret_cast<int*>(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<void *>(ker_vec_mul);
kerNodeParams1.gridDim = dim3(blocks);
kerNodeParams1.blockDim = dim3(threadsPerBlock);
kerNodeParams1.sharedMemBytes = 0;
kerNodeParams1.kernelParams = reinterpret_cast<void**>(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<void *>(ker_vec_mul);
kerNodeParams2.gridDim = dim3(blocks);
kerNodeParams2.blockDim = dim3(threadsPerBlock);
kerNodeParams2.sharedMemBytes = 0;
kerNodeParams2.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kerNodeParams2.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&vec_mul2, graph[2], nullptr, 0,
&kerNodeParams2));
void* kernelArgs3[] = {&C1_d, &C2_d};
kerNodeParams3.func =
reinterpret_cast<void *>(ker_vec_add);
kerNodeParams3.gridDim = dim3(blocks);
kerNodeParams3.blockDim = dim3(threadsPerBlock);
kerNodeParams3.sharedMemBytes = 0;
kerNodeParams3.kernelParams = reinterpret_cast<void**>(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<void*>(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<void*>(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<void *>(ker_vec_sqr);
kerNodeParams3.gridDim = dim3(blocks);
kerNodeParams3.blockDim = dim3(threadsPerBlock);
kerNodeParams3.sharedMemBytes = 0;
kerNodeParams3.kernelParams = reinterpret_cast<void**>(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<hipGraphNode_t *>(
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<hipGraphNode_t *>(
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<void *>(ker_vec_add)) {
nodeIdx = idx;
break;
}
}
}
if (updatetype == updateGraphNodeTests::updateFunKerNodParamTest) {
nodeParam.func = reinterpret_cast<void *>(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<void *>(ker_vec_sub);
kerNodeParams3.gridDim = dim3(blocks);
kerNodeParams3.blockDim = dim3(threadsPerBlock);
kerNodeParams3.sharedMemBytes = 0;
kerNodeParams3.kernelParams = reinterpret_cast<void**>(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<void *>(ker_vec_sub);
kerNodeParams3.gridDim = dim3(blocks);
kerNodeParams3.blockDim = dim3(threadsPerBlock);
kerNodeParams3.sharedMemBytes = 0;
kerNodeParams3.kernelParams = reinterpret_cast<void**>(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<hipGraphNode_t> 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<int>(&inputVec_d1, &outputVec_d1, nullptr,
&inputVec_h1, &outputVec_h1, nullptr, size, false);
HipTest::initArrays<int>(&inputVec_d2, &outputVec_d2, nullptr,
&inputVec_h2, &outputVec_h2, nullptr, size, false);
HipTest::initArrays<int>(&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<void*>(&inputVec_d1),
reinterpret_cast<void*>(&outputVec_d1),
reinterpret_cast<void*>(&size)};
kerNodeParams1.func = reinterpret_cast<void*>(HipTest::vector_square<int>);
kerNodeParams1.gridDim = dim3(blocks);
kerNodeParams1.blockDim = dim3(threadsPerBlock);
kerNodeParams1.sharedMemBytes = 0;
kerNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs1);
kerNodeParams1.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&vecSqr1, graph2, nullptr, 0,
&kerNodeParams1));
void* kernelArgs2[] = {reinterpret_cast<void*>(&inputVec_d2),
reinterpret_cast<void*>(&outputVec_d2),
reinterpret_cast<void*>(&size)};
kerNodeParams2.func = reinterpret_cast<void*>(HipTest::vector_square<int>);
kerNodeParams2.gridDim = dim3(blocks);
kerNodeParams2.blockDim = dim3(threadsPerBlock);
kerNodeParams2.sharedMemBytes = 0;
kerNodeParams2.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kerNodeParams2.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&vecSqr2, graph2, nullptr, 0,
&kerNodeParams2));
void* kernelArgs3[] = {reinterpret_cast<void*>(&inputVec_d3),
reinterpret_cast<void*>(&outputVec_d3),
reinterpret_cast<void*>(&size)};
kerNodeParams3.func = reinterpret_cast<void*>(HipTest::vector_square<int>);
kerNodeParams3.gridDim = dim3(blocks);
kerNodeParams3.blockDim = dim3(threadsPerBlock);
kerNodeParams3.sharedMemBytes = 0;
kerNodeParams3.kernelParams = reinterpret_cast<void**>(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<int>(inputVec_d1, outputVec_d1, nullptr,
inputVec_h1, outputVec_h1, nullptr, false);
HipTest::freeArrays<int>(inputVec_d2, outputVec_d2, nullptr,
inputVec_h2, outputVec_h2, nullptr, false);
HipTest::freeArrays<int>(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;
}
projects/hip-tests/catch/unit/memory/hipArrayCommon.hh
-6
Просмотреть файл
@@ -67,12 +67,6 @@ template <typename T> void checkDataIsAscending(const std::vector<T>& 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<int, 2> max2D;
projects/hip-tests/catch/unit/memory/hipArrayCreate.cc
+2 -15
Просмотреть файл
@@ -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);
}
}
projects/hip-tests/catch/unit/memory/hipFree.cc
+6 -6
Просмотреть файл
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include <hip_array_common.hh>
#include "hipArrayCommon.hh"
#include "DriverContext.hh"
#include <utils.hh>
/*
* 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<TestType>();
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<TestType>();
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
projects/hip-tests/catch/unit/memory/hipMalloc3D.cc
+1 -15
Просмотреть файл
@@ -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);
}
}
projects/hip-tests/catch/unit/memory/hipMemGetInfo.cc
-556
Просмотреть файл
@@ -17,563 +17,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <thread>
#include <vector>
/*
* 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<void**>(&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<void**>(&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<void**>(&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<void**>(&A_mem), Malloc1Size));
// allocate an extra quarter of free mem
auto Malloc2Size = Malloc1Size >> 1;
fixAllocSize(Malloc2Size);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&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<void**>(&A_mem), MallocSize));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&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<void**>(&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<TestType>();
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<TestType>();
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<size_t>(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<void**>(&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<void**>(&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<void**>(&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<void**>(&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<void**>(&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<void**>(&A_mem), MallocSize)); });
std::thread t2(
[&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast<void**>(&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<void**>(&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<void**>(&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<void**>(&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<void**>(&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};
projects/hip-tests/catch/unit/memory/hipMemcpySync.cc
+2 -2
Просмотреть файл
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
#include "MemUtils.hh"
#include <utils.hh>
/*
* 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<uint64_t, std::milli> delay = 100ms;
HipTest::runKernelForDuration(delay, stream);
LaunchDelayKernel(delay, stream);
memcpyCheck(type, memType, aPtr.first, data, fillerData, async, stream, fromHost);
checkForSync(stream, async, type, fromHost);
projects/hip-tests/catch/unit/memory/hipMemsetSync.cc
+15 -12
Просмотреть файл
@@ -21,6 +21,7 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <utils.hh>
/*
* 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
@@ -51,8 +52,8 @@ struct MultiDData {
// set of helper functions to tidy the nested switch statements
template <typename T>
static std::pair<T*,T*> deviceMallocHelper(memSetType memType, size_t dataW, size_t dataH, size_t dataD,
size_t& dataPitch) {
static std::pair<T*, T*> 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{};
@@ -88,7 +89,8 @@ static std::pair<T*,T*> deviceMallocHelper(memSetType memType, size_t dataW, siz
}
template <typename T>
static std::pair<T*, T*> hostMallocHelper(size_t dataW, size_t dataH, size_t dataD, size_t& dataPitch) {
static std::pair<T*, T*> 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;
@@ -100,7 +102,8 @@ static std::pair<T*, T*> hostMallocHelper(size_t dataW, size_t dataH, size_t dat
}
template <typename T>
static std::pair<T*, T*> hostRegisteredHelper(size_t dataW, size_t dataH, size_t dataD, size_t& dataPitch) {
static std::pair<T*, T*> 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];
@@ -406,7 +409,7 @@ void runTests(allocType type, memSetType memsetType, MultiDData data, hipStream_
std::pair<T*, T*> aPtr = initMemory<T>(type, memsetType, data);
using namespace std::chrono_literals;
const std::chrono::duration<uint64_t, std::milli> delay = 100ms;
HipTest::runKernelForDuration(delay, stream);
LaunchDelayKernel(delay, stream);
memsetCheck(aPtr.first, testValue, memsetType, data, async, stream);
if (async || type == allocType::deviceMalloc) {
@@ -447,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<char>(type, memset_type, data);
}
@@ -460,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<int8_t, TestType>::value) {
memset_type = memSetType::hipMemsetD8;
@@ -482,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<char>(mallocType, memset_type, data);
}
@@ -497,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<char>(mallocType, memset_type, data);
}
projects/hip-tests/catch/unit/stream/hipStreamACb_MultiThread.cc
+3 -29
Просмотреть файл
@@ -25,6 +25,7 @@ multiple Threads.
#include <hip_test_common.hh>
#include <atomic>
#include <utils.hh>
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,
projects/hip-tests/catch/unit/stream/hipStreamCreateWithFlags.cc
+3 -2
Просмотреть файл
@@ -19,6 +19,7 @@ THE SOFTWARE.
#include <chrono>
#include <hip_test_common.hh>
#include <utils.hh>
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);
}
projects/hip-tests/catch/unit/stream/hipStreamDestroy.cc
+2 -1
Просмотреть файл
@@ -18,6 +18,7 @@ THE SOFTWARE.
*/
#include <chrono>
#include <hip_test_common.hh>
#include <utils.hh>
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);
projects/hip-tests/catch/unit/stream/hipStreamQuery.cc
+4 -5
Просмотреть файл
@@ -19,7 +19,7 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include "streamCommon.hh"
#include <utils.hh>
/**
* @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));
projects/hip-tests/catch/unit/stream/hipStreamSynchronize.cc
+9 -9
Просмотреть файл
@@ -19,7 +19,7 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include "streamCommon.hh"
#include <utils.hh>
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);
projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc
+4 -48
Просмотреть файл
@@ -25,7 +25,7 @@ Unit_hipStreamWaitEvent_DifferentStreams - Test waiting for an event on a differ
*/
#include <hip_test_common.hh>
#include <utils.hh>
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));
projects/hip-tests/catch/unit/vector_types/CMakeLists.txt
+29
Просмотреть файл
@@ -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)
projects/hip-tests/catch/unit/vector_types/dim3.cc
+261
Просмотреть файл
@@ -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 <hip_test_common.hh>
/**
* @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<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::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<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::max());
uint32_t value_y =
GENERATE(std::numeric_limits<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::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<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::max());
uint32_t value_y =
GENERATE(std::numeric_limits<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::max());
uint32_t value_z =
GENERATE(std::numeric_limits<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::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<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::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<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::max());
uint32_t value_y =
GENERATE(std::numeric_limits<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::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<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::max());
uint32_t value_y =
GENERATE(std::numeric_limits<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::max());
uint32_t value_z =
GENERATE(std::numeric_limits<uint32_t>::min(), std::numeric_limits<uint32_t>::max() / 2,
std::numeric_limits<uint32_t>::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);
}
projects/hip-tests/samples/0_Intro/bit_extract/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/0_Intro/bit_extract/Makefile
+2 -1
Просмотреть файл
@@ -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)
projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp
+10 -19
Просмотреть файл
@@ -23,16 +23,7 @@ THE SOFTWARE.
#include <stdio.h>
#include <iostream>
#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");
projects/hip-tests/samples/0_Intro/module_api/CMakeLists.txt
+2
Просмотреть файл
@@ -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 ()
projects/hip-tests/samples/0_Intro/module_api/Makefile
+5 -4
Просмотреть файл
@@ -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
projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp
+14 -13
Просмотреть файл
@@ -24,6 +24,7 @@ THE SOFTWARE.
#include <iostream>
#include <fstream>
#include <vector>
#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;
}
projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp
+14 -19
Просмотреть файл
@@ -25,6 +25,7 @@ THE SOFTWARE.
#include <iostream>
#include <fstream>
#include <vector>
#include "hip_helper.h"
#ifdef __HIP_PLATFORM_AMD__
#include <hip/hip_ext.h>
@@ -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;
}
projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp
+13 -18
Просмотреть файл
@@ -26,6 +26,7 @@ THE SOFTWARE.
#include <fstream>
#include <vector>
#include <hip/hip_hcc.h>
#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;
}
projects/hip-tests/samples/0_Intro/module_api_global/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/0_Intro/module_api_global/Makefile
+2 -1
Просмотреть файл
@@ -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 $@
projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp
+12 -12
Просмотреть файл
@@ -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);
projects/hip-tests/samples/1_Utils/hipDispatchLatency/CMakeLists.txt
+2
Просмотреть файл
@@ -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 ()
projects/hip-tests/samples/1_Utils/hipDispatchLatency/Makefile
+2 -1
Просмотреть файл
@@ -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
projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp
+9 -18
Просмотреть файл
@@ -22,6 +22,7 @@ THE SOFTWARE.
#ifdef __HIP_PLATFORM_AMD__
#include "hip/hip_ext.h"
#endif
#include "hip_helper.h"
#include <iostream>
#include <fstream>
#include <chrono>
@@ -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<char>& 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<float, TOTAL_RUN_COUNT> results;
@@ -103,13 +94,13 @@ void hipModuleLaunchKernel_enqueue_rate(const std::vector<char>& 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<double, std::milli>(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<char>& buffer, std::atomi
{
//resources necessary for this thread
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
checkHipErrors(hipStreamCreate(&stream));
std::array<float, TOTAL_RUN_COUNT> results;
//synchronize all threads, before running
@@ -131,7 +122,7 @@ void hipLaunchKernelGGL_enqueue_rate(const std::vector<char>& buffer, std::atomi
results[i] = std::chrono::duration<double, std::milli>(stop - start).count();
}
print_timing("Thread ID : " + std::to_string(tid) + " , " + "hipLaunchKernelGGL enqueue rate", results);
HIPCHECK(hipStreamDestroy(stream));
checkHipErrors(hipStreamDestroy(stream));
}
// Simple thread pool
projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
+19 -18
Просмотреть файл
@@ -21,6 +21,7 @@ THE SOFTWARE.
#ifdef __HIP_PLATFORM_AMD__
#include "hip/hip_ext.h"
#endif
#include "hip_helper.h"
#include <iostream>
#include <chrono>
#include <algorithm>
@@ -66,19 +67,19 @@ void print_timing(std::string test, const std::array<float, TOTAL_RUN_COUNT> &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<float, TOTAL_RUN_COUNT> 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, &params, nullptr);
checkHipErrors(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, &params, nullptr));
auto stop = std::chrono::high_resolution_clock::now();
results[i] = std::chrono::duration<float, std::milli>(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));
}
projects/hip-tests/samples/1_Utils/hipInfo/CMakeLists.txt
+2
Просмотреть файл
@@ -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
projects/hip-tests/samples/1_Utils/hipInfo/Makefile
+2 -1
Просмотреть файл
@@ -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
projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp
+8 -21
Просмотреть файл
@@ -23,6 +23,7 @@ THE SOFTWARE.
#include <iostream>
#include <iomanip>
#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);
}
projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Makefile
+2 -1
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp
+8 -7
Просмотреть файл
@@ -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);
projects/hip-tests/samples/2_Cookbook/10_inline_asm/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/10_inline_asm/Makefile
+2 -2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp
+22 -21
Просмотреть файл
@@ -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);
projects/hip-tests/samples/2_Cookbook/11_texture_driver/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/11_texture_driver/Makefile
+2 -1
Просмотреть файл
@@ -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 $@
projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
+17 -26
Просмотреть файл
@@ -24,21 +24,12 @@ THE SOFTWARE.
#include <iostream>
#include <fstream>
#include <vector>
#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<typename T,
typename std::enable_if<std::is_arithmetic<T>::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<T>();
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<char>(module, "texChar", "tex2dKernelChar");
testResult = testResult && runTest<short>(module, "texShort", "tex2dKernelShort");
testResult = testResult && runTest<int>(module, "texInt", "tex2dKernelInt");
@@ -255,7 +246,7 @@ int main(int argc, char** argv) {
testResult = testResult && runTest<int4>(module, "texInt4", "tex2dKernelInt4");
testResult = testResult && runTest<float4>(module, "texFloat4", "tex2dKernelFloat4");
HIP_CHECK(hipModuleUnload(module));
checkHipErrors(hipModuleUnload(module));
printf("texture2dDrv %s ...\n", testResult ? "PASSED" : "FAILED");
return testResult ? EXIT_SUCCESS : EXIT_FAILURE;
}
projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/MatrixTranspose.cpp
+8 -7
Просмотреть файл
@@ -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);
projects/hip-tests/samples/2_Cookbook/13_occupancy/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/13_occupancy/Makefile
+2 -2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp
+22 -27
Просмотреть файл
@@ -19,14 +19,9 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#include <iostream>
#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);
projects/hip-tests/samples/2_Cookbook/14_gpu_arch/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/14_gpu_arch/Makefile
+2 -2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp
+4 -10
Просмотреть файл
@@ -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<int32_t*>(malloc(NBytes));
int32_t* hOutput = static_cast<int32_t*>(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
projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/Makefile
+3 -2
Просмотреть файл
@@ -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:
projects/hip-tests/samples/2_Cookbook/16_assembly_to_executable/square.cpp
+10 -18
Просмотреть файл
@@ -19,15 +19,7 @@ THE SOFTWARE.
#include <stdio.h>
#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"
/* 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<N; i++)
{
@@ -70,12 +62,12 @@ int main(int argc, char *argv[])
}
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));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
@@ -84,12 +76,12 @@ int main(int argc, char *argv[])
vector_square <<<blocks, threadsPerBlock>>> (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<N; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
CHECK(hipErrorUnknown);
checkHipErrors(hipErrorUnknown);
}
}
printf ("PASSED!\n");
projects/hip-tests/samples/2_Cookbook/17_llvm_ir_to_executable/Makefile
+3 -2
Просмотреть файл
@@ -32,6 +32,7 @@ LLVM_MC=$(HIP_PATH)/llvm/bin/llvm-mc
CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/llvm/bin/clang-offload-bundler
LLVM_AS=$(HIP_PATH)/llvm/bin/llvm-as
LLVM_DIS=$(HIP_PATH)/llvm/bin/llvm-dis
INCLUDES := -I../../common
SRCS=square.cpp
@@ -60,8 +61,8 @@ GPU_ARCH9=gfx1103
all: src_to_ir bc_to_ll ll_to_bc ir_to_exec
src_to_ir:
$(HIPCC) -c -emit-llvm --cuda-host-only -target x86_64-linux-gnu -o $(SQ_HOST_BC) $(SRCS)
$(HIPCC) -c -emit-llvm --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) $(INCLUDES) -c -emit-llvm --cuda-host-only -target x86_64-linux-gnu -o $(SQ_HOST_BC) $(SRCS)
$(HIPCC) $(INCLUDES) -c -emit-llvm --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)
# By default, the LLVM IR Bitcode file names will be:
# square-hip-amdgcn-amd-amdhsa-gfx900.bc
projects/hip-tests/samples/2_Cookbook/17_llvm_ir_to_executable/square.cpp
+10 -18
Просмотреть файл
@@ -19,15 +19,7 @@ THE SOFTWARE.
#include <stdio.h>
#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"
/* 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<N; i++)
{
@@ -70,12 +62,12 @@ int main(int argc, char *argv[])
}
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));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
@@ -84,12 +76,12 @@ int main(int argc, char *argv[])
vector_square <<<blocks, threadsPerBlock>>> (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<N; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
CHECK(hipErrorUnknown);
checkHipErrors(hipErrorUnknown);
}
}
printf ("PASSED!\n");
projects/hip-tests/samples/2_Cookbook/18_cmake_hip_device/CMakeLists.txt
+4
Просмотреть файл
@@ -22,11 +22,15 @@ project(cmake_hip_device_test)
cmake_minimum_required(VERSION 3.10.2)
include_directories(../../common)
# Find hip
find_package(hip REQUIRED)
# Create the excutable
add_executable(test_cpp square.cpp)
target_include_directories(test_cpp PRIVATE ../../common)
# Link with HIP
target_link_libraries(test_cpp hip::device)
projects/hip-tests/samples/2_Cookbook/18_cmake_hip_device/square.cpp
+10 -19
Просмотреть файл
@@ -22,16 +22,7 @@ THE SOFTWARE.
#include <stdio.h>
#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"
/*
* 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<N; i++)
{
@@ -72,12 +63,12 @@ int main(int argc, char *argv[])
}
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));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
@@ -86,12 +77,12 @@ int main(int argc, char *argv[])
hipLaunchKernelGGL(vector_square, 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");
printf ("info: checkHipErrors result\n");
for (size_t i=0; i<N; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
CHECK(hipErrorUnknown);
checkHipErrors(hipErrorUnknown);
}
}
printf ("PASSED!\n");
projects/hip-tests/samples/2_Cookbook/19_cmake_lang/CMakeLists.txt
+3
Просмотреть файл
@@ -10,5 +10,8 @@ add_executable(test_fortran TestFortran.F90)
add_executable(test_cpp MatrixTranspose.cpp)
target_link_libraries(test_cpp PUBLIC hip::device)
target_include_directories(test_cpp PRIVATE ../../common)
# Assuming to build a C/C++-to-Fortran library binding.
target_link_libraries(test_fortran PUBLIC hip::device)
projects/hip-tests/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp
+8 -8
Просмотреть файл
@@ -24,7 +24,7 @@ THE SOFTWARE.
// hip header file
#include "hip/hip_runtime.h"
#include "hip_helper.h"
#define WIDTH 1024
@@ -61,7 +61,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 +78,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 +90,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 +110,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);
projects/hip-tests/samples/2_Cookbook/1_hipEvent/CMakeLists.txt
+2
Просмотреть файл
@@ -40,5 +40,7 @@ set(CMAKE_BUILD_TYPE Release)
# Create the excutable
add_executable(hipEvent hipEvent.cpp)
target_include_directories(hipEvent PRIVATE ../../common)
# Link with HIP
target_link_libraries(hipEvent hip::host)
projects/hip-tests/samples/2_Cookbook/1_hipEvent/Makefile
+2 -1
Просмотреть файл
@@ -30,6 +30,7 @@ TARGET=hcc
SOURCES = hipEvent.cpp
OBJECTS = $(SOURCES:.cpp=.o)
INCLUDES := -I../../common
EXECUTABLE=./hipEvent
@@ -38,7 +39,7 @@ EXECUTABLE=./hipEvent
all: $(EXECUTABLE) test
CXXFLAGS =-g
CXXFLAGS =-g $(INCLUDES)
CXX=$(HIPCC)
projects/hip-tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp
+22 -21
Просмотреть файл
@@ -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);
@@ -146,8 +147,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);
projects/hip-tests/samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt
+2
Просмотреть файл
@@ -30,5 +30,7 @@ find_package(hip REQUIRED)
# Create the excutable
add_executable(square square.cpp)
target_include_directories(square PRIVATE ../../common)
# Link with HIP
target_link_libraries(square hip::device)
projects/hip-tests/samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp
+9 -18
Просмотреть файл
@@ -22,16 +22,7 @@ THE SOFTWARE.
#include <stdio.h>
#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"
/*
* 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<N; i++)
{
@@ -72,12 +63,12 @@ int main(int argc, char *argv[])
}
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));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
@@ -86,12 +77,12 @@ int main(int argc, char *argv[])
hipLaunchKernelGGL(vector_square, 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++) {
if (C_h[i] != A_h[i] * A_h[i]) {
CHECK(hipErrorUnknown);
checkHipErrors(hipErrorUnknown);
}
}
printf ("PASSED!\n");
projects/hip-tests/samples/2_Cookbook/22_cmake_hip_lang/CMakeLists.txt
+2
Просмотреть файл
@@ -25,3 +25,5 @@ project(cmake_hip_lang_support VERSION 1.0
LANGUAGES HIP)
# Create the executable
add_executable(square square.hip)
target_include_directories(square PRIVATE ../../common)
projects/hip-tests/samples/2_Cookbook/22_cmake_hip_lang/square.hip
+9 -18
Просмотреть файл
@@ -22,16 +22,7 @@ THE SOFTWARE.
#include <stdio.h>
#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"
/*
* 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<N; i++)
{
@@ -72,12 +63,12 @@ int main(int argc, char *argv[])
}
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));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
@@ -86,12 +77,12 @@ int main(int argc, char *argv[])
hipLaunchKernelGGL(vector_square, 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++) {
if (C_h[i] != A_h[i] * A_h[i]) {
CHECK(hipErrorUnknown);
checkHipErrors(hipErrorUnknown);
}
}
printf ("PASSED!\n");
projects/hip-tests/samples/2_Cookbook/23_cmake_hiprtc/CMakeLists.txt
+2
Просмотреть файл
@@ -34,3 +34,5 @@ add_executable(test saxpy.cpp)
target_link_libraries(test hiprtc::hiprtc)
# Link with HIP
target_link_libraries(test hip::device)
target_include_directories(test PRIVATE ../../common)
projects/hip-tests/samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp
+16 -15
Просмотреть файл
@@ -22,6 +22,7 @@ THE SOFTWARE.
#include <hip/hiprtc.h>
#include <hip/hip_runtime.h>
#include <hip_helper.h>
#include <cassert>
#include <cstddef>
@@ -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;
}
projects/hip-tests/samples/2_Cookbook/3_shared_memory/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/3_shared_memory/Makefile
+2 -1
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp
+8 -8
Просмотреть файл
@@ -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);
projects/hip-tests/samples/2_Cookbook/4_shfl/CMakeLists.txt
+2
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile
+2 -1
Просмотреть файл
@@ -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)
projects/hip-tests/samples/2_Cookbook/4_shfl/shfl.cpp
+8 -8
Просмотреть файл
@@ -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);
projects/hip-tests/samples/2_Cookbook/5_2dshfl/2dshfl.cpp
+8 -7
Просмотреть файл
@@ -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);
projects/hip-tests/samples/2_Cookbook/5_2dshfl/CMakeLists.txt
+2
Просмотреть файл
@@ -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)

Некоторые файлы не были показаны из-за слишком большого количества измененных файлов Показать больше