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

EXSWCPHIPT-133 - Added support for offsetting to arrays to hipStreamValue tests (#2766)

Просмотр исходного кода
Этот коммит содержится в:
Finlay
2022-08-08 06:17:42 +01:00
коммит произвёл GitHub
родитель e47d977614
Коммит cc1f3702e7
4 изменённых файлов: 368 добавлений и 268 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
catch/unit/stream/CMakeLists.txt
+2 -1
Просмотреть файл
@@ -43,4 +43,5 @@ endif()
hip_add_exe_to_target(NAME StreamTest
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests)
TEST_TARGET_NAME build_tests
COMPILE_OPTIONS -std=c++17)
catch/unit/stream/hipStreamCreateWithPriority.cc
+2 -2
Просмотреть файл
@@ -509,8 +509,8 @@ bool validateStreamPrioritiesWithEvents() {
#define OP(x) \
free(src_h_##x); \
free(dst_h_##x); \
hipFree(src_d_##x); \
hipFree(dst_d_##x);
HIP_CHECK(hipFree(src_d_##x)); \
HIP_CHECK(hipFree(dst_d_##x));
OP(low)
OP(normal)
OP(high)
catch/unit/stream/hipStreamValue.cc
+358 -259
Просмотреть файл
@@ -17,276 +17,377 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <memory>
#include <type_traits>
constexpr unsigned int writeFlag = 0;
#define DEFINE_HIP_STREAM_VALUE(TYPE, BITS, ...) hipStream##TYPE##Value##BITS(__VA_ARGS__)
template <typename UIntT, typename... Args> auto waitFunc(Args... args) {
if constexpr (std::is_same<UIntT, uint32_t>::value) {
return hipStreamWaitValue32(args...);
} else {
return hipStreamWaitValue64(args...);
}
};
#define CHECK_HIP_STREAM_VALUE(TYPE, BITS, ...) \
HIP_CHECK(DEFINE_HIP_STREAM_VALUE(TYPE, BITS, __VA_ARGS__));
template <typename UIntT, typename... Args> auto writeFunc(Args... args) {
if constexpr (std::is_same<UIntT, uint32_t>::value) {
return hipStreamWriteValue32(args...);
} else {
return hipStreamWriteValue64(args...);
}
};
#define NEG_TEST_ERROR_CHECK(TYPE, BITS, errorCode, ...) \
HIP_CHECK_ERROR(DEFINE_HIP_STREAM_VALUE(TYPE, BITS, __VA_ARGS__), errorCode);
// Random predefined 32 and 64 bit values
using value32_t = std::integral_constant<uint32_t, 0x70F0F0FF>;
using value64_t = std::integral_constant<uint64_t, 0x7FFF0000FFFF0000>;
template <typename UIntT>
using testValue =
typename std::conditional<std::is_same<UIntT, uint32_t>::value, value32_t, value64_t>::type;
#if HT_AMD
// Random predefiend 32 and 64 bit values
constexpr uint32_t value32 = 0x70F0F0FF;
constexpr uint64_t value64 = 0x7FFF0000FFFF0000;
constexpr uint32_t DATA_INIT = 0x1234;
constexpr uint32_t DATA_UPDATE = 0X4321;
template <typename intT> struct TEST_WAIT {
using uintT = typename std::make_unsigned<intT>::type;
int compareOp;
uintT mask;
uintT waitValue;
intT signalValueFail;
intT signalValuePass;
template <typename UIntT> struct TEST_WAIT {
static_assert(std::is_same<UIntT, uint32_t>::value or std::is_same<UIntT, uint64_t>::value,
"only implemented for 32 bit and 64 bit unsigned integers");
unsigned int compareOp;
UIntT mask = ~static_cast<UIntT>(0);
UIntT waitValue;
UIntT signalValueFail;
UIntT signalValuePass;
TEST_WAIT(int compareOp, uintT waitValue, intT signalValueFail, intT signalValuePass)
TEST_WAIT(unsigned int compareOp, UIntT waitValue, UIntT signalValueFail, UIntT signalValuePass)
: compareOp{compareOp},
waitValue{waitValue},
signalValueFail{signalValueFail},
signalValuePass{signalValuePass} {
mask = static_cast<uintT>(0xFFFFFFFFFFFFFFFF);
}
signalValuePass{signalValuePass} {}
TEST_WAIT(int compareOp, uintT mask, uintT waitValue, intT signalValueFail, intT signalValuePass)
TEST_WAIT(unsigned int compareOp, UIntT mask, UIntT waitValue, UIntT signalValueFail,
UIntT signalValuePass)
: compareOp{compareOp},
mask{mask},
waitValue{waitValue},
signalValueFail{signalValueFail},
signalValuePass{signalValuePass} {}
};
typedef TEST_WAIT<int32_t> TEST_WAIT32;
typedef TEST_WAIT<int64_t> TEST_WAIT64;
using TEST_WAIT32 = TEST_WAIT<uint32_t>;
using TEST_WAIT64 = TEST_WAIT<uint64_t>;
bool streamWaitValueSupported() {
int device_num = 0;
HIP_CHECK(hipGetDeviceCount(&device_num));
int waitValueSupport;
for (int device_id = 0; device_id < device_num; ++device_id) {
HIP_CHECK(hipSetDevice(device_id));
waitValueSupport = 0;
HIP_CHECK(hipDeviceGetAttribute(&waitValueSupport, hipDeviceAttributeCanUseStreamWaitValue,
device_id));
int waitValueSupport = 0;
auto getAttributeError = hipDeviceGetAttribute(
&waitValueSupport, hipDeviceAttributeCanUseStreamWaitValue, device_id);
if (getAttributeError != hipSuccess) {
HipTest::HIP_SKIP_TEST("attribute not supported");
return false;
}
if (waitValueSupport == 1) return true;
}
return false;
}
// hipStreamWriteValue Tests
TEST_CASE("Unit_hipStreamValue_Write") {
int64_t* signalPtr;
// The different types of memory that can be used with hipStream[Wait|Write]
enum class PtrType { HostPtr, DevicePtr, DevicePtrToHost, Signal };
hipStream_t stream{nullptr};
HIP_CHECK(hipStreamCreate(&stream));
// Helper class to expose the pointer that is used with hipStream[Write|Wait]Value and also store a
// unique pointer with the deleter to simplify cleanup
// Also includes functions to update and get the value directly
template <PtrType type, typename UIntT, typename UniquePtrWithDeleter> class TestPtr {
// This stores the memory that must be deleted, as well as the deleter
UniquePtrWithDeleter ptrToDelete;
// Allocate Host Memory
auto hostPtr64 = std::unique_ptr<uint64_t>(new uint64_t(1));
auto hostPtr32 = std::unique_ptr<uint32_t>(new uint32_t(1));
public:
// The pointer that should be used with hipStream[Write|Wait]Value
UIntT* ptr;
// Register Host Memory
HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0));
HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0));
TestPtr(UIntT* ptr, UniquePtrWithDeleter ptrToDelete)
: ptrToDelete(std::move(ptrToDelete)), ptr(ptr) {}
// Register Signal Memory
HIP_CHECK(hipExtMallocWithFlags((void**)&signalPtr, 8, hipMallocSignalMemory));
// Initialise Data
*signalPtr = 0x0;
*hostPtr64 = 0x0;
*hostPtr32 = 0x0;
SECTION("Registered host memory hipStreamWriteValue32") {
INFO("Test writting to registered host pointer using hipStreamWriteValue32");
HIP_CHECK(hipStreamWriteValue32(stream, hostPtr32.get(), value32, writeFlag));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_ASSERT(*hostPtr32 == value32);
// directly retrieve the value from wherever it was allocated
UIntT getValue(size_t offset = 0) {
if constexpr (type == PtrType::Signal || type == PtrType::HostPtr ||
type == PtrType::DevicePtrToHost) {
return ptrToDelete.get()[offset];
} else {
static_assert(type == PtrType::DevicePtr, "Expected DevicePtr");
UIntT value;
HIP_CHECK(hipMemcpy(&value, ptr + offset, sizeof(UIntT), hipMemcpyDeviceToHost));
return value;
}
}
SECTION("Registered host memory hipStreamWriteValue64") {
INFO("Test writting to registered host pointer using hipStreamWriteValue32");
HIP_CHECK(hipStreamWriteValue64(stream, hostPtr64.get(), value64, writeFlag));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_ASSERT(*hostPtr64 == value64);
// directly set the value wherever it was allocated
void setValue(UIntT value, size_t offset = 0) {
if constexpr (type == PtrType::Signal || type == PtrType::DevicePtrToHost ||
type == PtrType::HostPtr) {
ptrToDelete.get()[offset] = value;
} else {
// hipMemcpy causes deadlock, so use hipStreamWriteValue
static_assert(type == PtrType::DevicePtr, "Expected DevicePtr");
hipStream_t stream;
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
HIP_CHECK(writeFunc<UIntT>(stream, ptr + offset, value, writeFlag));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
}
}
};
// Test writting device pointer
void* devicePtr64;
void* devicePtr32;
HIP_CHECK(hipHostGetDevicePointer((void**)&devicePtr64, hostPtr64.get(), 0));
HIP_CHECK(hipHostGetDevicePointer((void**)&devicePtr32, hostPtr32.get(), 0));
// Reset values
*hostPtr64 = 0x0;
*hostPtr32 = 0x0;
// required for the static assert
template <PtrType> inline constexpr bool AMD_ACTIVE = HT_AMD == 1;
SECTION("Device Memory hipStreamWriteValue32") {
INFO("Test writting to device pointer using hipStreamWriteValue32");
HIP_CHECK(hipStreamWriteValue32(stream, devicePtr32, value32, writeFlag));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_ASSERT(*hostPtr32 == value32);
template <PtrType type, typename UIntT> auto allocMem() {
constexpr std::size_t arraySize = 1024;
if constexpr (type == PtrType::Signal) {
static_assert(std::is_same<UIntT, uint64_t>::value,
"signal memory should only be used with 64bit memory");
// Allocate Signal Memory
uint64_t* signalPtr{};
static_assert(AMD_ACTIVE<type>,
"nvidia backend compiler doesn't like hipExtMallocWithFlags, even in this "
"constexpr branch");
#if HT_AMD
// 8 is the only acceptable size
HIP_CHECK(
hipExtMallocWithFlags(reinterpret_cast<void**>(&signalPtr), 8, hipMallocSignalMemory));
#endif
// Init Memory
*signalPtr = 0;
auto freeStuff = [](uint64_t* sPtr) { HIP_CHECK(hipFree(sPtr)); };
return TestPtr<type, UIntT, std::unique_ptr<uint64_t, decltype(freeStuff)>>{
signalPtr, std::unique_ptr<uint64_t, decltype(freeStuff)>(signalPtr, freeStuff)};
} else if constexpr (type == PtrType::DevicePtrToHost) {
auto hostPtr = new UIntT[arraySize];
// Register Host Memory
HIP_CHECK(hipHostRegister(hostPtr, sizeof(UIntT) * arraySize, 0));
// Init memory
std::fill(hostPtr, hostPtr + arraySize, 0);
UIntT* devicePtr;
// Test writing device pointer
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&devicePtr), hostPtr, 0));
auto freeStuff = [](UIntT* ptr) {
HIP_CHECK(hipHostUnregister(ptr));
delete[] ptr;
};
return TestPtr<type, UIntT, std::unique_ptr<UIntT[], decltype(freeStuff)>>{
devicePtr, std::unique_ptr<UIntT[], decltype(freeStuff)>(hostPtr, freeStuff)};
} else if constexpr (type == PtrType::HostPtr) {
auto hostPtr = new UIntT[arraySize];
// Register Host Memory
HIP_CHECK(hipHostRegister(hostPtr, sizeof(UIntT) * arraySize, 0));
// Init memory
std::fill(hostPtr, hostPtr + arraySize, 0);
auto freeStuff = [](UIntT* ptr) {
HIP_CHECK(hipHostUnregister(ptr));
delete[] ptr;
};
return TestPtr<type, UIntT, std::unique_ptr<UIntT[], decltype(freeStuff)>>{
hostPtr, std::unique_ptr<UIntT[], decltype(freeStuff)>(hostPtr, freeStuff)};
} else {
static_assert(type == PtrType::DevicePtr, "Expected DevicePtr");
UIntT* devicePtr;
HIP_CHECK(hipMalloc(&devicePtr, sizeof(UIntT) * arraySize));
HIP_CHECK(hipMemset(devicePtr, 0, sizeof(UIntT) * arraySize));
auto freeStuff = [](UIntT* ptr) { HIP_CHECK(hipFree(ptr)); };
return TestPtr<type, UIntT, std::unique_ptr<UIntT, decltype(freeStuff)>>{
devicePtr, std::unique_ptr<UIntT, decltype(freeStuff)>(devicePtr, freeStuff)};
}
SECTION("Device Memory hipStreamWriteValue64") {
INFO("Test writting to device pointer using hipStreamWriteValue64");
HIP_CHECK(hipStreamWriteValue64(stream, devicePtr64, value64, writeFlag));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_ASSERT(*hostPtr64 == value64);
}
// Test Writing to Signal Memory
SECTION("Signal Memory hipStreamWriteValue64") {
INFO("Test writting to signal memory using hipStreamWriteValue64");
HIP_CHECK(hipStreamWriteValue64(stream, signalPtr, value64, writeFlag));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_ASSERT(*signalPtr == value64);
}
// Cleanup
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipHostUnregister(hostPtr64.get()));
HIP_CHECK(hipHostUnregister(hostPtr32.get()));
HIP_CHECK(hipFree(signalPtr));
}
// hipStreamWaitValue Tests
template <bool isBlocking, typename intT, typename TEST_T>
void initData(intT* dataPtr, int64_t* signalPtr, TEST_T tc, std::vector<hipEvent_t>& events) {
// Initialize memory to be waited on
*signalPtr = isBlocking ? tc.signalValueFail : tc.signalValuePass;
// Initialize host pointers
dataPtr[0] = DATA_INIT;
dataPtr[1] = DATA_INIT;
hipEvent_t firstWriteEvent{nullptr};
hipEvent_t secondWriteEvent{nullptr};
HIP_CHECK(hipEventCreate(&firstWriteEvent));
HIP_CHECK(hipEventCreate(&secondWriteEvent));
events.push_back(firstWriteEvent);
events.push_back(secondWriteEvent);
}
template <bool isBlocking, typename intT, typename TEST_T>
void syncAndCheckData(hipStream_t stream, intT* dataPtr, int64_t* signalPtr, TEST_T tc,
std::vector<hipEvent_t>& events) {
// Ensure first part of host memory is updated
HIP_CHECK(hipStreamWaitEvent(stream, events[0], 0));
HIP_ASSERT(dataPtr[0] == DATA_UPDATE);
if (isBlocking) {
// Ensure second part of host memory isn't updated yet
HIP_ASSERT(hipEventQuery(events[1]) == hipErrorNotReady);
HIP_ASSERT(dataPtr[1] == DATA_INIT);
// Update value to release stream
*signalPtr = tc.signalValuePass;
// allows the creation of a list of offsets while avoiding it for signal memory
template <PtrType type> constexpr auto get_offsets() {
if constexpr (type == PtrType::Signal) {
return std::array<size_t, 1>{0};
} else {
return std::array<size_t, 6>{0, 1, 2, 3, 31, 1023};
}
HIP_CHECK(hipStreamSynchronize(stream));
HIP_ASSERT(hipEventQuery(events[1]) == hipSuccess);
// Finally ensure that second part of host memory is updated
HIP_ASSERT(dataPtr[1] == DATA_UPDATE);
}
template <typename intT> void cleanup(hipStream_t& stream, intT* dataPtr, int64_t* signalPtr) {
// Cleanup
HIP_CHECK(hipFree(signalPtr));
HIP_CHECK(hipHostUnregister(dataPtr));
HIP_CHECK(hipStreamDestroy(stream));
}
template <typename UIntT, PtrType ptrTypeValue> struct TestParams {
using UIntType = UIntT;
constexpr static PtrType ptrType = ptrTypeValue;
};
template <typename intT, bool isBlocking, typename TEST_T> void testWait(TEST_T tc) {
#if HT_AMD
TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams<uint32_t, PtrType::HostPtr>),
(TestParams<uint32_t, PtrType::DevicePtr>),
(TestParams<uint32_t, PtrType::DevicePtrToHost>),
(TestParams<uint64_t, PtrType::HostPtr>),
(TestParams<uint64_t, PtrType::DevicePtr>),
(TestParams<uint64_t, PtrType::DevicePtrToHost>),
(TestParams<uint64_t, PtrType::Signal>)) {
#else
TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams<uint32_t, PtrType::HostPtr>),
(TestParams<uint32_t, PtrType::DevicePtr>),
(TestParams<uint32_t, PtrType::DevicePtrToHost>),
(TestParams<uint64_t, PtrType::HostPtr>),
(TestParams<uint64_t, PtrType::DevicePtr>),
(TestParams<uint64_t, PtrType::DevicePtrToHost>)) {
#endif
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-126");
return;
#endif
if (!streamWaitValueSupported()) {
UNSCOPED_INFO(" hipStreamWaitValue: not supported on this device , skipping ...");
HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
return;
}
// Initialize stream
using UIntT = typename TestType::UIntType;
constexpr auto ptrType = TestType::ptrType;
constexpr auto writeValue = testValue<UIntT>::value;
hipStream_t stream{nullptr};
HIP_CHECK(hipStreamCreate(&stream));
const auto offsets = get_offsets<ptrType>();
const auto offset = GENERATE_COPY(from_range(std::begin(offsets), std::end(offsets)));
CAPTURE(offset);
// Allocate Host Memory
std::unique_ptr<intT> dataPtr(new intT(2));
auto ptr = allocMem<ptrType, UIntT>();
UIntT* target = ptr.ptr + offset;
HIP_CHECK(writeFunc<UIntT>(stream, target, writeValue, writeFlag));
HIP_CHECK(hipStreamSynchronize(stream));
REQUIRE(ptr.getValue(offset) == writeValue);
// Register Host Memory
HIP_CHECK(hipHostRegister(&(dataPtr.get()[0]), sizeof(intT), 0));
HIP_CHECK(hipHostRegister(&(dataPtr.get()[1]), sizeof(intT), 0));
// Cleanup
HIP_CHECK(hipStreamDestroy(stream));
}
// Allocate Signal Memory
int64_t* signalPtr;
HIP_CHECK(hipExtMallocWithFlags((void**)&signalPtr, 8, hipMallocSignalMemory));
template <bool isBlocking, typename UIntT, typename TestPtr>
void syncAndCheckData(hipStream_t stream, UIntT* dataPtr, TestPtr signalPtr, size_t offset,
TEST_WAIT<UIntT> tc, std::array<hipEvent_t, 2>& events) {
// Ensure first part of host memory is updated
HIP_CHECK(hipEventSynchronize(events[0]));
REQUIRE(dataPtr[0] == DATA_UPDATE);
std::vector<hipEvent_t> events;
initData<isBlocking>(dataPtr.get(), signalPtr, tc, events);
if (std::is_same<intT, int32_t>::value) {
CHECK_HIP_STREAM_VALUE(Write, 32, stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag)
HIP_CHECK(hipEventRecord(events[0], stream));
if (static_cast<uint32_t>(tc.mask) != 0xFFFFFFFF) {
CHECK_HIP_STREAM_VALUE(Wait, 32, stream, signalPtr, static_cast<uint32_t>(tc.waitValue),
tc.compareOp, static_cast<uint32_t>(tc.mask));
} else {
CHECK_HIP_STREAM_VALUE(Wait, 32, stream, signalPtr, tc.waitValue, tc.compareOp);
}
CHECK_HIP_STREAM_VALUE(Write, 32, stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag)
} else {
CHECK_HIP_STREAM_VALUE(Write, 64, stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag)
HIP_CHECK(hipEventRecord(events[0], stream));
if (tc.mask != 0xFFFFFFFFFFFFFFFF) {
CHECK_HIP_STREAM_VALUE(Wait, 64, stream, signalPtr, tc.waitValue, tc.compareOp, tc.mask);
} else {
CHECK_HIP_STREAM_VALUE(Wait, 64, stream, signalPtr, tc.waitValue, tc.compareOp);
}
CHECK_HIP_STREAM_VALUE(Write, 64, stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag)
if constexpr (isBlocking) {
// Ensure second part of host memory isn't updated yet
HIP_CHECK_ERROR(hipEventQuery(events[1]), hipErrorNotReady);
REQUIRE(dataPtr[1] == DATA_INIT);
// Update value to release stream
signalPtr.setValue(tc.signalValuePass, offset);
}
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipEventQuery(events[1]));
// Finally ensure that second part of host memory is updated
REQUIRE(dataPtr[1] == DATA_UPDATE);
}
template <typename TestType, bool isBlocking>
void testWait(TEST_WAIT<typename TestType::UIntType> tc) {
if (!streamWaitValueSupported()) {
HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
return;
}
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-128");
return;
#endif
using UIntT = typename TestType::UIntType;
constexpr auto ptrType = TestType::ptrType;
constexpr UIntT defaultMask = ~static_cast<UIntT>(0);
// Initialize stream
hipStream_t stream{};
HIP_CHECK(hipStreamCreate(&stream));
// Allocate Host Memory
auto dataPtr = std::make_unique<UIntT[]>(2);
// Register Host Memory
HIP_CHECK(hipHostRegister(dataPtr.get(), sizeof(UIntT), 0));
HIP_CHECK(hipHostRegister(dataPtr.get() + 1, sizeof(UIntT), 0));
std::fill(dataPtr.get(), dataPtr.get() + 2, DATA_INIT);
std::array<hipEvent_t, 2> events;
HIP_CHECK(hipEventCreate(&events[0]));
HIP_CHECK(hipEventCreate(&events[1]));
const auto offsets = get_offsets<ptrType>();
const auto offset = GENERATE_COPY(from_range(std::begin(offsets), std::end(offsets)));
auto waitPtr = allocMem<ptrType, UIntT>();
UIntT* const target = waitPtr.ptr + offset;
waitPtr.setValue(isBlocking ? tc.signalValueFail : tc.signalValuePass, offset);
HIP_CHECK(writeFunc<UIntT>(stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag));
HIP_CHECK(hipEventRecord(events[0], stream));
if (tc.mask != defaultMask) {
HIP_CHECK(waitFunc<UIntT>(stream, target, tc.waitValue, tc.compareOp, tc.mask));
} else {
HIP_CHECK(waitFunc<UIntT>(stream, target, tc.waitValue, tc.compareOp));
}
HIP_CHECK(writeFunc<UIntT>(stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag));
HIP_CHECK(hipEventRecord(events[1], stream));
syncAndCheckData<isBlocking>(stream, dataPtr.get(), signalPtr, tc, events);
cleanup(stream, dataPtr.get(), signalPtr);
}
#undef CHECK_HIP_STREAM_VALUE
syncAndCheckData<isBlocking>(stream, dataPtr.get(), std::move(waitPtr), offset, tc, events);
// Cleanup
HIP_CHECK(hipEventDestroy(events[0]));
HIP_CHECK(hipEventDestroy(events[1]));
HIP_CHECK(hipHostUnregister(dataPtr.get()));
HIP_CHECK(hipHostUnregister(dataPtr.get() + 1));
HIP_CHECK(hipStreamDestroy(stream));
}
// TEMPLATE_TEST_CASE wasn't working within a macro, so sections were used instead
#define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32(suffix, test_t) \
TEST_CASE("Unit_hipStreamValue_Wait32_Blocking_" + std::string(suffix)) { \
testWait<int32_t, true>(test_t); \
SECTION("HostPtr") { testWait<TestParams<uint32_t, PtrType::HostPtr>, true>(test_t); } \
SECTION("DevicePtr") { testWait<TestParams<uint32_t, PtrType::DevicePtr>, true>(test_t); } \
SECTION("DevicePtrToHost") { \
testWait<TestParams<uint32_t, PtrType::DevicePtrToHost>, true>(test_t); \
} \
} \
TEST_CASE("Unit_hipStreamValue_Wait32_NonBlocking_" + std::string(suffix)) { \
testWait<int32_t, false>(test_t); \
SECTION("HostPtr") { testWait<TestParams<uint32_t, PtrType::HostPtr>, false>(test_t); } \
SECTION("DevicePtr") { testWait<TestParams<uint32_t, PtrType::DevicePtr>, false>(test_t); } \
SECTION("DevicePtrToHost") { \
testWait<TestParams<uint32_t, PtrType::DevicePtrToHost>, false>(test_t); \
} \
}
// Using Mask
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte_1",
TEST_WAIT64( // mask will ignore few MSB bits
hipStreamWaitValueGte, 0x0000FFFFFFFFFFFF,
0x000000007FFF0001, 0x7FFF00007FFF0000,
0x000000007FFF0001))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte_2",
TEST_WAIT64(hipStreamWaitValueGte, 0xF, 0x4, 0x3, 0x6))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte",
TEST_WAIT32(hipStreamWaitValueGte, 0xF, 0x4, 0x3, 0x6))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Eq_1",
TEST_WAIT64( // mask will ignore few MSB bits
hipStreamWaitValueEq, 0x0000FFFFFFFFFFFF,
0x000000000FFF0001, 0x7FFF00000FFF0000,
0x7F0000000FFF0001))
TEST_WAIT32( // mask will ignore few MSB bits
hipStreamWaitValueEq, 0x0000FFFF, 0x00000001,
0x0FFF0000, 0x0FFF0001))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Eq_2",
TEST_WAIT64(hipStreamWaitValueEq, 0xFF, 0x11, 0x25, 0x11))
TEST_WAIT32(hipStreamWaitValueEq, 0xFF, 0x11, 0x25, 0x11))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_And",
TEST_WAIT64( // mask will discard bits 8 to 11
TEST_WAIT32( // mask will discard bits 8 to 11
hipStreamWaitValueAnd, 0xFF, 0xF4A, 0xF35, 0X02))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Nor_1",
TEST_WAIT64( // mask is set to ignore the sign bit.
hipStreamWaitValueNor, 0x7FFFFFFFFFFFFFFF,
0x7FFFFFFFFFFFF247, 0x7FFFFFFFFFFFFdbd,
0x7FFFFFFFFFFFFdb5))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Nor_2",
TEST_WAIT64( // mask is set to apply NOR for bits 0 to 3.
hipStreamWaitValueNor, 0xF, 0x7E, 0x7D, 0x76))
// Not Using Mask
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_Eq",
@@ -299,19 +400,47 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_And",
TEST_WAIT32(hipStreamWaitValueAnd, 0x70F0F0F0, 0x0F0F0F0F,
0X1F0F0F0F))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_Nor",
TEST_WAIT32(hipStreamWaitValueNor, 0x7AAAAAAA,
static_cast<int32_t>(0x85555555),
static_cast<int32_t>(0x9AAAAAAA)))
TEST_WAIT32(hipStreamWaitValueNor, 0x7AAAAAAA, 0x85555555,
0x9AAAAAAA))
#undef DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32
#if HT_AMD
// TEMPLATE_TEST_CASE wasn't working within a macro, so sections were used instead
#define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64(suffix, test_t) \
TEST_CASE("Unit_hipStreamValue_Wait64_Blocking_" + std::string(suffix)) { \
testWait<int64_t, true>(test_t); \
SECTION("HostPtr") { testWait<TestParams<uint64_t, PtrType::HostPtr>, true>(test_t); } \
SECTION("DevicePtr") { testWait<TestParams<uint64_t, PtrType::DevicePtr>, true>(test_t); } \
SECTION("DevicePtrToHost") { \
testWait<TestParams<uint64_t, PtrType::DevicePtrToHost>, true>(test_t); \
} \
SECTION("Signal") { testWait<TestParams<uint64_t, PtrType::Signal>, true>(test_t); } \
} \
TEST_CASE("Unit_hipStreamValue_Wait64_NonBlocking_" + std::string(suffix)) { \
testWait<int64_t, false>(test_t); \
SECTION("HostPtr") { testWait<TestParams<uint64_t, PtrType::HostPtr>, false>(test_t); } \
SECTION("DevicePtr") { testWait<TestParams<uint64_t, PtrType::DevicePtr>, false>(test_t); } \
SECTION("DevicePtrToHost") { \
testWait<TestParams<uint64_t, PtrType::DevicePtrToHost>, false>(test_t); \
} \
SECTION("Signal") { testWait<TestParams<uint64_t, PtrType::Signal>, false>(test_t); } \
}
#else
#define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64(suffix, test_t) \
TEST_CASE("Unit_hipStreamValue_Wait64_Blocking_" + std::string(suffix)) { \
SECTION("HostPtr") { testWait<TestParams<uint64_t, PtrType::HostPtr>, true>(test_t); } \
SECTION("DevicePtr") { testWait<TestParams<uint64_t, PtrType::DevicePtr>, true>(test_t); } \
SECTION("DevicePtrToHost") { \
testWait<TestParams<uint64_t, PtrType::DevicePtrToHost>, true>(test_t); \
} \
} \
TEST_CASE("Unit_hipStreamValue_Wait64_NonBlocking_" + std::string(suffix)) { \
SECTION("HostPtr") { testWait<TestParams<uint64_t, PtrType::HostPtr>, false>(test_t); } \
SECTION("DevicePtr") { testWait<TestParams<uint64_t, PtrType::DevicePtr>, false>(test_t); } \
SECTION("DevicePtrToHost") { \
testWait<TestParams<uint64_t, PtrType::DevicePtrToHost>, false>(test_t); \
} \
}
#endif
// Using Mask
@@ -332,14 +461,6 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Eq_2",
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_And",
TEST_WAIT64( // mask will discard bits 8 to 11
hipStreamWaitValueAnd, 0xFF, 0xF4A, 0xF35, 0X02))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Nor_1",
TEST_WAIT64( // mask is set to ignore the sign bit.
hipStreamWaitValueNor, 0x7FFFFFFFFFFFFFFF,
0x7FFFFFFFFFFFF247, 0x7FFFFFFFFFFFFdbd,
0x7FFFFFFFFFFFFdb5))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Nor_2",
TEST_WAIT64( // mask is set to apply NOR for bits 0 to 3.
hipStreamWaitValueNor, 0xF, 0x7E, 0x7D, 0x76))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_Gte",
TEST_WAIT64(hipStreamWaitValueGte, 0x7FFFFFFFFFFF0001,
@@ -352,94 +473,72 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_And",
0x0F0F0F0F0F0F0F0F, 0X1F0F0F0F0F0F0F0F))
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_Nor",
TEST_WAIT64(hipStreamWaitValueNor, 0x4724724747247247,
static_cast<int64_t>(0xbddbddbdbddbddbd),
static_cast<int64_t>(0xbddbddbdbddbddb3)))
0xbddbddbdbddbddbd, 0xbddbddbdbddbddb3))
#undef DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64
#endif
// Negative Tests
TEST_CASE("Unit_hipStreamValue_Negative_InvalidMemory") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-96");
return;
#endif
if (!streamWaitValueSupported()) {
HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
return;
}
hipStream_t stream{nullptr};
HIP_CHECK(hipStreamCreate(&stream));
REQUIRE(stream != nullptr);
// Allocate Host Memory
auto hostPtr32 = std::unique_ptr<uint32_t>(new uint32_t(1));
auto hostPtr64 = std::unique_ptr<uint64_t>(new uint64_t(1));
// Register Host Memory
HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0));
HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0));
// Set dummy data
*hostPtr64 = 0x0;
*hostPtr32 = 0x0;
auto compareOp = hipStreamWaitValueGte;
const auto compareOp = hipStreamWaitValueGte;
const auto expectedError = hipErrorInvalidValue;
// Memory pointer negative tests
INFO("Testing Invalid Memory Pointer for hipStreamWriteValue32");
NEG_TEST_ERROR_CHECK(Write, 32, hipErrorNotSupported, stream, nullptr, 0, writeFlag)
INFO("Testing Invalid Memory Pointer for hipStreamWriteValue64");
NEG_TEST_ERROR_CHECK(Write, 64, hipErrorNotSupported, stream, nullptr, 0, writeFlag)
INFO("Testing Invalid Memory Pointer for hipStreamWaitValue32");
NEG_TEST_ERROR_CHECK(Wait, 32, hipErrorNotSupported, stream, nullptr, 0, compareOp)
INFO("Testing Invalid Memory Pointer for hipStreamWaitValue64");
NEG_TEST_ERROR_CHECK(Wait, 64, hipErrorNotSupported, stream, nullptr, 0, compareOp)
SECTION("Invalid Memory Pointer for hipStreamWriteValue32") {
HIP_CHECK_ERROR(hipStreamWriteValue32(stream, nullptr, 0, writeFlag), expectedError);
}
SECTION("Invalid Memory Pointer for hipStreamWriteValue64") {
HIP_CHECK_ERROR(hipStreamWriteValue64(stream, nullptr, 0, writeFlag), expectedError);
}
SECTION("Invalid Memory Pointer for hipStreamWaitValue32") {
HIP_CHECK_ERROR(hipStreamWaitValue32(stream, nullptr, 0, compareOp), expectedError);
}
SECTION("Invalid Memory Pointer for hipStreamWaitValue32") {
HIP_CHECK_ERROR(hipStreamWaitValue64(stream, nullptr, 0, compareOp), expectedError);
}
// Cleanup
HIP_CHECK(hipHostUnregister(hostPtr32.get()));
HIP_CHECK(hipHostUnregister(hostPtr64.get()));
HIP_CHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag") {
TEMPLATE_TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag", "", uint32_t, uint64_t) {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-96");
return;
#endif
if (!streamWaitValueSupported()) {
HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
return;
}
hipStream_t stream{nullptr};
HIP_CHECK(hipStreamCreate(&stream));
REQUIRE(stream != nullptr);
// Allocate Host Memory
auto hostPtr32 = std::unique_ptr<uint32_t>(new uint32_t(1));
auto hostPtr64 = std::unique_ptr<uint64_t>(new uint64_t(1));
auto hostPtr = std::make_unique<TestType>();
// Register Host Memory
HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0));
HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0));
HIP_CHECK(hipHostRegister(hostPtr.get(), sizeof(TestType), 0));
// Set dummy data
*hostPtr64 = 0x0;
*hostPtr32 = 0x0;
*hostPtr = 0x0;
/* EXSWCPHIPT-96 */
INFO("Testing Invalid flag for hipStreamWaitValue32");
NEG_TEST_ERROR_CHECK(Wait, 32, hipErrorNotSupported, stream, hostPtr32.get(), 0, -1)
INFO("Testing Invalid flag for hipStreamWaitValue64");
NEG_TEST_ERROR_CHECK(Wait, 64, hipErrorNotSupported, stream, hostPtr64.get(), 0, -1)
HIP_CHECK_ERROR(waitFunc<TestType>(stream, hostPtr.get(), 0, -1), hipErrorInvalidValue);
// Cleanup
HIP_CHECK(hipHostUnregister(hostPtr32.get()));
HIP_CHECK(hipHostUnregister(hostPtr64.get()));
HIP_CHECK(hipHostUnregister(hostPtr.get()));
HIP_CHECK(hipStreamDestroy(stream));
}
#undef NEG_TEST_ERROR_CHECK
catch/unit/stream/hipStreamWithCUMask.cc
+6 -6
Просмотреть файл
@@ -149,8 +149,8 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_ValidateCallbackFunc") {
HIP_CHECK(hipGetDeviceProperties(&props, 0));
createDefaultCUMask(&defaultCUMask, props.multiProcessorCount);
hipExtStreamCreateWithCUMask(&mystream, defaultCUMask.size(),
defaultCUMask.data());
HIP_CHECK(hipExtStreamCreateWithCUMask(&mystream, defaultCUMask.size(),
defaultCUMask.data()));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice,
mystream));
const unsigned blocks = GRIDSIZE;
@@ -244,7 +244,7 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") {
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, streams[0], dA[0], dC[0], N);
hipDeviceSynchronize();
HIP_CHECK(hipDeviceSynchronize());
auto single_end = std::chrono::steady_clock::now();
std::chrono::duration<double> single_kernel_time = single_end - single_start;
@@ -264,7 +264,7 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") {
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, streams[np], dA[np], dC[np], N);
}
hipDeviceSynchronize();
HIP_CHECK(hipDeviceSynchronize());
auto all_end = std::chrono::steady_clock::now();
std::chrono::duration<double> all_kernel_time = all_end - all_start;
@@ -288,8 +288,8 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") {
delete [] hA;
delete [] hC;
for (int np = 0; np < KNumPartition; np++) {
hipFree(dC[np]);
hipFree(dA[np]);
HIP_CHECK(hipFree(dC[np]));
HIP_CHECK(hipFree(dA[np]));
HIP_CHECK(hipStreamDestroy(streams[np]));
}
}