https://github.com/Microsoft/CNTK
Tip revision: 8fef21be1d80b02849548dd05993abcebafd4837 authored by Dong Yu on 09 August 2016, 21:51:45 UTC
support minibatchSize (for larger models that will be out of memory without it) and frameMode (for RNNNode based Windowed RNN) for HTKMLFWriter (and HTKMLFReader).
support minibatchSize (for larger models that will be out of memory without it) and frameMode (for RNNNode based Windowed RNN) for HTKMLFWriter (and HTKMLFReader).
Tip revision: 8fef21b
MatrixQuantizerGPU.cu
#include "stdafx.h"
#include "MatrixQuantizerGPU.h"
#include "MatrixQuantizer_kernel.cu"
#include "GPUMatrix.h"
#include "GPUDataTransferer.h"
#pragma comment(lib, "cudart.lib") // instruct linker to reference these libs
#pragma comment(lib, "cublas.lib")
#pragma comment(lib, "cusparse.lib")
#pragma comment(lib, "curand.lib")
#pragma warning(disable : 4267) // conversion from 'size_t' to 'unsigned int'; happens in CUDA <<<a,b>>> syntax if a and b are size_t
#pragma warning(disable : 4127) // conditional expression is constant; "if (sizeof(ElemType)==sizeof(float))" triggers this
#pragma warning(disable : 4702) // unreachable code; triggered for unknown reasons
namespace Microsoft { namespace MSR { namespace CNTK {
// CUDA failed
// Since the outer code sometimes does not recover properly, as an option we log and die right away.
// This is needed for our GCD farm which has intermittent CUDA errors that sometimes cause the DBN tool, when running with MPI, to hang instead of terminating.
void cudafail(const char* msg)
{
// TODO: get from an env variable
bool dieoncudafailure = false;
if (!dieoncudafailure)
{
RuntimeError("%s", msg);
}
fprintf(stderr, "%s\n", msg);
fprintf(stderr, "cudafail: terminating\n"), fflush(stderr);
#ifdef WIN32
TerminateProcess(GetCurrentProcess(), EXIT_FAILURE); // fail the hard way to ensure it won't hang elsewhere
#else
exit(1);
#endif
}
// allows to write cudaFunction() || "error" (CUDA runtime)
static
#ifdef WIN32
__declspec(noinline)
#endif
void
operator||(cudaError_t rc, const char* msg)
{
if (rc != cudaSuccess)
{
char buf[1000];
sprintf_s(buf, 1000, "%s: %s (cuda error %d)", msg, cudaGetErrorString(rc), rc);
cudafail(buf);
}
}
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::Sync()
{
cudaDeviceSynchronize() || "cudaDeviceSynchronize failed";
}
// wait until stream has completed all scheduled operations
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::SyncStream(cudaStream_t stream)
{
cudaStreamSynchronize(stream) || "cudaStreamSynchronize failed";
}
// same but for event
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::SyncEvent(cudaEvent_t ev)
{
auto rc = cudaEventQuery(ev);
if (rc != cudaErrorNotReady)
{
// if Event is ready then no need to wait
rc || "cudaEventQuery failed";
return;
}
// we must wait
cudaEventSynchronize(ev) || "cudaEventSynchronize failed";
}
//streams
template <class ElemType>
cudaStream_t MatrixQuantizerGPU<ElemType>::m_computeStream = NULL;
template <class ElemType>
cudaStream_t MatrixQuantizerGPU<ElemType>::m_fetchStream = NULL;
template <class ElemType>
cudaStream_t MatrixQuantizerGPU<ElemType>::m_assignStream = NULL;
template <class ElemType>
cudaStream_t MatrixQuantizerGPU<ElemType>::GetComputeStream()
{
return m_computeStream;
}
template <class ElemType>
cudaStream_t MatrixQuantizerGPU<ElemType>::GetFetchStream()
{
return m_fetchStream;
}
template <class ElemType>
cudaStream_t MatrixQuantizerGPU<ElemType>::GetAssignStream()
{
return m_assignStream;
}
/////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
// computestream: the stream the caller issued the quant op on
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::RecordQuantizeCompleteEvent(cudaStream_t computestream) const
{
// schedule to flag the quantize-complete event (on main stream)
cudaEventRecord(m_quantizeCompleteEvent, computestream) || "cudaEventRecord failed";
// when running synchronously (for time measurements), then we (CPU) wait right here
if (m_forceSync)
{
SyncStream(computestream);
}
}
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::SyncQuantizeCompleEventAndFetchAndRecordFetchCompleteEvent(char* cpuBuffer, char* gpuBuffer, size_t size) const
{
// schedule fetch stream to wait until the last quantize op is complete, i.e. the data in the buffer is now valid
// wait until commencement
cudaStreamWaitEvent(GetFetchStream(), m_quantizeCompleteEvent, 0 /*flags 'must be 0'*/) || "cudaStreamWaitEvent failed";
// schedule to fetch that quantized data into CPU buffer (on a separate transfer stream)
cudaMemcpyAsync(cpuBuffer, gpuBuffer, size, cudaMemcpyDeviceToHost, GetFetchStream()) || "cudaMemcpyAsync failed";
cudaEventRecord(m_fetchCompleteEvent, GetFetchStream()) || "cudaEventRecord failed"; // for next GPU operation
// when running synchronously (for time measurements), then we (CPU) wait right here
if (m_forceSync)
{
SyncStream(GetFetchStream());
}
}
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::SyncAssignCompleteEvent(cudaStream_t computestream) const
{
// schedule to wait for the assign-complete event (on main/compute stream) --CPU buffer free once main stream does anything after this
cudaStreamWaitEvent(computestream, m_assignCompleteEvent, 0 /*flags 'must be 0'*/) || "cudaStreamWaitEvent failed";
// Note that the NVidia doc says somewhat confusingly:
// * If \p stream is NULL, any future work submitted in any stream will wait for
// * \p event to complete before beginning execution. This effectively creates a
// * barrier for all future work submitted to the device on this thread.
// -> it says that this may bring the whole machinery to stall. Or does cudaStreamWaitEvent() honor cudaStreamNonBlocking?
// According to NVidia (Jiri Kraus), this works as expected.
}
template <class ElemType>
QuantizedMatrix<ElemType>& MatrixQuantizerGPU<ElemType>::GetTempGPUQuantizedMatrix(size_t numRows, size_t numCols, size_t nBits, bool& newlyAllocated)
{
newlyAllocated = false;
// Check if the existing one is good for our needs
if ((m_tempGPUQuantizedMatrix != nullptr) && (m_tempGPUQuantizedMatrix->GetNumBits() == nBits) && (m_tempGPUQuantizedMatrix->GetNumRows() >= numRows) && (m_tempGPUQuantizedMatrix->GetNumCols() >= numCols))
{
return *m_tempGPUQuantizedMatrix;
}
if (m_tempGPUQuantizedMatrix != nullptr)
{
delete m_tempGPUQuantizedMatrix;
m_tempGPUQuantizedMatrix = nullptr;
}
m_tempGPUQuantizedMatrix = new QuantizedMatrix<ElemType>(numRows, numCols, nBits, (short) this->GetDeviceId());
newlyAllocated = true;
return *m_tempGPUQuantizedMatrix;
}
/////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
///cpubuffer should be page-locked memory allocated, otherwise CUDA will not be efficient (hence we don't use STL)
template <class ElemType>
MatrixQuantizerGPU<ElemType>::MatrixQuantizerGPU(int deviceId, bool useDedicatedComputeStream, bool forceSync /*= false*/)
: MatrixQuantizerImpl<ElemType>(deviceId), m_quantizeCompleteEvent(NULL), m_fetchCompleteEvent(NULL), m_tempMatrixZeroingCompleteEvent(NULL), m_assignCompleteEvent(NULL), m_forceSync(forceSync), m_tempGPUQuantizedMatrix(nullptr), m_quantizeOpIncludedFetch(false)
{
PrepareDevice(this->GetDeviceId());
// events
// Note: Do NOT use cudaEventBlockingSync (which supposedly yields the process)--it will totally break cudaEventSynchronize(), causing it to take 50 or 100 ms randomly.
cudaEventCreateWithFlags(&m_tempMatrixZeroingCompleteEvent, cudaEventDisableTiming) || "cudaEventCreateWithFlags failed";
cudaEventCreateWithFlags(&m_quantizeCompleteEvent, cudaEventDisableTiming) || "cudaEventCreateWithFlags failed";
cudaEventCreateWithFlags(&m_fetchCompleteEvent, cudaEventDisableTiming) || "cudaEventCreateWithFlags failed";
cudaEventCreateWithFlags(&m_assignCompleteEvent, cudaEventDisableTiming) || "cudaEventCreateWithFlags failed";
#pragma warning(disable : 4127)
if (useDedicatedComputeStream && (m_computeStream == NULL))
{
cudaStreamCreateWithFlags(&m_computeStream, cudaStreamNonBlocking) || "cudaStreamCreateWithFlags failed";
}
if (m_fetchStream == NULL)
{
cudaStreamCreateWithFlags(&m_fetchStream, cudaStreamNonBlocking) || "cudaStreamCreateWithFlags failed";
cudaStreamCreateWithFlags(&m_assignStream, cudaStreamNonBlocking) || "cudaStreamCreateWithFlags failed";
}
}
template <class ElemType>
MatrixQuantizerGPU<ElemType>::~MatrixQuantizerGPU()
{
if (nullptr != m_tempGPUQuantizedMatrix)
{
delete m_tempGPUQuantizedMatrix;
m_tempGPUQuantizedMatrix = nullptr;
}
// BUGBUG: we don't destroy our streams (they are static variables); we need a static destructor, I am too lazy now
// TODO: Check for error code and throw if !std::uncaught_exception()
cudaEventDestroy(m_assignCompleteEvent);
cudaEventDestroy(m_fetchCompleteEvent);
cudaEventDestroy(m_quantizeCompleteEvent);
cudaEventDestroy(m_tempMatrixZeroingCompleteEvent);
}
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::QuantizeAsync(const Matrix<ElemType>& inMatrix, const Matrix<ElemType>& inResidual, QuantizedMatrix<ElemType>& outQMatrix, Matrix<ElemType>& outResidual, bool zeroThresholdFor1Bit)
{
// Verify various input matrix parameter's dimensions
assert((inMatrix.GetNumRows() == outQMatrix.GetNumRows()) && (inMatrix.GetNumCols() == outQMatrix.GetNumCols()));
assert((inMatrix.GetNumRows() == inResidual.GetNumRows()) && (inMatrix.GetNumCols() == inResidual.GetNumCols()));
assert((inMatrix.GetNumRows() == outResidual.GetNumRows()) && (inMatrix.GetNumCols() == outResidual.GetNumCols()));
size_t nBits = outQMatrix.GetNumBits();
PrepareDevice(this->GetDeviceId());
if (m_forceSync)
{
Sync();
}
bool GPUMatrixNewlyAllocated = false;
QuantizedMatrix<ElemType>& outQMatrixGPU = (outQMatrix.GetDeviceId() == CPUDEVICE) ? GetTempGPUQuantizedMatrix(outQMatrix.GetNumRows(), outQMatrix.GetNumCols(), nBits, GPUMatrixNewlyAllocated) : outQMatrix;
// If we newly allocated the target GPU matrix then the aysnc zeroing of the matrix is still in procgress on
// the main compute stream. We must synchroniz with the mail compute stream in case the quantization
// compute stream is different from the main compute stream
if (GPUMatrixNewlyAllocated && (GetComputeStream() != GetStream()))
{
cudaEventRecord(m_tempMatrixZeroingCompleteEvent, GetStream()) || "cudaEventRecord failed";
cudaStreamWaitEvent(GetComputeStream(), m_tempMatrixZeroingCompleteEvent, 0 /*flags 'must be 0'*/) || "cudaStreamWaitEvent failed";
}
// Do the quantization on compute sstream and insert event into stream
_QuantizeMatrix<ElemType>(inMatrix.Data(), inResidual.Data(),
inMatrix.GetNumRows(), inMatrix.GetNumCols(),
outQMatrixGPU.Buffer(), nBits, GetComputeStream(),
outResidual.Data(), zeroThresholdFor1Bit);
RecordQuantizeCompleteEvent(GetComputeStream());
// copy from gpu to cpu if needed
m_quantizeOpIncludedFetch = false;
if (outQMatrix.GetDeviceId() == CPUDEVICE)
{
SyncQuantizeCompleEventAndFetchAndRecordFetchCompleteEvent(outQMatrix.Buffer(), outQMatrixGPU.Buffer(), outQMatrixGPU.GetSize());
m_quantizeOpIncludedFetch = true;
}
}
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::WaitQuantizeAsyncDone()
{
PrepareDevice(this->GetDeviceId());
if (m_quantizeOpIncludedFetch)
{
SyncEvent(m_fetchCompleteEvent);
}
else
{
SyncEvent(m_quantizeCompleteEvent);
}
}
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::UnquantizeAsync(QuantizedMatrix<ElemType>& inQMatrix, Matrix<ElemType>& outMatrix, bool add /*= false*/)
{
// The outMatrix should be on the same GPU as m_inMatrix
assert(outMatrix.GetDeviceId() == this->GetDeviceId());
PrepareDevice(this->GetDeviceId());
size_t nBits = inQMatrix.GetNumBits();
// Verify input matrix parameter's dimensions
assert((inQMatrix.GetNumRows() == outMatrix.GetNumRows()) && (inQMatrix.GetNumCols() == outMatrix.GetNumCols()));
bool GPUMatrixNewlyAllocated = false;
QuantizedMatrix<ElemType>& inQMatrixGPU = (inQMatrix.GetDeviceId() == CPUDEVICE) ? GetTempGPUQuantizedMatrix(inQMatrix.GetNumRows(), inQMatrix.GetNumCols(), nBits, GPUMatrixNewlyAllocated) : inQMatrix;
if (inQMatrix.GetDeviceId() == CPUDEVICE)
{
// If the intermediate GPU Matrix was newly allocated, we need to wait for its zeroing to finish
// before assigning the inQMatrix contents
if (GPUMatrixNewlyAllocated)
{
cudaEventRecord(m_tempMatrixZeroingCompleteEvent, GetStream()) || "cudaEventRecord failed";
cudaStreamWaitEvent(GetAssignStream(), m_tempMatrixZeroingCompleteEvent, 0 /*flags 'must be 0'*/) || "cudaStreamWaitEvent failed";
}
// schedule assign to GPU (on transfer stream)
cudaMemcpyAsync(inQMatrixGPU.Buffer(), inQMatrix.Buffer(), inQMatrix.GetSize(), cudaMemcpyHostToDevice, GetAssignStream()) || "cudaMemcpyAsync failed";
// schedule to flag the assign-complete event
cudaEventRecord(m_assignCompleteEvent, GetAssignStream()) || "cudaEventRecord failed"; // for subsequent GPU operation to consume this buffer
if (m_forceSync)
{
SyncStream(GetAssignStream());
}
// let the computing stream wait for the assign complete
SyncAssignCompleteEvent(GetComputeStream());
}
// do the actually unquantization
_UnquantizeMatrix(inQMatrixGPU.Buffer(), inQMatrixGPU.GetSize(),
outMatrix.Data(), outMatrix.GetNumRows(), outMatrix.GetNumCols(),
nBits, add, GetComputeStream());
// Record the event of unquantization
RecordQuantizeCompleteEvent(GetComputeStream());
}
template <class ElemType>
void MatrixQuantizerGPU<ElemType>::WaitUnquantizeAsyncDone()
{
PrepareDevice(this->GetDeviceId());
SyncEvent(m_quantizeCompleteEvent);
}
//explicit
template class MatrixQuantizerGPU<float>;
template class MatrixQuantizerGPU<double>;
GPUMatrixComputeStreamEvent::GPUMatrixComputeStreamEvent(int deviceId)
: MatrixComputeStreamEvent(deviceId)
{
// Note: Do NOT use cudaEventBlockingSync (which supposedly yields the process)--it will totally break cudaEventSynchronize(), causing it to take 50 or 100 ms randomly.
cudaEventCreateWithFlags(&m_mainGPUComputeStreamCUDAEvent, cudaEventDisableTiming) || "cudaEventCreateWithFlags failed";
// Record an event on the main GPU compute stream
cudaEventRecord(m_mainGPUComputeStreamCUDAEvent, GetStream()) || "cudaEventRecord failed";
}
GPUMatrixComputeStreamEvent::~GPUMatrixComputeStreamEvent()
{
// TODO: Check for error code and throw if !std::uncaught_exception()
cudaEventDestroy(m_mainGPUComputeStreamCUDAEvent) || "cudaEventDestroy failed";
}
void GPUMatrixComputeStreamEvent::SynchronizeEvent()
{
cudaEventSynchronize(m_mainGPUComputeStreamCUDAEvent) || "cudaEventSynchronize failed";
}
template <typename ElemType>
void GPUMatrixComputeStreamEvent::SynchronizeQuantizationComputeStreamWithEvent()
{
cudaStreamWaitEvent(MatrixQuantizerGPU<ElemType>::GetComputeStream(), m_mainGPUComputeStreamCUDAEvent, 0 /*flags 'must be 0'*/) || "cudaStreamWaitEvent failed";
}
template <typename ElemType>
void GPUMatrixComputeStreamEvent::SynchronizeDataTransferFetchStreamWithEvent()
{
cudaStreamWaitEvent(GPUDataTransferer<ElemType>::GetFetchStream(), m_mainGPUComputeStreamCUDAEvent, 0 /*flags 'must be 0'*/) || "cudaStreamWaitEvent failed";
}
// Explicit template instantiations
template void GPUMatrixComputeStreamEvent::SynchronizeQuantizationComputeStreamWithEvent<float>();
template void GPUMatrixComputeStreamEvent::SynchronizeQuantizationComputeStreamWithEvent<double>();
template void GPUMatrixComputeStreamEvent::SynchronizeDataTransferFetchStreamWithEvent<float>();
template void GPUMatrixComputeStreamEvent::SynchronizeDataTransferFetchStreamWithEvent<double>();
} } }