diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index beb84bdb6f..8516d7520d 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -148,6 +148,8 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) hip_status = hipErrorMemoryAllocation; } + //printf (" hipMalloc allocated %p\n", *ptr); + return ihipLogStatus(hip_status); } diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index f400e27799..e044b6b578 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -74,7 +74,7 @@ hipError_t hipStreamCreate(hipStream_t *stream) } -#ifndef USE_AV_COPY +#if USE_AV_COPY==0 //--- /** * @bug This function conservatively waits for all work in the specified stream to complete. @@ -93,7 +93,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int bool fastWait = false; -#ifdef USE_AV_COPY +#if USE_AV_COPY if (stream != hipStreamNull) { printf ("HIP: wait locked stream\n"); stream->locked_waitEvent(event); diff --git a/projects/hip/tests/src/hipPointerAttrib.cpp b/projects/hip/tests/src/hipPointerAttrib.cpp index ab620eb022..f0d04f81bf 100644 --- a/projects/hip/tests/src/hipPointerAttrib.cpp +++ b/projects/hip/tests/src/hipPointerAttrib.cpp @@ -28,11 +28,12 @@ THE SOFTWARE. #include #ifdef __HIP_PLATFORM_HCC__ -//#include "hcc_detail/AM.h" #include "hc_am.hpp" #endif +#define USE_AV_COPY (__hcc_workweek__ >= 16351) + size_t Nbytes = 0; //================================================================================================= @@ -409,11 +410,21 @@ void thread_noise_generator(int iters, size_t numBuffers, Dir addDir, Dir remove if (addDir == Up) { for (char *p = basePtr; p=0; p-=bufferSize) { +#if USE_AV_COPY + hc::AmPointerInfo info(p, p, bufferSize, acc, false, false); + hc::am_memtracker_add(p, info); +#else hc::am_memtracker_add(p, bufferSize, acc, false); +#endif } } diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp index 63685bb509..7bdae6a96f 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp @@ -62,7 +62,7 @@ Streamer::Streamer(size_t numElements) : template void Streamer::runAsync() { - printf ("testing: %s numElements=%zu\n", __func__, _numElements); + printf ("testing: %s numElements=%zu size=%6.2fMB\n", __func__, _numElements, _numElements * sizeof(T) / 1024.0/1024.0); unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, _numElements); hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, _stream, _A_d, _B_d, _C_d, _numElements); HIPCHECK(hipEventRecord(_event, _stream)); @@ -117,4 +117,6 @@ int main(int argc, char *argv[]) } HIPCHECK(hipDeviceSynchronize()); + + passed(); }