Luo Tao
8 years ago
commit
5b81f5da74
@ -0,0 +1,48 @@
|
||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
||||
|
||||
Licensed under the Apache License, Version 2.0 (the "License");
|
||||
you may not use this file except in compliance with the License.
|
||||
You may obtain a copy of the License at
|
||||
|
||||
http://www.apache.org/licenses/LICENSE-2.0
|
||||
|
||||
Unless required by applicable law or agreed to in writing, software
|
||||
distributed under the License is distributed on an "AS IS" BASIS,
|
||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
See the License for the specific language governing permissions and
|
||||
limitations under the License. */
|
||||
|
||||
#ifndef HL_BATCH_NORM_H_
|
||||
#define HL_BATCH_NORM_H_
|
||||
|
||||
#include "hl_base.h"
|
||||
|
||||
/**
|
||||
* @brief batch norm inferece.
|
||||
*
|
||||
* @param[in] input input data.
|
||||
* @param[out] output output data.
|
||||
* @param[in] scale batch normalization scale parameter (in original
|
||||
* paper scale is referred to as gamma).
|
||||
* @param[in] bias batch normalization bias parameter (in original
|
||||
* paper scale is referred to as beta).
|
||||
* @param[in] estimatedMean
|
||||
* @param[in] estimatedVar The moving mean and variance
|
||||
* accumulated during the training phase are passed
|
||||
* as inputs here.
|
||||
* @param[in] epsilon Epsilon value used in the batch
|
||||
* normalization formula.
|
||||
*/
|
||||
extern void hl_batch_norm_cuda_inference(const real* input,
|
||||
real* output,
|
||||
const real* scale,
|
||||
const real* bias,
|
||||
const real* estimatedMean,
|
||||
const real* estimatedVar,
|
||||
const double epsilon,
|
||||
size_t batchSize,
|
||||
size_t channel,
|
||||
size_t height,
|
||||
size_t width);
|
||||
|
||||
#endif // HL_BATCH_NORM_H_
|
||||
@ -0,0 +1,66 @@
|
||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
||||
|
||||
Licensed under the Apache License, Version 2.0 (the "License");
|
||||
you may not use this file except in compliance with the License.
|
||||
You may obtain a copy of the License at
|
||||
|
||||
http://www.apache.org/licenses/LICENSE-2.0
|
||||
|
||||
Unless required by applicable law or agreed to in writing, software
|
||||
distributed under the License is distributed on an "AS IS" BASIS,
|
||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
See the License for the specific language governing permissions and
|
||||
limitations under the License. */
|
||||
|
||||
#include "hl_batch_norm.h"
|
||||
|
||||
__global__ void batchNormInference(real* output,
|
||||
const real* input,
|
||||
const real* scale,
|
||||
const real* bias,
|
||||
const real* estimatedMean,
|
||||
const real* estimatedVar,
|
||||
const double epsilon,
|
||||
size_t batchSize,
|
||||
size_t channel,
|
||||
size_t height,
|
||||
size_t width) {
|
||||
const int tid = threadIdx.x;
|
||||
const int num = channel * height * width;
|
||||
const int batch = blockIdx.x;
|
||||
for (int i = tid; i < num; i += blockDim.x) {
|
||||
const int c = i / (height * width);
|
||||
const int id = batch * num + i;
|
||||
real val = input[id] - estimatedMean[c];
|
||||
val /= sqrt(estimatedVar[c] + epsilon);
|
||||
val *= scale[c];
|
||||
val += bias[c];
|
||||
output[id] = val;
|
||||
}
|
||||
}
|
||||
|
||||
void hl_batch_norm_cuda_inference(const real* input,
|
||||
real* output,
|
||||
const real* scale,
|
||||
const real* bias,
|
||||
const real* estimatedMean,
|
||||
const real* estimatedVar,
|
||||
const double epsilon,
|
||||
size_t batchSize,
|
||||
size_t channel,
|
||||
size_t height,
|
||||
size_t width) {
|
||||
batchNormInference<<<batchSize, 256, 0, STREAM_DEFAULT>>>(output,
|
||||
input,
|
||||
scale,
|
||||
bias,
|
||||
estimatedMean,
|
||||
estimatedVar,
|
||||
epsilon,
|
||||
batchSize,
|
||||
channel,
|
||||
height,
|
||||
width);
|
||||
|
||||
CHECK_SYNC("hl_batch_norm_cuda_inference failed!");
|
||||
}
|
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue