You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@singa.apache.org by wa...@apache.org on 2019/11/20 03:02:43 UTC
[singa] branch master updated: SINGA-490 Optimization on GPU Malloc
and Set Cudastream
This is an automated email from the ASF dual-hosted git repository.
wangwei pushed a commit to branch master
in repository https://gitbox.apache.org/repos/asf/singa.git
The following commit(s) were added to refs/heads/master by this push:
new 1538a26 SINGA-490 Optimization on GPU Malloc and Set Cudastream
new 10ff205 Merge pull request #555 from chrishkchris/SINGA-490_2
1538a26 is described below
commit 1538a26859ea1a3d98a08fe6f31d6c499751f33b
Author: Chris Yeung <ch...@yahoo.com.hk>
AuthorDate: Tue Nov 5 20:21:46 2019 +0800
SINGA-490 Optimization on GPU Malloc and Set Cudastream
---
python/singa/autograd.py | 3 +-
src/core/device/cuda_gpu.cc | 16 ++++++---
src/core/tensor/math_kernel.cu | 76 +++++++++++++++++++++---------------------
3 files changed, 51 insertions(+), 44 deletions(-)
diff --git a/python/singa/autograd.py b/python/singa/autograd.py
index d6b7553..5428065 100644
--- a/python/singa/autograd.py
+++ b/python/singa/autograd.py
@@ -919,7 +919,8 @@ class SoftMaxCrossEntropy(Operation):
def backward(self, dy=1.0):
dx = singa.SoftmaxCrossEntropyBwd(self.p, self.t)
- return singa.DivFloat(dx, float(self.p.shape()[0]))
+ dx /= float(self.p.shape()[0])
+ return dx
def softmax_cross_entropy(x, t):
diff --git a/src/core/device/cuda_gpu.cc b/src/core/device/cuda_gpu.cc
index 623f8c3..1312ece 100644
--- a/src/core/device/cuda_gpu.cc
+++ b/src/core/device/cuda_gpu.cc
@@ -61,21 +61,28 @@ CudaGPU::CudaGPU(int id, std::shared_ptr<DeviceMemPool> pool)
void CudaGPU::Setup() {
lang_ = kCuda;
ctx_.stream = NULL; // use the default sync stream
+
// TODO(wangwei) create one handle for each steam?
+ // Preserse for future use instead of default sync stream, for concurrency
+ // cudaStreamCreate(&ctx_.stream);
+
CUDA_CHECK(cudaSetDevice(id_));
// use curandCreateGeneratorHost for CudaHost device
CURAND_CHECK(
curandCreateGenerator(&ctx_.curand_generator, CURAND_RNG_PSEUDO_DEFAULT));
+ CURAND_CHECK(curandSetStream(ctx_.curand_generator, ctx_.stream));
auto seed = std::chrono::system_clock::now().time_since_epoch().count();
SetRandSeed(seed);
// TODO(wangwei) if one generator per stream, then need diff offset per gen?
CURAND_CHECK(curandSetGeneratorOffset(ctx_.curand_generator, 0));
CUBLAS_CHECK(cublasCreate(&(ctx_.cublas_handle)));
+ CUBLAS_CHECK(cublasSetStream(ctx_.cublas_handle, ctx_.stream));
#ifdef USE_CUDNN
// TODO(wangwei) create one handle for each stream?
auto status = cudnnCreate(&ctx_.cudnn_handle);
CHECK_EQ(status, CUDNN_STATUS_SUCCESS) << cudnnGetErrorString(status);
+ cudnnSetStream(ctx_.cudnn_handle, ctx_.stream);
#endif // USE_CUDNN
}
@@ -88,9 +95,8 @@ void CudaGPU::DoExec(function<void(Context*)>&& fn, int executor) { fn(&ctx_); }
void CudaGPU::CopyToFrom(void* dst, const void* src, size_t nBytes,
CopyDirection direction, Context* ctx) {
- cudaMemcpy(dst, src, nBytes, copyKind[direction]);
- // TODO(wangwei) use async copy
- // cudaMemcpyAsync(dst, src, nBytes,cudaMemcpyDefault, ctx_.stream);
+ // cudaMemcpy(dst, src, nBytes, copyKind[direction]);
+ cudaMemcpyAsync(dst, src, nBytes, copyKind[direction], ctx_.stream);
}
size_t CudaGPU::GetAllocatedMem() {
@@ -108,8 +114,8 @@ void* CudaGPU::Malloc(int size) {
if (size > 0) {
CUDA_CHECK(cudaSetDevice(id_));
pool_->Malloc((void**)&ptr, size);
- // TODO(wangwei) remove the memset.
- CUDA_CHECK(cudaMemset(ptr, 0, size));
+ // Comment out for future analysis: without cnmem
+ // CUDA_CHECK(cudaMemsetAsync(ptr, 0, size, ctx_.stream));
}
return ptr;
}
diff --git a/src/core/tensor/math_kernel.cu b/src/core/tensor/math_kernel.cu
index 2ce6531..bd39c64 100644
--- a/src/core/tensor/math_kernel.cu
+++ b/src/core/tensor/math_kernel.cu
@@ -357,7 +357,7 @@ __global__ void KernelSoftmaxCrossEntropyBwd(const bool int_target, const size_t
} \
} \
void fn(const size_t n, const float *in, float *out, cudaStream_t s) { \
- kernelfn <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out); \
+ kernelfn <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out); \
}
GenUnaryCudaKernel(cos,KernelCos,cosf);
@@ -379,127 +379,127 @@ GenUnaryCudaKernel(atanh,KernelAtanh,atanhf);
// ********************************
void set(const size_t n, const float v, float *out, cudaStream_t s) {
- KernelSet <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, v, out);
+ KernelSet <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, v, out);
}
void abs(const size_t n, const float *in, float *out, cudaStream_t s) {
- KernelAbs <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out);
+ KernelAbs <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out);
}
void sign(const size_t n, const float *in, float *out, cudaStream_t s) {
- KernelSign <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out);
+ KernelSign <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out);
}
void exp(const size_t n, const float *in, float *out, cudaStream_t s) {
- KernelExp <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out);
+ KernelExp <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out);
}
void log(const size_t n, const float *in, float *out, cudaStream_t s) {
- KernelLog <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out);
+ KernelLog <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out);
}
void sqrt(const size_t n, const float *in, float *out, cudaStream_t s) {
- KernelSqrt <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out);
+ KernelSqrt <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out);
}
void square(const size_t n, const float *in, float *out, cudaStream_t s) {
- KernelSquare <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out);
+ KernelSquare <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out);
}
void relu(const size_t n, const float *in, float *out, cudaStream_t s) {
- KernelRelu <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out);
+ KernelRelu <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out);
}
void sigmoid(const size_t n, const float *in, float *out, cudaStream_t s) {
- KernelSigmoid <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out);
+ KernelSigmoid <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out);
}
void softplus(const size_t n, const float *in, float *out, cudaStream_t s) {
- KernelSoftplus <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, out);
+ KernelSoftplus <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, out);
}
void clamp(const size_t n, const float low, const float high, const float *in,
float *out, cudaStream_t s) {
- KernelClamp <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, low, high, in, out);
+ KernelClamp <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, low, high, in, out);
}
void pow(const size_t n, const float *in, const float x, float *out,
cudaStream_t s) {
- KernelPow <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, x, out);
+ KernelPow <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, x, out);
}
void add(const size_t n, const float *in, const float x, float *out,
cudaStream_t s) {
- KernelAdd <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, x, out);
+ KernelAdd <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, x, out);
}
void mult(const size_t n, const float *in, const float x, float *out,
cudaStream_t s) {
- KernelMult <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in, x, out);
+ KernelMult <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in, x, out);
}
void div(const size_t n, const float x, const float *in, float *out,
cudaStream_t s) {
- KernelDiv <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, x, in, out);
+ KernelDiv <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, x, in, out);
}
void threshold(const size_t n, const float x, const float *in, float *out,
cudaStream_t s) {
- KernelThreshold <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, x, in, out);
+ KernelThreshold <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, x, in, out);
}
void gt(const size_t num, const float *in, const float x, float *out,
cudaStream_t s) {
- KernelGT <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF>>> (num, in, x, out);
+ KernelGT <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (num, in, x, out);
}
void gt(const size_t num, const float *in1, const float *in2, float *out,
cudaStream_t s) {
- KernelBGT <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF>>> (num, in1, in2, out);
+ KernelBGT <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (num, in1, in2, out);
}
void ge(const size_t num, const float *in, const float x, float *out,
cudaStream_t s) {
- KernelGE <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF>>> (num, in, x, out);
+ KernelGE <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (num, in, x, out);
}
void ge(const size_t num, const float *in1, const float *in2, float *out,
cudaStream_t s) {
- KernelBGE <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF>>> (num, in1, in2, out);
+ KernelBGE <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (num, in1, in2, out);
}
void lt(const size_t num, const float *in, const float x, float *out,
cudaStream_t s) {
- KernelLT <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF>>> (num, in, x, out);
+ KernelLT <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (num, in, x, out);
}
void lt(const size_t num, const float *in1, const float *in2, float *out,
cudaStream_t s) {
- KernelBLT <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF>>> (num, in1, in2, out);
+ KernelBLT <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (num, in1, in2, out);
}
void le(const size_t num, const float *in, const float x, float *out,
cudaStream_t s) {
- KernelLE <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF>>> (num, in, x, out);
+ KernelLE <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (num, in, x, out);
}
void le(const size_t num, const float *in1, const float *in2, float *out,
cudaStream_t s) {
- KernelBLE <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF>>> (num, in1, in2, out);
+ KernelBLE <<<ceil(num / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (num, in1, in2, out);
}
void pow(const size_t n, const float *in1, const float *in2, float *out,
cudaStream_t s) {
- KernelPow <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in1, in2, out);
+ KernelPow <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in1, in2, out);
}
void add(const size_t n, const float *in1, const float *in2, float *out,
cudaStream_t s) {
- KernelAdd <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in1, in2, out);
+ KernelAdd <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in1, in2, out);
}
void sub(const size_t n, const float *in1, const float *in2, float *out,
cudaStream_t s) {
- KernelSub <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in1, in2, out);
+ KernelSub <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in1, in2, out);
}
void mult(const size_t n, const float *in1, const float *in2, float *out,
cudaStream_t s) {
- KernelMult <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in1, in2, out);
+ KernelMult <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in1, in2, out);
}
void div(const size_t n, const float *in1, const float *in2, float *out,
cudaStream_t s) {
- KernelDiv <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (n, in1, in2, out);
+ KernelDiv <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (n, in1, in2, out);
}
/*
@@ -513,42 +513,42 @@ void sum(const size_t n, const float *in, float *out, cudaStream_t s) {
void ComputeCrossEntropy(const bool int_target, size_t batchsize, const size_t dim, const float *p,
const int *t, float *loss, cudaStream_t stream) {
- KernelComputeCrossEntropy <<<ceil(batchsize / CU1DBLOCKF), CU1DBLOCKF>>>
+ KernelComputeCrossEntropy <<<ceil(batchsize / CU1DBLOCKF), CU1DBLOCKF, 0, stream>>>
(int_target, batchsize, dim, p, t, loss);
}
void SoftmaxCrossEntropyBwd(const bool int_target, size_t batchsize, const size_t dim, const float *p,
const int *t, float *grad, cudaStream_t stream) {
- KernelSoftmaxCrossEntropyBwd <<<ceil(batchsize / CU1DBLOCKF), CU1DBLOCKF>>>
+ KernelSoftmaxCrossEntropyBwd <<<ceil(batchsize / CU1DBLOCKF), CU1DBLOCKF, 0, stream>>>
(int_target, batchsize, dim, p, t, grad);
}
void RowMax(const size_t nrow, const size_t ncol, const float *inPtr,
float *outPtr, cudaStream_t stream) {
- KernelRowMax <<<ceil(nrow / CU1DBLOCKF), CU1DBLOCKF>>>(nrow, ncol, inPtr, outPtr);
+ KernelRowMax <<<ceil(nrow / CU1DBLOCKF), CU1DBLOCKF, 0, stream>>>(nrow, ncol, inPtr, outPtr);
}
/*
void square_grad(int n, const float *in, float *out, cudaStream_t s) {
- kernel_square_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (in, out, n);
+ kernel_square_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (in, out, n);
}
void tanh_grad(int n, const float *in, float *out, cudaStream_t s) {
- kernel_tanh_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (in, out, n);
+ kernel_tanh_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (in, out, n);
}
void relu_grad(int n, const float *in, float *out, cudaStream_t s) {
- kernel_relu_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (in, out, n);
+ kernel_relu_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (in, out, n);
}
void sigmoid_grad(int n, const float *in, float *out, cudaStream_t s) {
- kernel_sigmoid_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (in, out, n);
+ kernel_sigmoid_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (in, out, n);
}
void softplus_grad(int n, const float *in, float *out, cudaStream_t s) {
- kernel_softplus_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (in, out, n);
+ kernel_softplus_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF, 0, s>>> (in, out, n);
}