AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

SWDEV-179105 - Stream Operations: Add swqupport for Wait and Write

فهرست منبع 
Change-Id: I9d6d0665d12b62fe705ce1569a0e8264a4f23ab7


[ROCm/clr commit: fd0f49503f]
This commit is contained in:
Ravi C Akkenapally
2020-12-14 15:28:01 -08:00
والد f6a1072ffe
کامیت 4cf7a5799a
9فایلهای تغییر یافته به همراه555 افزوده شده و 2 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h
+119
مشاهده پرونده
@@ -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
/**
* @}
*/
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
projects/clr/hipamd/include/hip/hip_runtime_api.h
+5
مشاهده پرونده
@@ -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 {
projects/clr/hipamd/rocclr/CMakeLists.txt
+1
مشاهده پرونده
@@ -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
projects/clr/hipamd/rocclr/hip_device.cpp
+3
مشاهده پرونده
@@ -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);
}
projects/clr/hipamd/rocclr/hip_device_runtime.cpp
+3
مشاهده پرونده
@@ -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);
}
projects/clr/hipamd/rocclr/hip_hcc.map.in
+4
مشاهده پرونده
@@ -172,6 +172,10 @@ global:
hipStreamQuery;
hipStreamSynchronize;
hipStreamWaitEvent;
hipStreamWaitValue32;
hipStreamWaitValue64;
hipStreamWriteValue32;
hipStreamWriteValue64;
__hipPopCallConfiguration;
__hipPushCallConfiguration;
__hipRegisterFatBinary;
projects/clr/hipamd/rocclr/hip_memory.cpp
+12 -2
مشاهده پرونده
@@ -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) {
projects/clr/hipamd/rocclr/hip_stream_ops.cpp
+129
مشاهده پرونده
@@ -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 <hip/hip_runtime.h>
#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));
}
projects/clr/hipamd/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp
+279
مشاهده پرونده
@@ -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 <unistd.h>
#include <hip/hip_runtime.h>
#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<int32_t>(tc.waitValue),
static_cast<uint32_t>(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<int32_t>(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();
}