From 4cf7a5799ac2f4f29e7eae68a08db88185606164 Mon Sep 17 00:00:00 2001 From: Ravi C Akkenapally Date: Mon, 14 Dec 2020 15:28:01 -0800 Subject: [PATCH] SWDEV-179105 - Stream Operations: Add swqupport for Wait and Write Change-Id: I9d6d0665d12b62fe705ce1569a0e8264a4f23ab7 [ROCm/clr commit: fd0f49503fdcf0e836c4d898d0d6e66c0d45a159] --- .../include/hip/amd_detail/hip_runtime_api.h | 119 ++++++++ .../clr/hipamd/include/hip/hip_runtime_api.h | 5 + projects/clr/hipamd/rocclr/CMakeLists.txt | 1 + projects/clr/hipamd/rocclr/hip_device.cpp | 3 + .../clr/hipamd/rocclr/hip_device_runtime.cpp | 3 + projects/clr/hipamd/rocclr/hip_hcc.map.in | 4 + projects/clr/hipamd/rocclr/hip_memory.cpp | 14 +- projects/clr/hipamd/rocclr/hip_stream_ops.cpp | 129 ++++++++ .../streamOperations/hipstream_operations.cpp | 279 ++++++++++++++++++ 9 files changed, 555 insertions(+), 2 deletions(-) create mode 100644 projects/clr/hipamd/rocclr/hip_stream_ops.cpp create mode 100644 projects/clr/hipamd/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h index 6301251a98..6b04b7a57a 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h @@ -176,6 +176,7 @@ enum hipLimit_t { #define hipDeviceMallocDefault 0x0 #define hipDeviceMallocFinegrained 0x1 ///< Memory is allocated in fine grained region of device. +#define hipMallocSignalMemory 0x2 ///< Memory represents a HSA signal. //! Flags that can be used with hipHostRegister #define hipHostRegisterDefault 0x0 ///< Memory is Mapped and Portable @@ -216,6 +217,12 @@ enum hipLimit_t { // Flags that can be used with hipExtLaunch Set of APIs #define hipExtAnyOrderLaunch 0x01 ///< AnyOrderLaunch of kernels +// Flags to be used with hipStreamWaitValue32 and hipStreamWaitValue64 +#define hipStreamWaitValueGte 0x0 +#define hipStreamWaitValueEq 0x1 +#define hipStreamWaitValueAnd 0x2 +#define hipStreamWaitValueNor 0x3 + /* * @brief HIP Memory Advise values * @enum @@ -1221,6 +1228,118 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback */ +/** + *------------------------------------------------------------------------------------------------- + *------------------------------------------------------------------------------------------------- + * @defgroup Stream Memory Operations + * @{ + * This section describes Stream Memory Wait and Write functions of HIP runtime API. + */ + + +/** + * @brief Enqueues a wait command to the stream. + * + * @param [in] stream - Stream identifier + * @param [in] ptr - Pointer to memory object allocated using 'hipMallocSignalMemory' flag. + * @param [in] value - Value to be used in compare operation + * @param [in] mask - Mask to be applied on value at memory before it is compared with value + * @param [in] flags - Defines the compare operation, supported values are hipStreamWaitValueGte + * hipStreamWaitValueEq, hipStreamWaitValueAnd and hipStreamWaitValueNor. + * + * @returns #hipSuccess, #hipErrorInvalidValue + * + * Enqueues a wait command to the stream, all operations enqueued on this stream after this, will + * not execute until the defined wait condition is true. + * + * hipStreamWaitValueGte: waits until *ptr&mask >= value + * hipStreamWaitValueEq : waits until *ptr&mask == value + * hipStreamWaitValueAnd: waits until ((*ptr&mask) & value) != 0 + * hipStreamWaitValueNor: waits until ~((*ptr&mask) | (value&mask)) != 0 + * + * @note when using 'hipStreamWaitValueNor', mask is applied on both 'value' and '*ptr'. + * + * @note Support for hipStreamWaitValue32 can be queried using 'hipDeviceGetAttribute()' and + * 'hipDeviceAttributeCanUseStreamWaitValue' flag. + * + * @see hipExtMallocWithFlags, hipFree, hipStreamWaitValue64, hipStreamWriteValue64, + * hipStreamWriteValue32, hipDeviceGetAttribute + */ + +hipError_t hipStreamWaitValue32(hipStream_t stream, void* ptr, int32_t value, uint32_t mask, unsigned int flags); + +/** + * @brief Enqueues a wait command to the stream. + * + * @param [in] stream - Stream identifier + * @param [in] ptr - Pointer to memory object allocated using 'hipMallocSignalMemory' flag. + * @param [in] value - Value to be used in compare operation + * @param [in] mask - Mask to be applied on value at memory before it is compared with value. + * @param [in] flags - Defines the compare operation, supported values are hipStreamWaitValueGte + * hipStreamWaitValueEq, hipStreamWaitValueAnd and hipStreamWaitValueNor. + * + * @returns #hipSuccess, #hipErrorInvalidValue + * + * Enqueues a wait command to the stream, all operations enqueued on this stream after this, will + * not execute until the defined wait condition is true. + * + * hipStreamWaitValueGte: waits until *ptr&mask >= value + * hipStreamWaitValueEq : waits until *ptr&mask == value + * hipStreamWaitValueAnd: waits until ((*ptr&mask) & value) != 0 + * hipStreamWaitValueNor: waits until ~((*ptr&mask) | (value&mask)) != 0 + * + * @note when using 'hipStreamWaitValueNor', mask is applied on both 'value' and '*ptr'. + * + * @note Support for hipStreamWaitValue64 can be queried using 'hipDeviceGetAttribute()' and + * 'hipDeviceAttributeCanUseStreamWaitValue' flag. + * + * @see hipExtMallocWithFlags, hipFree, hipStreamWaitValue32, hipStreamWriteValue64, + * hipStreamWriteValue32, hipDeviceGetAttribute + */ + +hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, int64_t value, uint64_t mask, unsigned int flags); + +/** + * @brief Enqueues a write command to the stream. + * + * @param [in] stream - Stream identifier + * @param [in] ptr - Pointer to a GPU accessible memory object. + * @param [in] value - Value to be written + * + * @returns #hipSuccess, #hipErrorInvalidValue + * + * Enqueues a write command to the stream, write operation is performed after all earlier commands + * on this stream have completed the execution. + * + * @see hipExtMallocWithFlags, hipFree, hipStreamWriteValue32, hipStreamWaitValue32, + * hipStreamWaitValue64 + */ +hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, int32_t value); + +/** + * @brief Enqueues a write command to the stream. + * + * @param [in] stream - Stream identifier + * @param [in] ptr - Pointer to a GPU accessible memory object. + * @param [in] value - Value to be written + * + * @returns #hipSuccess, #hipErrorInvalidValue + * + * Enqueues a write command to the stream, write operation is performed after all earlier commands + * on this stream have completed the execution. + * + * @see hipExtMallocWithFlags, hipFree, hipStreamWriteValue32, hipStreamWaitValue32, + * hipStreamWaitValue64 + */ +hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, int64_t value); + + +// end doxygen Stream Memory Operations +/** + * @} + */ + + /** *------------------------------------------------------------------------------------------------- *------------------------------------------------------------------------------------------------- diff --git a/projects/clr/hipamd/include/hip/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hip_runtime_api.h index 6a281b43c4..792b794afa 100644 --- a/projects/clr/hipamd/include/hip/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hip_runtime_api.h @@ -147,6 +147,8 @@ typedef struct hipDeviceProp_t { int pageableMemoryAccess; ///< Device supports coherently accessing pageable memory ///< without calling hipHostRegister on it int pageableMemoryAccessUsesHostPageTables; ///< Device accesses pageable memory via the host's page tables + int canUseStreamWaitValue; ///< '1' if Device supports hipStreamWaitValue32() and + ///< hipStreamWaitValue64() , '0' otherwise. } hipDeviceProp_t; @@ -366,6 +368,9 @@ typedef enum hipDeviceAttribute_t { /// without calling hipHostRegister on it hipDeviceAttributePageableMemoryAccessUsesHostPageTables, ///< Device accesses pageable memory via /// the host's page tables + hipDeviceAttributeCanUseStreamWaitValue ///< '1' if Device supports hipStreamWaitValue32() and + ///< hipStreamWaitValue64() , '0' otherwise. + } hipDeviceAttribute_t; enum hipComputeMode { diff --git a/projects/clr/hipamd/rocclr/CMakeLists.txt b/projects/clr/hipamd/rocclr/CMakeLists.txt index fefb1fa3fd..28177b7c75 100755 --- a/projects/clr/hipamd/rocclr/CMakeLists.txt +++ b/projects/clr/hipamd/rocclr/CMakeLists.txt @@ -133,6 +133,7 @@ add_library(hip64 OBJECT hip_platform.cpp hip_profile.cpp hip_stream.cpp + hip_stream_ops.cpp hip_surface.cpp hip_texture.cpp hip_activity.cpp diff --git a/projects/clr/hipamd/rocclr/hip_device.cpp b/projects/clr/hipamd/rocclr/hip_device.cpp index 3488bc21bd..8725edcf24 100644 --- a/projects/clr/hipamd/rocclr/hip_device.cpp +++ b/projects/clr/hipamd/rocclr/hip_device.cpp @@ -235,6 +235,9 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device ) deviceProps.pageableMemoryAccess = info.hmmCpuMemoryAccessible_; deviceProps.pageableMemoryAccessUsesHostPageTables = info.hostUnifiedMemory_; + // hipStreamWaitValue64() and hipStreamWaitValue32() support + deviceProps.canUseStreamWaitValue = info.aqlBarrierValue_; + *props = deviceProps; HIP_RETURN(hipSuccess); } diff --git a/projects/clr/hipamd/rocclr/hip_device_runtime.cpp b/projects/clr/hipamd/rocclr/hip_device_runtime.cpp index 8bf4c7829e..ee82b04329 100755 --- a/projects/clr/hipamd/rocclr/hip_device_runtime.cpp +++ b/projects/clr/hipamd/rocclr/hip_device_runtime.cpp @@ -314,6 +314,9 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) case hipDeviceAttributePageableMemoryAccessUsesHostPageTables: *pi = prop.pageableMemoryAccessUsesHostPageTables; break; + case hipDeviceAttributeCanUseStreamWaitValue: + *pi = prop.canUseStreamWaitValue; + break; default: HIP_RETURN(hipErrorInvalidValue); } diff --git a/projects/clr/hipamd/rocclr/hip_hcc.map.in b/projects/clr/hipamd/rocclr/hip_hcc.map.in index 7f4bffdedc..f556085cf6 100755 --- a/projects/clr/hipamd/rocclr/hip_hcc.map.in +++ b/projects/clr/hipamd/rocclr/hip_hcc.map.in @@ -172,6 +172,10 @@ global: hipStreamQuery; hipStreamSynchronize; hipStreamWaitEvent; + hipStreamWaitValue32; + hipStreamWaitValue64; + hipStreamWriteValue32; + hipStreamWriteValue64; __hipPopCallConfiguration; __hipPushCallConfiguration; __hipRegisterFatBinary; diff --git a/projects/clr/hipamd/rocclr/hip_memory.cpp b/projects/clr/hipamd/rocclr/hip_memory.cpp index d8a089dcc2..dbbbbb0ad0 100755 --- a/projects/clr/hipamd/rocclr/hip_memory.cpp +++ b/projects/clr/hipamd/rocclr/hip_memory.cpp @@ -262,11 +262,21 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_API(hipExtMallocWithFlags, ptr, sizeBytes, flags); - if (flags != hipDeviceMallocDefault && flags != hipDeviceMallocFinegrained) { + unsigned int ihipFlags = 0; + if (flags == hipDeviceMallocDefault) { + ihipFlags = 0; + } else if (flags == hipDeviceMallocFinegrained) { + ihipFlags = CL_MEM_SVM_ATOMICS; + } else if (flags == hipMallocSignalMemory) { + ihipFlags = CL_MEM_SVM_ATOMICS | CL_MEM_SVM_FINE_GRAIN_BUFFER | ROCCLR_MEM_HSA_SIGNAL_MEMORY; + if (sizeBytes != 8) { + HIP_RETURN(hipErrorInvalidValue); + } + } else { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(ihipMalloc(ptr, sizeBytes, (flags & hipDeviceMallocFinegrained)? CL_MEM_SVM_ATOMICS: 0), (ptr != nullptr)? *ptr : nullptr); + HIP_RETURN(ihipMalloc(ptr, sizeBytes, ihipFlags), (ptr != nullptr)? *ptr : nullptr); } hipError_t hipMalloc(void** ptr, size_t sizeBytes) { diff --git a/projects/clr/hipamd/rocclr/hip_stream_ops.cpp b/projects/clr/hipamd/rocclr/hip_stream_ops.cpp new file mode 100644 index 0000000000..213d8fe638 --- /dev/null +++ b/projects/clr/hipamd/rocclr/hip_stream_ops.cpp @@ -0,0 +1,129 @@ +/* Copyright (c) 2015-present Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#include +#include "hip_internal.hpp" +#include "platform/command_utils.hpp" + +hipError_t ihipStreamOperation(hipStream_t stream, cl_command_type cmdType, void* ptr, + int64_t value, uint64_t mask, unsigned int flags, size_t sizeBytes) { + size_t offset = 0; + unsigned int outFlags = 0; + + amd::Memory* memory = getMemoryObject(ptr, offset); + if (!memory) { + return hipErrorInvalidValue; + } + + // NOTE: 'mask' is only used in Wait operation, 'sizeBytes' is only used in Write operation + // 'flags' for now used only for Wait, but in future there will usecases for Write too. + + if (cmdType == ROCCLR_COMMAND_STREAM_WAIT_VALUE) { + // Wait is only supported on SignalMemory objects + if (!(memory->getMemFlags() & ROCCLR_MEM_HSA_SIGNAL_MEMORY)) { + return hipErrorInvalidValue; + } + switch (flags) { + case hipStreamWaitValueGte: + outFlags = ROCCLR_STREAM_WAIT_VALUE_GTE; + break; + case hipStreamWaitValueEq: + outFlags = ROCCLR_STREAM_WAIT_VALUE_EQ; + break; + case hipStreamWaitValueAnd: + outFlags = ROCCLR_STREAM_WAIT_VALUE_AND; + break; + case hipStreamWaitValueNor: + outFlags = ROCCLR_STREAM_WAIT_VALUE_NOR; + break; + default: + ShouldNotReachHere(); + break; + } + } else if (cmdType != ROCCLR_COMMAND_STREAM_WRITE_VALUE) { + return hipErrorInvalidValue; + } + + amd::HostQueue* queue = hip::getQueue(stream); + amd::Command::EventWaitList waitList; + + amd::StreamOperationCommand* command = + new amd::StreamOperationCommand(*queue, cmdType, waitList, *memory->asBuffer(), + value, mask, outFlags, offset, sizeBytes); + + if (command == nullptr) { + return hipErrorOutOfMemory; + } + command->enqueue(); + command->release(); + return hipSuccess; +} + +hipError_t hipStreamWaitValue32(hipStream_t stream, void* ptr, int32_t value, uint32_t mask, + unsigned int flags) { + HIP_INIT_API(hipStreamWaitValue32, stream, ptr, value, mask, flags); + // NOTE: ptr corresponds to a HSA Signal memeory which is 64 bits. + // 32 bit value and mask are converted to 64-bit values. + HIP_RETURN_DURATION(ihipStreamOperation( + stream, + ROCCLR_COMMAND_STREAM_WAIT_VALUE, + ptr, + value, + mask, + flags, + 0)); // sizeBytes un-used for wait, set it to 0 +} + +hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, int64_t value, uint64_t mask, + unsigned int flags) { + HIP_INIT_API(hipStreamWaitValue64, stream, ptr, value, mask, flags); + HIP_RETURN_DURATION(ihipStreamOperation( + stream, + ROCCLR_COMMAND_STREAM_WAIT_VALUE, + ptr, + value, + mask, + flags, + 0)); // sizeBytes un-used for wait, set it to 0 +} + +hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, int32_t value) { + HIP_INIT_API(hipStreamWriteValue32, stream, ptr, value); + HIP_RETURN_DURATION(ihipStreamOperation( + stream, + ROCCLR_COMMAND_STREAM_WRITE_VALUE, + ptr, + value, + 0, // mask un-used set it to 0 + 0, // flags un-used for now set it to 0 + 4)); +} + +hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, int64_t value) { + HIP_INIT_API(hipStreamWriteValue64, stream, ptr, value); + HIP_RETURN_DURATION(ihipStreamOperation( + stream, + ROCCLR_COMMAND_STREAM_WRITE_VALUE, + ptr, + value, + 0, // mask un-used set it to 0 + 0, // flags un-used for now set it to 0 + 8)); +} diff --git a/projects/clr/hipamd/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp b/projects/clr/hipamd/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp new file mode 100644 index 0000000000..733b65098a --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp @@ -0,0 +1,279 @@ +/* +Copyright (c) 2020-present 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +// Test Description: +/* + This unit test is written to test Stream Write and Stream Wait API. + Stream Write: + Both 32 and 64 bit version of this APIs are tested by writing a specific value and checking + the correctness. Various mememory objects (host registered, device and Signal Memory) are tested. + Stream Wait: + Wait API is tested using two memory locations (DataPr and SignalPtr). Following + commands are executed for each type of wait operaitons (GEQ, EQ, AND and NOR) in the order + specified. + 1. CPU : An intial values are written to DataPtr and SignalPtr + 2. GPU : Wait operation (with false condition that blocks the stream) is enqued. + 3. GPU : Write operation on DataPtr to update its value is enqued. + 4. CPU : A query or CPU wait to make sure all commands are processed by GPU. + 5. CPU : streamQuery is performed to make sure it is not finshed executing the commands, + since step-2 is blocking. + 6. CPU : A new value is written to SignalPtr memory that make wait condition defined in + step-2 to be true. This causes step-3 to be executed. + 7. CPU : Synchronize the stream and read value at DataPtr, it should be equal to updated + value (step-3). +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 EXCLUDE_HIP_PLATFORM nvidia + * TEST: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" + +// Random predefiend 32 and 64 bit values +constexpr int32_t value32 = 0x70F0F0FF; +constexpr int64_t value64 = 0x7FFF0000FFFF0000; + +constexpr float SLEEP_MS = 100; +void testWrite() { + + int64_t* signalPtr; + + hipStream_t stream; + hipStreamCreate(&stream); + + int64_t* host_ptr64 = (int64_t *) malloc(sizeof(int64_t)); + int32_t* host_ptr32 = (int32_t *) malloc(sizeof(int32_t)); + std::cout << " hipStreamWriteValue: testing ... \n"; + + HIPCHECK(hipExtMallocWithFlags((void **)&signalPtr, 8, hipMallocSignalMemory)); + + void* device_ptr64; + void* device_ptr32; + + *host_ptr64 = 0x0; + *host_ptr32 = 0x0; + *signalPtr = 0x0; + + hipHostRegister(host_ptr64, sizeof(int64_t), 0); + hipHostRegister(host_ptr32, sizeof(int32_t), 0); + + // Test writting registered host pointer + HIPCHECK(hipStreamWriteValue64(stream, host_ptr64, value64)); + HIPCHECK(hipStreamWriteValue32(stream, host_ptr32, value32)); + hipStreamSynchronize(stream); + + HIPASSERT(*host_ptr64 == value64); + HIPASSERT(*host_ptr32 == value32); + + // Test writting device pointer + hipHostGetDevicePointer((void**)&device_ptr64, host_ptr64, 0); + hipHostGetDevicePointer((void**)&device_ptr32, host_ptr32, 0); + + // Reset values + *host_ptr64 = 0x0; + *host_ptr32 = 0x0; + + HIPCHECK(hipStreamWriteValue64(stream, device_ptr64, value64)); + HIPCHECK(hipStreamWriteValue32(stream, device_ptr32, value32)); + hipStreamSynchronize(stream); + + HIPASSERT(*host_ptr64 == value64); + HIPASSERT(*host_ptr32 == value32); + + // Test Writing to Signal Memory + HIPCHECK(hipStreamWriteValue64(stream, signalPtr, value64)); + hipStreamSynchronize(stream); + + HIPASSERT(*signalPtr == value64); + + // Cleanup + hipStreamDestroy(stream); + hipHostUnregister(host_ptr64); + hipHostUnregister(host_ptr32); + HIPCHECK(hipFree(signalPtr)); + free(host_ptr32); + free(host_ptr64); +} + +bool streamWaitValueSupported() { + int device_num = 0; + HIPCHECK(hipGetDeviceCount(&device_num)); + int waitValueSupport; + for (int device_id = 0; device_id < device_num; ++device_id) { + HIPCHECK(hipSetDevice(device_id)); + waitValueSupport = 0; + HIPCHECK(hipDeviceGetAttribute(&waitValueSupport, hipDeviceAttributeCanUseStreamWaitValue, + device_id)); + if (waitValueSupport == 1) return true; + } + return false; +} + +void testWait() { + int64_t* signalPtr; + // random data values + int32_t DATA_INIT = 0x1234; + int32_t DATA_UPDATE = 0X4321; + + struct TEST_WAIT { + int compareOp; + uint64_t mask; + int64_t waitValue; + int64_t signalValueFail; + int64_t signalValuePass; + }; + + + TEST_WAIT testCases[] = { + { + // mask will ignore few MSB bits + hipStreamWaitValueGte, + 0x0000FFFFFFFFFFFF, + 0x000000007FFF0001, + 0x7FFF00007FFF0000, + 0x000000007FFF0001 + }, + { + hipStreamWaitValueGte, + 0xF, + 0x4, + 0x3, + 0x6 + }, + { + // mask will ignore few MSB bits + hipStreamWaitValueEq, + 0x0000FFFFFFFFFFFF, + 0x000000000FFF0001, + 0x7FFF00000FFF0000, + 0x7F0000000FFF0001 + }, + { + hipStreamWaitValueEq, + 0xFF, + 0x11, + 0x25, + 0x11 + }, + { + // mask will discard bits 8 to 11 + hipStreamWaitValueAnd, + 0xFF, + 0xF4A, + 0xF35, + 0X02 + }, + { + // mask is set to ignore the sign bit. + hipStreamWaitValueNor, + 0x7FFFFFFFFFFFFFFF, + 0x7FFFFFFFFFFFF247, + 0x7FFFFFFFFFFFFdbd, + 0x7FFFFFFFFFFFFdb5 + }, + { + // mask is set to apply NOR for bits 0 to 3. + hipStreamWaitValueNor, + 0xF, + 0x7E, + 0x7D, + 0x76 + } + }; + + if (!streamWaitValueSupported()) { + std::cout << " hipStreamWaitValue: not supported on this device , skipping ... \n"; + return; + } + std::cout << " hipStreamWaitValue: testing ... \n"; + hipStream_t stream; + hipStreamCreate(&stream); + + HIPCHECK(hipExtMallocWithFlags((void **)&signalPtr, 8, hipMallocSignalMemory)); + int64_t* dataPtr64 = (int64_t *) malloc(sizeof(int64_t)); + int32_t* dataPtr32 = (int32_t *) malloc(sizeof(int32_t)); + hipHostRegister(dataPtr64, sizeof(int64_t), 0); + hipHostRegister(dataPtr32, sizeof(int32_t), 0); + + // We run all test cases twice + // Run-1: streamWait is blocking (wait conditions is false) + // Run-2: streamWait is non-blocking (wait condition is true) + for (int run = 0; run < 2; run++) { + bool isBlocking = run == 0; + + for (const auto& tc : testCases) { + *signalPtr = isBlocking ? tc.signalValueFail : tc.signalValuePass; + *dataPtr64 = DATA_INIT; + + HIPCHECK(hipStreamWaitValue64(stream, signalPtr, tc.waitValue, tc.mask, tc.compareOp)); + HIPCHECK(hipStreamWriteValue64(stream, dataPtr64, DATA_UPDATE)); + + if (isBlocking) { + // Trigger an implict flush and verify stream has pending work. + HIPASSERT(hipStreamQuery(stream) == hipErrorNotReady); + + // update signal to unblock the wait. + *signalPtr = tc.signalValuePass; + } + // HIPASSERT(hipStreamQuery(stream) == hipSuccess); + hipStreamSynchronize(stream); + HIPASSERT(*dataPtr64 == DATA_UPDATE); + + // 32-bit API + *signalPtr = isBlocking ? tc.signalValueFail : tc.signalValuePass; + *dataPtr32 = DATA_INIT; + + HIPCHECK(hipStreamWaitValue32(stream, signalPtr, static_cast(tc.waitValue), + static_cast(tc.mask), tc.compareOp)); + HIPCHECK(hipStreamWriteValue32(stream, dataPtr32, DATA_UPDATE)); + + if (isBlocking) { + // For DEBUG only + // usleep(500); + // HIPASSERT(*dataPtr32 == DATA_INIT); + + // Trigger an implict flush and verify stream has pending work. + HIPASSERT(hipStreamQuery(stream) == hipErrorNotReady); + + // update signal to unblock the wait. + *signalPtr = static_cast(tc.signalValuePass); + } + hipStreamSynchronize(stream); + HIPASSERT(*dataPtr32 == DATA_UPDATE); + } + } + + // Cleanup + HIPCHECK(hipFree(signalPtr)); + hipHostUnregister(dataPtr64); + hipHostUnregister(dataPtr32); + free(dataPtr64); + free(dataPtr32); + hipStreamDestroy(stream); +} + + +int main() { + testWrite(); + testWait(); + passed(); +}