From 352ae7e65b6a74bcd768aa88b866a44c713284c8 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Wed, 26 Oct 2016 15:35:44 +0000
Subject: [PATCH] ADAM
---
src/blas_kernels.cu | 37 ++++++++++++++++++++++++++++---------
1 files changed, 28 insertions(+), 9 deletions(-)
diff --git a/src/blas_kernels.cu b/src/blas_kernels.cu
index 59ec005..684e66d 100644
--- a/src/blas_kernels.cu
+++ b/src/blas_kernels.cu
@@ -140,6 +140,21 @@
}
+__global__ void adam_kernel(int N, float *x, float *m, float *v, float B1, float B2, float rate, float eps, int t)
+{
+ int index = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if (index >= N) return;
+
+ x[index] = x[index] - (rate * sqrt(1.-pow(B2, t)) / (1.-pow(B1, t)) * m[index] / (sqrt(v[index]) + eps));
+ //if(index == 0) printf("%f %f %f %f\n", m[index], v[index], (rate * sqrt(1.-pow(B2, t)) / (1.-pow(B1, t)) * m[index] / (sqrt(v[index]) + eps)));
+}
+
+extern "C" void adam_gpu(int n, float *x, float *m, float *v, float B1, float B2, float rate, float eps, int t)
+{
+ adam_kernel<<<cuda_gridsize(n), BLOCK>>>(n, x, m, v, B1, B2, rate, eps, t);
+ check_error(cudaPeekAtLastError());
+}
+
__global__ void normalize_kernel(int N, float *x, float *mean, float *variance, int batch, int filters, int spatial)
{
int index = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
@@ -693,31 +708,35 @@
}
-__global__ void softmax_kernel(int n, int batch, float *input, float temp, float *output)
+__device__ void softmax_device(int n, float *input, float temp, float *output)
{
- int b = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- if(b >= batch) return;
-
int i;
float sum = 0;
float largest = -INFINITY;
for(i = 0; i < n; ++i){
- int val = input[i+b*n];
+ int val = input[i];
largest = (val>largest) ? val : largest;
}
for(i = 0; i < n; ++i){
- sum += exp(input[i+b*n]/temp-largest/temp);
+ sum += exp(input[i]/temp-largest/temp);
}
sum = (sum != 0) ? largest/temp+log(sum) : largest-100;
for(i = 0; i < n; ++i){
- output[i+b*n] = exp(input[i+b*n]/temp-sum);
+ output[i] = exp(input[i]/temp-sum);
}
}
-extern "C" void softmax_gpu(float *input, int n, int groups, float temp, float *output, cudaStream_t stream)
+__global__ void softmax_kernel(int n, int offset, int batch, float *input, float temp, float *output)
+{
+ int b = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if(b >= batch) return;
+ softmax_device(n, input + b*offset, temp, output + b*offset);
+}
+
+extern "C" void softmax_gpu(float *input, int n, int offset, int groups, float temp, float *output)
{
int inputs = n;
int batch = groups;
- softmax_kernel<<<cuda_gridsize(batch), BLOCK, 0, stream>>>(inputs, batch, input, temp, output);
+ softmax_kernel<<<cuda_gridsize(batch), BLOCK>>>(inputs, offset, batch, input, temp, output);
check_error(cudaPeekAtLastError());
}
--
Gitblit v1.10.0