AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Cleanup hipEvent. (Intermediate checkpoint)

Browse Source 
Support hipEventDisableSystemRelease flag.
Update test.
Remove stray printf


[ROCm/hip commit: 620eb30691]
This commit is contained in:
Ben Sander
2017-05-26 14:48:27 -05:00
parent c88190d698
commit e24627aa80
3 changed files with 100 additions and 36 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/src/hip_event.cpp
+16 -12
View File
@@ -82,7 +82,7 @@ void ihipEvent_t::setTimestamp()
}
if (_state != hipEventStatusRecorded) {
printf (" not ready isReady0=%d val=%d isReady1=%d\n", isReady0, val, isReady1);
//printf (" not ready isReady0=%d val=%d isReady1=%d\n", isReady0, val, isReady1);
}
}
@@ -92,7 +92,10 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
hipError_t e = hipSuccess;
// TODO-IPC - support hipEventInterprocess.
unsigned supportedFlags = hipEventDefault | hipEventBlockingSync | hipEventDisableTiming;
unsigned supportedFlags = hipEventDefault
| hipEventBlockingSync
| hipEventDisableTiming
| hipEventDisableSystemRelease;
if ((flags & ~supportedFlags) == 0) {
ihipEvent_t *eh = new ihipEvent_t(flags);
@@ -197,20 +200,18 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
{
HIP_INIT_API(ms, start, stop);
ihipEvent_t *start_eh = start;
ihipEvent_t *stop_eh = stop;
start->setTimestamp();
stop->setTimestamp();
hipError_t status = hipSuccess;
*ms = 0.0f;
if (start_eh && stop_eh) {
if ((start_eh->_state == hipEventStatusRecorded) && (stop_eh->_state == hipEventStatusRecorded)) {
if (start && stop) {
// refresh status:
if ((start->_state == hipEventStatusRecorded) && (stop->_state == hipEventStatusRecorded)) {
// Common case, we have good information for both events.
int64_t tickDiff = (stop_eh->timestamp() - start_eh->timestamp());
int64_t tickDiff = (stop->timestamp() - start->timestamp());
uint64_t freqHz;
hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz);
@@ -223,13 +224,16 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
}
} else if ((start_eh->_state == hipEventStatusRecording) ||
(stop_eh->_state == hipEventStatusRecording)) {
} else if ((start->_state == hipEventStatusRecording) ||
(stop->_state == hipEventStatusRecording)) {
status = hipErrorNotReady;
} else if ((start_eh->_state == hipEventStatusUnitialized) ||
(stop_eh->_state == hipEventStatusUnitialized)) {
} else if ((start->_state == hipEventStatusUnitialized) ||
(stop->_state == hipEventStatusUnitialized)) {
status = hipErrorInvalidResourceHandle;
}
} else {
status = hipErrorInvalidResourceHandle;
}
return ihipLogStatus(status);
projects/hip/src/hip_hcc.cpp
+3 -3
View File
@@ -49,7 +49,7 @@ THE SOFTWARE.
// needs HCC change for hc::no_scope
#define USE_NO_SCOPE 0
#define USE_NO_SCOPE 1
//=================================================================================================
//Global variables:
@@ -331,10 +331,10 @@ void ihipStream_t::locked_recordEvent(hipEvent_t event)
LockedAccessor_StreamCrit_t crit(_criticalData);
#if USE_NO_SCOPE
printf ("create_marker, flags = %x\n", event->_flags);
//printf ("create_marker, flags = %x\n", event->_flags);
event->_marker = crit->_av.create_marker((event->_flags & hipEventDisableSystemRelease) ? hc::no_scope : hc::system_scope);
#else
event->_marker = crit->_av.create_marker();
event->_marker = crit->_av.create_marker((event->_flags & hipEventDisableSystemRelease) ? hc::accelerator_scope : hc::system_scope);
#endif
};
projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp
+81 -21
View File
@@ -42,6 +42,63 @@ __global__ void Set(int *Ad, int val){
Ad[tx] = val;
}
#define SYNC_EVENT 0
#define SYNC_STREAM 1
#define SYNC_DEVICE 2
std::vector<std::string> syncMsg = {"event", "stream", "device"};
void CheckHostPointer(int numElements, int *ptr, int syncMethod, std::string msg)
{
std::cerr << "test: CheckHostPointer " << msg
<< " ptr=" << ptr
<< " syncMethod=" << syncMsg[syncMethod] << "\n";
hipStream_t s;
hipEvent_t e;
// Init:
HIPCHECK(hipStreamCreate(&s));
HIPCHECK(hipEventCreateWithFlags(&e, hipEventDisableSystemRelease));
dim3 dimBlock(64,1,1);
dim3 dimGrid(numElements/dimBlock.x,1,1);
const int expected = 13;
// Init array to know state:
hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, 0x0, ptr, -42);
HIPCHECK(hipDeviceSynchronize());
hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, s, ptr, expected);
HIPCHECK(hipEventRecord(e, s));
// Host waits for event :
switch (syncMethod) {
case SYNC_EVENT:
HIPCHECK(hipEventSynchronize(e));
break;
case SYNC_STREAM:
HIPCHECK(hipStreamSynchronize(s));
break;
case SYNC_DEVICE:
HIPCHECK(hipDeviceSynchronize());
break;
default:
assert(0);
};
for (int i=0; i<numElements; i++) {
if (ptr[i] != expected) {
printf ("mismatch at %d: %d != %d\n", i, ptr[i], expected);
assert(ptr[i] == expected);
}
}
HIPCHECK(hipStreamDestroy(s));
HIPCHECK(hipEventDestroy(e));
};
int main(){
@@ -86,38 +143,41 @@ int main(){
}
{
int *A, *B;
int numElements = 1024*16;
size_t sizeBytes = numElements * sizeof (int);
#ifdef __HIP_PLATFORM_HCC__
HIPCHECK_API(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent|hipHostMallocNonCoherent), hipErrorInvalidValue);
assert (A == 0);
#ifdef __HIP_PLATFORM_HCC__
{
// Stimulate error condition:
int *A = &numElements;
HIPCHECK_API(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent|hipHostMallocNonCoherent), hipErrorInvalidValue);
assert (A == 0);
}
#endif
HIPCHECK(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent));
hipStream_t s;
hipEvent_t e;
// Init:
HIPCHECK(hipStreamCreate(&s));
HIPCHECK(hipEventCreateWithFlags(&e, 0));
dim3 dimBlock(64,1,1);
dim3 dimGrid(numElements/dimBlock.x,1,1);
{
int *A = nullptr;
HIPCHECK(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocNonCoherent));
const char *ptrType = "non-coherent"; // TODO
//CheckHostPointer(numElements, A, SYNC_DEVICE, ptrType);
//CheckHostPointer(numElements, A, SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, SYNC_EVENT, ptrType);
}
// Init array to know state:
hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, 0x0, A, -42);
HIPCHECK(hipDeviceSynchronize());
if (0) { // TODO, remove me
int *A = nullptr;
HIPCHECK(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent));
const char *ptrType = "coherent";
CheckHostPointer(numElements, A, SYNC_DEVICE, ptrType);
CheckHostPointer(numElements, A, SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, SYNC_EVENT, ptrType);
}
hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, s, A, 13);
HIPCHECK(hipEventRecord(e, s));
// Host waits for event :
HIPCHECK(hipEventSynchronize(e));
// check result?
HIPCHECK(hipHostMalloc((void**)&B, sizeBytes, hipHostMallocNonCoherent));
}
passed();