changed copyright year from 2016 to 2017 in src directory
Change-Id: Idb97db509b2b4b1656b2df7a14a62ade38c9d574
[ROCm/clr commit: e9ff23e5f9]
This commit is contained in:
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -54,5 +54,3 @@ __device__ double __hiloint2double (int hi, int lo) {
|
||||
s.s2.lo = lo;
|
||||
return s.d;
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -479,4 +479,3 @@ hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop )
|
||||
}
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -43,7 +43,7 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
|
||||
eh->_stream = NULL;
|
||||
eh->_flags = flags;
|
||||
eh->_timestamp = 0;
|
||||
*event = eh;
|
||||
*event = eh;
|
||||
} else {
|
||||
e = hipErrorInvalidValue;
|
||||
}
|
||||
@@ -186,5 +186,3 @@ hipError_t hipEventQuery(hipEvent_t event)
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -279,12 +279,12 @@ inline void ihipStream_t::ensureHaveQueue(LockedAccessor_StreamCrit_t &streamCri
|
||||
// TODO
|
||||
auto needyCritPtr = this->_criticalData.mlock();
|
||||
|
||||
// Second test to ensure we still need to steal the queue - another thread may have
|
||||
// Second test to ensure we still need to steal the queue - another thread may have
|
||||
// snuck in here and already solved the issue.
|
||||
if (!needyCritPtr->_hasQueue) {
|
||||
needyCritPtr->_av = this->_ctx->stealActiveQueue(ctxCrit, this);
|
||||
}
|
||||
|
||||
|
||||
streamCrit->_hasQueue = true;
|
||||
}
|
||||
assert(streamCrit->_hasQueue);
|
||||
@@ -394,7 +394,7 @@ LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand()
|
||||
}
|
||||
|
||||
this->ensureHaveQueue(crit);
|
||||
|
||||
|
||||
|
||||
|
||||
return crit;
|
||||
@@ -944,10 +944,10 @@ ihipCtx_t::stealActiveQueue(LockedAccessor_CtxCrit_t &ctxCrit, ihipStream_t *nee
|
||||
uint64_t *p = (uint64_t*)(&victimCritPtr->_av);
|
||||
*p = 0; // damage the victim av so attempt to use it will fault.
|
||||
|
||||
(*iter)->_criticalData.munlock();
|
||||
(*iter)->_criticalData.munlock();
|
||||
return av;
|
||||
}
|
||||
(*iter)->_criticalData.munlock();
|
||||
}
|
||||
(*iter)->_criticalData.munlock();
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -1296,7 +1296,7 @@ void ihipInit()
|
||||
tokenize(HIP_LAUNCH_BLOCKING_KERNELS, ',', &g_hipLaunchBlockingKernels);
|
||||
}
|
||||
READ_ENV_I(release, HIP_API_BLOCKING, 0, "Make HIP APIs 'host-synchronous', so they block until completed. Impacts hipMemcpyAsync, hipMemsetAsync." );
|
||||
|
||||
|
||||
|
||||
READ_ENV_I(release, HIP_MAX_QUEUES, 0, "Maximum number of queues that this app will use per-device. Additional streams will share the specified number of queues. 0=no limit.");
|
||||
|
||||
@@ -1320,8 +1320,8 @@ void ihipInit()
|
||||
|
||||
|
||||
READ_ENV_I(release, HIP_WAIT_MODE, 0, "Force synchronization mode. 1= force yield, 2=force spin, 0=defaults specified in application");
|
||||
READ_ENV_I(release, HIP_FORCE_P2P_HOST, 0, "Force use of host/staging copy for peer-to-peer copies.1=always use copies, 2=always return false for hipDeviceCanAccessPeer");
|
||||
READ_ENV_I(release, HIP_FORCE_SYNC_COPY, 0, "Force all copies (even hipMemcpyAsync) to use sync copies");
|
||||
READ_ENV_I(release, HIP_FORCE_P2P_HOST, 0, "Force use of host/staging copy for peer-to-peer copies.1=always use copies, 2=always return false for hipDeviceCanAccessPeer");
|
||||
READ_ENV_I(release, HIP_FORCE_SYNC_COPY, 0, "Force all copies (even hipMemcpyAsync) to use sync copies");
|
||||
|
||||
// TODO - review, can we remove this?
|
||||
READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Max number of inflight kernels per stream before active synchronization is forced.");
|
||||
@@ -2026,7 +2026,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a
|
||||
stream = device->_defaultStream;
|
||||
}
|
||||
|
||||
*av = stream->locked_getAv(); // TODO - review.
|
||||
*av = stream->locked_getAv(); // TODO - review.
|
||||
|
||||
hipError_t err = hipSuccess;
|
||||
return ihipLogStatus(err);
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -415,15 +415,15 @@ public:
|
||||
|
||||
ihipStreamCriticalBase_t<StreamMutex> * mlock() { LockedBase<MUTEX_TYPE>::lock(); return this;};
|
||||
|
||||
void munlock() {
|
||||
void munlock() {
|
||||
tprintf(DB_SYNC, "munlocking criticalData=%p for %s...\n", this, ToString(this->_parent).c_str());
|
||||
LockedBase<MUTEX_TYPE>::unlock();
|
||||
LockedBase<MUTEX_TYPE>::unlock();
|
||||
};
|
||||
|
||||
ihipStreamCriticalBase_t<StreamMutex> * mtry_lock() {
|
||||
ihipStreamCriticalBase_t<StreamMutex> * mtry_lock() {
|
||||
bool gotLock = LockedBase<MUTEX_TYPE>::try_lock() ;
|
||||
tprintf(DB_SYNC, "mtry_locking=%d criticalData=%p for %s...\n", gotLock, this, ToString(this->_parent).c_str());
|
||||
return gotLock ? this: nullptr;
|
||||
return gotLock ? this: nullptr;
|
||||
};
|
||||
|
||||
public:
|
||||
@@ -683,7 +683,7 @@ public: // Functions:
|
||||
ihipCtx_t(ihipDevice_t *device, unsigned deviceCnt, unsigned flags); // note: calls constructor for _criticalData
|
||||
~ihipCtx_t();
|
||||
|
||||
// Functions which read or write the critical data are named locked_.
|
||||
// Functions which read or write the critical data are named locked_.
|
||||
// (might be better called "locking_"
|
||||
// ihipCtx_t does not use recursive locks so the ihip implementation must avoid calling a locked_ function from within a locked_ function.
|
||||
// External functions which call several locked_ functions will acquire and release the lock for each function. if this occurs in
|
||||
@@ -697,7 +697,7 @@ public: // Functions:
|
||||
hc::accelerator_view stealActiveQueue(LockedAccessor_CtxCrit_t &ctxCrit, ihipStream_t *needyStream);
|
||||
hc::accelerator_view createOrStealQueue(LockedAccessor_CtxCrit_t &ctxCrit);
|
||||
|
||||
ihipCtxCritical_t &criticalData() { return _criticalData; };
|
||||
ihipCtxCritical_t &criticalData() { return _criticalData; };
|
||||
|
||||
const ihipDevice_t *getDevice() const { return _device; };
|
||||
int getDeviceNum() const { return _device->_deviceId; };
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -1087,7 +1087,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
|
||||
hsa_amd_ipc_memory_attach(&handle->ipc_handle, handle->psize, 1, agent, devPtr);
|
||||
if(hsa_status != HSA_STATUS_SUCCESS)
|
||||
hipStatus = hipErrorMapBufferObjectFailed;
|
||||
#else
|
||||
#else
|
||||
hipStatus = hipErrorRuntimeOther;
|
||||
#endif
|
||||
return hipStatus;
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -403,5 +403,3 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
|
||||
}
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -31,7 +31,7 @@ THE SOFTWARE.
|
||||
// There are two flavors:
|
||||
// - one where contexts are specified with hipCtx_t type.
|
||||
// - one where contexts are specified with integer deviceIds, that are mapped to the primary context for that device.
|
||||
// The implementation contains a set of internal ihip* functions which operate on contexts. Then the
|
||||
// The implementation contains a set of internal ihip* functions which operate on contexts. Then the
|
||||
// public APIs are thin wrappers which call into this internal implementations.
|
||||
// TODO - actually not yet - currently the integer deviceId flavors just call the context APIs. need to fix.
|
||||
|
||||
@@ -46,16 +46,16 @@ hipError_t ihipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx
|
||||
|
||||
if (thisCtx == peerCtx) {
|
||||
*canAccessPeer = 0;
|
||||
tprintf(DB_MEM, "Can't be peer to self. (this=%s, peer=%s)\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str());
|
||||
tprintf(DB_MEM, "Can't be peer to self. (this=%s, peer=%s)\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str());
|
||||
} else if (HIP_FORCE_P2P_HOST & 0x2) {
|
||||
*canAccessPeer = false;
|
||||
tprintf(DB_MEM, "HIP_FORCE_P2P_HOST denies peer access this=%s peer=%s canAccessPeer=%d\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str(), *canAccessPeer);
|
||||
tprintf(DB_MEM, "HIP_FORCE_P2P_HOST denies peer access this=%s peer=%s canAccessPeer=%d\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str(), *canAccessPeer);
|
||||
} else {
|
||||
*canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc);
|
||||
tprintf(DB_MEM, "deviceCanAccessPeer this=%s peer=%s canAccessPeer=%d\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str(), *canAccessPeer);
|
||||
tprintf(DB_MEM, "deviceCanAccessPeer this=%s peer=%s canAccessPeer=%d\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str(), *canAccessPeer);
|
||||
}
|
||||
|
||||
} else {
|
||||
@@ -99,14 +99,14 @@ hipError_t ihipDisablePeerAccess (hipCtx_t peerCtx)
|
||||
LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
|
||||
bool changed = peerCrit->removePeerWatcher(peerCtx, thisCtx);
|
||||
if (changed) {
|
||||
tprintf(DB_MEM, "device %s disable access to memory allocated on peer:%s\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str());
|
||||
tprintf(DB_MEM, "device %s disable access to memory allocated on peer:%s\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str());
|
||||
// Update the peers for all memory already saved in the tracker:
|
||||
am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
|
||||
} else {
|
||||
err = hipErrorPeerAccessNotEnabled; // never enabled P2P access.
|
||||
}
|
||||
}
|
||||
}
|
||||
} else {
|
||||
err = hipErrorInvalidDevice;
|
||||
}
|
||||
@@ -133,8 +133,8 @@ hipError_t ihipEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
|
||||
// Add thisCtx to peerCtx's access list so that new allocations on peer will be made visible to this device:
|
||||
bool isNewPeer = peerCrit->addPeerWatcher(peerCtx, thisCtx);
|
||||
if (isNewPeer) {
|
||||
tprintf(DB_MEM, "device=%s can now see all memory allocated on peer=%s\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str());
|
||||
tprintf(DB_MEM, "device=%s can now see all memory allocated on peer=%s\n",
|
||||
thisCtx->toString().c_str(), peerCtx->toString().c_str());
|
||||
am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
|
||||
} else {
|
||||
err = hipErrorPeerAccessAlreadyEnabled;
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -214,7 +214,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
|
||||
{
|
||||
HIP_INIT_API(stream, callback, userData, flags);
|
||||
hipError_t e = hipSuccess;
|
||||
//--- explicitly synchronize stream to add callback routines
|
||||
//--- explicitly synchronize stream to add callback routines
|
||||
hipStreamSynchronize(stream);
|
||||
callback(stream, e, userData);
|
||||
return ihipLogStatus(e);
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 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
|
||||
@@ -55,7 +55,7 @@ inline std::string ToHexString(T v)
|
||||
|
||||
// This is the default which works for most types:
|
||||
template <typename T>
|
||||
inline std::string ToString(T v)
|
||||
inline std::string ToString(T v)
|
||||
{
|
||||
std::ostringstream ss;
|
||||
ss << v;
|
||||
@@ -65,7 +65,7 @@ inline std::string ToString(T v)
|
||||
|
||||
// hipEvent_t specialization. TODO - maybe add an event ID for debug?
|
||||
template <>
|
||||
inline std::string ToString(hipEvent_t v)
|
||||
inline std::string ToString(hipEvent_t v)
|
||||
{
|
||||
std::ostringstream ss;
|
||||
ss << v;
|
||||
@@ -74,7 +74,7 @@ inline std::string ToString(hipEvent_t v)
|
||||
|
||||
// hipEvent_t specialization. TODO - maybe add an event ID for debug?
|
||||
template <>
|
||||
inline std::string ToString(hipFunction_t v)
|
||||
inline std::string ToString(hipFunction_t v)
|
||||
{
|
||||
std::ostringstream ss;
|
||||
ss << "0x" << std::hex << v._object;
|
||||
@@ -85,7 +85,7 @@ inline std::string ToString(hipFunction_t v)
|
||||
|
||||
// hipStream_t
|
||||
template <>
|
||||
inline std::string ToString(hipStream_t v)
|
||||
inline std::string ToString(hipStream_t v)
|
||||
{
|
||||
std::ostringstream ss;
|
||||
if (v == NULL) {
|
||||
@@ -99,7 +99,7 @@ inline std::string ToString(hipStream_t v)
|
||||
|
||||
// hipMemcpyKind specialization
|
||||
template <>
|
||||
inline std::string ToString(hipMemcpyKind v)
|
||||
inline std::string ToString(hipMemcpyKind v)
|
||||
{
|
||||
switch(v) {
|
||||
CASE_STR(hipMemcpyHostToHost);
|
||||
@@ -113,14 +113,14 @@ inline std::string ToString(hipMemcpyKind v)
|
||||
|
||||
|
||||
template <>
|
||||
inline std::string ToString(hipError_t v)
|
||||
inline std::string ToString(hipError_t v)
|
||||
{
|
||||
return ihipErrorString(v);
|
||||
};
|
||||
|
||||
|
||||
// Catch empty arguments case
|
||||
inline std::string ToString()
|
||||
inline std::string ToString()
|
||||
{
|
||||
return ("");
|
||||
}
|
||||
@@ -129,8 +129,8 @@ inline std::string ToString()
|
||||
//---
|
||||
// C++11 variadic template - peels off first argument, converts to string, and calls itself again to peel the next arg.
|
||||
// Strings are automatically separated by comma+space.
|
||||
template <typename T, typename... Args>
|
||||
inline std::string ToString(T first, Args... args)
|
||||
template <typename T, typename... Args>
|
||||
inline std::string ToString(T first, Args... args)
|
||||
{
|
||||
return ToString(first) + ", " + ToString(args...) ;
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user