From a236942ca5c92e57151b6fc3e1ed6ecd866814bf Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Sat, 12 Mar 2016 04:04:57 -0600
Subject: [PATCH] Remove USE_AM_TRACKER (assume 1)

[ROCm/clr commit: de4bb0ca0e0f6ea60ccac78b72c2cba52d62f21b]
---
 projects/clr/hipamd/src/hip_hcc.cpp | 43 -----------------------------
 1 file changed, 43 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp
index 8f33206b38..6d4bafa0e2 100644
--- a/projects/clr/hipamd/src/hip_hcc.cpp
+++ b/projects/clr/hipamd/src/hip_hcc.cpp
@@ -45,13 +45,8 @@ THE SOFTWARE.
 
 
 
-#define USE_AM_TRACKER 1  /* >0 = use new AM memory tracker features. */
 #define USE_ROCR_V2    1  /* use the ROCR v2 async copy API with dst and src agents */
 
-#if (USE_AM_TRACKER) and (__hcc_workweek__ < 16074)
-#error (USE_AM_TRACKER requries HCC version of 16074 or newer)
-#endif
-
 
 
 
@@ -1357,12 +1352,10 @@ hipError_t hipDeviceReset(void)
     device->_streams.clear();
 
 
-#if USE_AM_TRACKER
     if (device) {
         am_memtracker_reset(device->_acc);
         device->reset(); // re-allocate required resources.
     }
-#endif
 
 	// TODO - reset all streams on the device.
 
@@ -1862,7 +1855,6 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr)
 
     hipError_t e = hipSuccess;
 
-#if USE_AM_TRACKER
     hc::accelerator acc;
     hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
     am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr);
@@ -1896,16 +1888,11 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr)
 
         e = hipErrorUnknown;
     }
-#else
-    e = hipErrorInvalidDevice;
-#endif
 
     return ihipLogStatus(e);
 }
 
 
-#if USE_AM_TRACKER
-// TODO - test this function:
 /**
  * @returns #hipSuccess,
  * @returns #hipErrorInvalidValue if flags are not 0
@@ -1934,7 +1921,6 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi
 
     return ihipLogStatus(e);
 }
-#endif
 
 
 
@@ -2031,9 +2017,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
         if (sizeBytes && (*ptr == NULL)) {
             hip_status = hipErrorMemoryAllocation;
         } else {
-#if USE_AM_TRACKER
             hc::am_memtracker_update(*ptr, device->_device_index, 0);
-#endif
         }
     } else {
         hip_status = hipErrorMemoryAllocation;
@@ -2057,9 +2041,7 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
         if (sizeBytes && (*ptr == NULL)) {
             hip_status = hipErrorMemoryAllocation;
         } else {
-#if USE_AM_TRACKER
             hc::am_memtracker_update(*ptr, device->_device_index, 0);
-#endif
         }
 
         tprintf (DB_MEM, "  %s: pinned ptr=%p\n", __func__, *ptr);
@@ -2320,7 +2302,6 @@ void StagingBuffer::CopyDeviceToHost(void* dst, const void* src, size_t sizeByte
 
 
 
-#if USE_AM_TRACKER
 void ihipSyncCopy(ihipStream_t *stream, void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
 {
     ihipDevice_t *device = stream->getDevice();
@@ -2429,7 +2410,6 @@ void ihipSyncCopy(ihipStream_t *stream, void* dst, const void* src, size_t sizeB
 
     }
 }
-#endif
 
 
 //---
@@ -2443,7 +2423,6 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
 
     hipError_t e = hipSuccess;
 
-#if USE_AM_TRACKER
     try {
         ihipSyncCopy(stream, dst, src, sizeBytes, kind);
     }
@@ -2452,11 +2431,6 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
     }
 
 
-#else
-    hc::am_copy(dst, src, sizeBytes);
-    e = hipSuccess;
-#endif
-
     if (HIP_LAUNCH_BLOCKING) {
         tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpy\n");
         stream->wait();
@@ -2466,11 +2440,6 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
 }
 
 
-#if USE_AM_TRACKER==0
-/**
- * @warning on HCC hipMemcpyAsync uses a synchronous copy.
- */
-#endif
 /**
  * @result #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidMemcpyDirection, #hipErrorInvalidValue
  * @warning on HCC hipMemcpyAsync does not support overlapped H2D and D2H copies.
@@ -2485,7 +2454,6 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
 
     stream = ihipSyncAndResolveStream(stream);
 
-#if USE_AM_TRACKER
     if (stream) {
         ihipDevice_t *device = stream->getDevice();
 
@@ -2560,12 +2528,6 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
     } else {
         e = hipErrorInvalidValue;
     }
-#else
-    // TODO-hsart This routine needs to ensure that dst and src are mapped on the GPU.
-    // This is a synchronous copy - remove and replace with code below when we have appropriate LOCK APIs.
-    hc::am_copy(dst, src, sizeBytes);
-#endif
-
 
     return ihipLogStatus(e);
 }
@@ -2650,15 +2612,10 @@ hipError_t hipMemGetInfo  (size_t *free, size_t *total)
         }
 
         if (free) {
-#if USE_AM_TRACKER
             // TODO - replace with kernel-level for reporting free memory:
             size_t deviceMemSize, hostMemSize, userMemSize;
             hc::am_memtracker_sizeinfo(hipDevice->_acc, &deviceMemSize, &hostMemSize, &userMemSize);
             *free =  hipDevice->_props.totalGlobalMem - deviceMemSize;
-#else
-            *free =  hipDevice->_props.totalGlobalMem * 0.5; // TODO
-            e=hipErrorInvalidValue;
-#endif
         }
 
     } else {