wehub-resource-sync
429 lines
17 KiB
Plaintext
429 lines
17 KiB
Plaintext
/* ******************************************************************************
|
|
*
|
|
*
|
|
* This program and the accompanying materials are made available under the
|
|
* terms of the Apache License, Version 2.0 which is available at
|
|
* https://www.apache.org/licenses/LICENSE-2.0.
|
|
*
|
|
* See the NOTICE file distributed with this work for additional
|
|
* information regarding copyright ownership.
|
|
* 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.
|
|
*
|
|
* SPDX-License-Identifier: Apache-2.0
|
|
******************************************************************************/
|
|
|
|
//
|
|
// @author Yurii Shyrma (iuriish@yahoo.com)
|
|
// @author sgazeos@gmail.com
|
|
// @author raver119@gmail.com
|
|
//
|
|
|
|
#include <helpers/ConstantTadHelper.h>
|
|
#include <helpers/PointersManager.h>
|
|
#include <helpers/ShapeUtils.h>
|
|
#include <ops/declarable/helpers/transforms.h>
|
|
|
|
#include "execution/cuda/LaunchDims.h"
|
|
|
|
|
|
namespace sd {
|
|
namespace ops {
|
|
namespace helpers {
|
|
|
|
//////////////////////////////////////////////////////////////////////////
|
|
template <typename T>
|
|
SD_KERNEL static void clipByNormCuda(const void* vClipNorm, const void* vNorm, const LongType* normShapeInfo,
|
|
void* vz, const LongType* zShapeInfo, const LongType* dimensions,
|
|
const LongType dimsLen,
|
|
const bool useAverage) {
|
|
const T clipNorm = *reinterpret_cast<const T*>(vClipNorm);
|
|
const T* norm = reinterpret_cast<const T*>(vNorm);
|
|
T* z = reinterpret_cast<T*>(vz);
|
|
|
|
__shared__ LongType zLen, tadLen, totalThreads;
|
|
__shared__ int zRank, normRank;
|
|
__shared__ const LongType *zShape;
|
|
__shared__ const LongType *zStride;
|
|
__shared__ const LongType *normStride;
|
|
|
|
if (threadIdx.x == 0) {
|
|
zLen = shape::length(zShapeInfo);
|
|
tadLen = zLen / shape::length(normShapeInfo);
|
|
totalThreads = gridDim.x * blockDim.x;
|
|
|
|
// Cache ranks
|
|
zRank = shape::rank(zShapeInfo);
|
|
normRank = shape::rank(normShapeInfo);
|
|
|
|
// Cache shapes and strides
|
|
zShape = shape::shapeOf(zShapeInfo);
|
|
zStride = shape::stride(zShapeInfo);
|
|
normStride = shape::stride(normShapeInfo);
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
LongType zCoords[SD_MAX_RANK], normCoords[SD_MAX_RANK];
|
|
|
|
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
|
|
|
|
for (LongType i = tid; i < zLen; i += totalThreads) {
|
|
INDEX2COORDS(i, zRank, zShape, zCoords);
|
|
|
|
// deduce norm coords
|
|
for (int j = 0; j < dimsLen; ++j) normCoords[j] = zCoords[dimensions[j]];
|
|
|
|
LongType normOffset, zOffset;
|
|
COORDS2INDEX(normRank, normStride, normCoords, normOffset);
|
|
COORDS2INDEX(zRank, zStride, zCoords, zOffset);
|
|
|
|
const T actualNorm = useAverage ? static_cast<T>(norm[normOffset]) / static_cast<T>(tadLen) : static_cast<T>(norm[normOffset]);
|
|
|
|
if (actualNorm > clipNorm) z[zOffset] *= static_cast<T>(clipNorm) / static_cast<T>(actualNorm);
|
|
}
|
|
}
|
|
//////////////////////////////////////////////////////////////////////////
|
|
template <typename T>
|
|
SD_HOST static void clipByNormCudaLauncher(const int blocksPerGrid, const int threadsPerBlock,
|
|
const cudaStream_t* stream, const void* vClipNorm, const void* vNorm,
|
|
const LongType* normShapeInfo, void* vz, const LongType* zShapeInfo,
|
|
const LongType* dimensions, const LongType dimsLen, const bool useAverage) {
|
|
clipByNormCuda<T><<<blocksPerGrid, threadsPerBlock, 512, *stream>>>(vClipNorm, vNorm, normShapeInfo, vz, zShapeInfo,
|
|
dimensions, dimsLen, useAverage);
|
|
sd::DebugHelper::checkGlobalErrorCode("clipByNorm failed");
|
|
|
|
}
|
|
|
|
//////////////////////////////////////////////////////////////////////////
|
|
void clipByNorm(LaunchContext* context, NDArray* input, NDArray* output, const std::vector<LongType>& dims,
|
|
NDArray* clipNorm, const bool isInplace, const bool useAverage) {
|
|
NDArray* z = nullptr;
|
|
|
|
if (isInplace) {
|
|
z = input;
|
|
} else {
|
|
output->assign(input);
|
|
z = output;
|
|
}
|
|
|
|
if (dims.empty()) {
|
|
std::vector<LongType> empty;
|
|
NDArray actualNorm = useAverage ? z->reduceAlongDimension(reduce::Norm2, &empty) / z->lengthOf()
|
|
: z->reduceAlongDimension(reduce::Norm2, &empty);
|
|
|
|
if (actualNorm.e<float>(0) > clipNorm->e<float>(0)) *z *= *clipNorm / actualNorm;
|
|
} else {
|
|
NDArray actualNorms = z->reduceAlongDimension(reduce::Norm2, &dims);
|
|
|
|
std::vector<LongType> *dimsToExclude = ShapeUtils::evalDimsToExclude(z->rankOf(), dims.size(),dims.data());
|
|
|
|
const int threadsPerBlock = SD_MAX_NUM_THREADS / 2;
|
|
const int blocksPerGrid = (z->lengthOf() + threadsPerBlock - 1) / threadsPerBlock;
|
|
|
|
PointersManager manager(context, "clipByNorm");
|
|
|
|
const LongType* dimensions = reinterpret_cast<const LongType*>(
|
|
manager.replicatePointer(dimsToExclude->data(), dimsToExclude->size() * sizeof(LongType)));
|
|
|
|
NDArray::prepareSpecialUse({z}, {z, &actualNorms, clipNorm});
|
|
|
|
|
|
BUILD_SINGLE_SELECTOR(z->dataType(), clipByNormCudaLauncher,
|
|
(blocksPerGrid,
|
|
threadsPerBlock,
|
|
context->getCudaStream(),
|
|
clipNorm->specialBuffer(),
|
|
actualNorms.specialBuffer(),
|
|
actualNorms.specialShapeInfo(),
|
|
z->specialBuffer(),
|
|
z->specialShapeInfo(),
|
|
dimensions,
|
|
dimsToExclude->size(),
|
|
useAverage),
|
|
SD_FLOAT_TYPES);
|
|
NDArray::registerSpecialUse({z}, {z, &actualNorms, clipNorm});
|
|
|
|
manager.synchronize();
|
|
delete dimsToExclude;
|
|
}
|
|
}
|
|
|
|
//////////////////////////////////////////////////////////////////////////
|
|
template <typename T>
|
|
SD_KERNEL static void clipByNormBpCuda(const void* vClipNorm, const void* vx, const LongType* xShapeInfo, // input
|
|
const void* vy, const LongType* yShapeInfo, // gradO
|
|
const void* vNorm, const LongType* normShapeInfo, const void* vSum,
|
|
const LongType* sumShapeInfo, void* vz,
|
|
const LongType* zShapeInfo, // gradI
|
|
const LongType* dimensions, const LongType dimsLen, const bool useAverage) {
|
|
const T clipNorm = *reinterpret_cast<const T*>(vClipNorm);
|
|
const T* norm = reinterpret_cast<const T*>(vNorm);
|
|
const T* sum = reinterpret_cast<const T*>(vSum);
|
|
const T* x = reinterpret_cast<const T*>(vx);
|
|
const T* y = reinterpret_cast<const T*>(vy);
|
|
T* z = reinterpret_cast<T*>(vz);
|
|
|
|
__shared__ LongType zLen, tadLen, totalThreads;
|
|
__shared__ bool sameOffsets;
|
|
__shared__ int zRank, yRank, normRank, sumRank, xRank;
|
|
__shared__ const LongType *zShape;
|
|
__shared__ const LongType *zStride;
|
|
__shared__ const LongType *yStride;
|
|
__shared__ const LongType *normStride;
|
|
__shared__ const LongType *sumStride;
|
|
__shared__ const LongType *xStride;
|
|
|
|
if (threadIdx.x == 0) {
|
|
zLen = shape::length(zShapeInfo);
|
|
tadLen = zLen / shape::length(normShapeInfo);
|
|
totalThreads = gridDim.x * blockDim.x;
|
|
|
|
sameOffsets = shape::haveSameShapeAndStrides(xShapeInfo, yShapeInfo, zShapeInfo);
|
|
|
|
// Cache ranks
|
|
zRank = shape::rank(zShapeInfo);
|
|
yRank = shape::rank(yShapeInfo);
|
|
normRank = shape::rank(normShapeInfo);
|
|
sumRank = shape::rank(sumShapeInfo);
|
|
xRank = shape::rank(xShapeInfo);
|
|
|
|
// Cache shapes and strides
|
|
zShape = shape::shapeOf(zShapeInfo);
|
|
zStride = shape::stride(zShapeInfo);
|
|
yStride = shape::stride(yShapeInfo);
|
|
normStride = shape::stride(normShapeInfo);
|
|
sumStride = shape::stride(sumShapeInfo);
|
|
xStride = shape::stride(xShapeInfo);
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
LongType zCoords[SD_MAX_RANK], normCoords[SD_MAX_RANK];
|
|
|
|
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
|
|
|
|
for (LongType i = tid; i < zLen; i += totalThreads) {
|
|
INDEX2COORDS(i, zRank, zShape, zCoords);
|
|
|
|
LongType zOffset, yOffset;
|
|
COORDS2INDEX(zRank, zStride, zCoords, zOffset);
|
|
if(sameOffsets) {
|
|
yOffset = zOffset;
|
|
} else {
|
|
COORDS2INDEX(yRank, yStride, zCoords, yOffset);
|
|
}
|
|
|
|
// deduce norm coords
|
|
for (int j = 0; j < dimsLen; ++j) normCoords[j] = zCoords[dimensions[j]];
|
|
|
|
LongType normOffset;
|
|
COORDS2INDEX(normRank, normStride, normCoords, normOffset);
|
|
|
|
const T actualNorm = useAverage ? norm[normOffset] / tadLen : norm[normOffset];
|
|
|
|
if (actualNorm > clipNorm) {
|
|
LongType sumOffset, xOffset;
|
|
COORDS2INDEX(sumRank, sumStride, normCoords, sumOffset);
|
|
if(sameOffsets) {
|
|
xOffset = zOffset;
|
|
} else {
|
|
COORDS2INDEX(xRank, xStride, zCoords, xOffset);
|
|
}
|
|
|
|
const T sumVal = sum[sumOffset];
|
|
z[zOffset] = (clipNorm / actualNorm) * y[yOffset] *
|
|
(static_cast<T>(1.f) - (x[xOffset] * sumVal) / (actualNorm * actualNorm));
|
|
} else {
|
|
z[zOffset] = y[yOffset];
|
|
}
|
|
}
|
|
}
|
|
//////////////////////////////////////////////////////////////////////////
|
|
template <typename T>
|
|
void clipByNormBp_(LaunchContext* context, NDArray* input, NDArray* gradO, NDArray* gradI,
|
|
const std::vector<LongType>& dims, NDArray* clipNorm, const bool useAverage) {
|
|
const int rank = input->rankOf();
|
|
|
|
auto actualNorms = input->reduceAlongDimension(reduce::Norm2, &dims);
|
|
|
|
if (actualNorms.lengthOf() == 1) {
|
|
const T norm = useAverage ? actualNorms.e<T>(0) / static_cast<T>(input->lengthOf()) : actualNorms.e<T>(0);
|
|
|
|
auto clipVal = clipNorm->e<T>(0);
|
|
|
|
if (norm > clipVal) {
|
|
const T sum = input->reduceNumber(reduce::Sum).e<T>(0); // reduce to scalar
|
|
const T factor1 = clipVal / norm;
|
|
const T factor2 = static_cast<T>(1.f) / (norm * norm); // 1 / (norm*norm*norm)
|
|
|
|
auto lambda = LAMBDA_TT(x, y, sum, factor1, factor2) {
|
|
return factor1 * y * (static_cast<T>(1.f) - factor2 * x * sum);
|
|
});
|
|
|
|
input->applyPairwiseLambda(gradO, lambda, gradI);
|
|
const_cast<NDArray*>(input)->applyPairwiseLambda(const_cast<NDArray*>(gradO), lambda, gradI);
|
|
} else
|
|
gradI->assign(gradO);
|
|
} else {
|
|
NDArray actualNorms = input->reduceAlongDimension(reduce::Norm2, &dims);
|
|
NDArray sums = input->reduceAlongDimension(reduce::Sum, &dims);
|
|
|
|
std::vector<LongType> *dimsToExclude = ShapeUtils::evalDimsToExclude(gradI->rankOf(), dims.size(),dims.data());
|
|
|
|
|
|
dim3 launchDims = clipDims(gradI->lengthOf());
|
|
PointersManager manager(context, "clipByNormBp");
|
|
|
|
const LongType* dimensions = reinterpret_cast<const LongType*>(
|
|
manager.replicatePointer(dimsToExclude->data(), dimsToExclude->size() * sizeof(LongType)));
|
|
|
|
NDArray::prepareSpecialUse({gradI}, {&actualNorms, &sums, clipNorm, input, gradO});
|
|
clipByNormBpCuda<T><<<launchDims.y, launchDims.x,launchDims.z, *context->getCudaStream()>>>(
|
|
clipNorm->specialBuffer(), input->specialBuffer(), input->specialShapeInfo(), gradO->specialBuffer(),
|
|
gradO->specialShapeInfo(), actualNorms.specialBuffer(), actualNorms.specialShapeInfo(), sums.specialBuffer(),
|
|
sums.specialShapeInfo(), gradI->specialBuffer(), gradI->specialShapeInfo(), dimensions, (LongType)dimsToExclude->size(),
|
|
useAverage);
|
|
sd::DebugHelper::checkGlobalErrorCode("clipByNorm failed");
|
|
|
|
NDArray::registerSpecialUse({gradI}, {&actualNorms, &sums, clipNorm, input, gradO});
|
|
|
|
manager.synchronize();
|
|
delete dimsToExclude;
|
|
}
|
|
}
|
|
BUILD_SINGLE_TEMPLATE( void clipByNormBp_,
|
|
(sd::LaunchContext * context, NDArray* input, NDArray* gradO, NDArray* gradI,
|
|
const std::vector<sd::LongType>& dimensions, NDArray* clipNorm, const bool useAverage),
|
|
SD_FLOAT_TYPES);
|
|
|
|
//////////////////////////////////////////////////////////////////////////
|
|
void clipByNormBp(LaunchContext* context, NDArray* input, NDArray* gradO, NDArray* gradI,
|
|
const std::vector<LongType>& dimensions, NDArray* clipNorm, const bool useAverage) {
|
|
NDArray casted = clipNorm->cast(input->dataType());
|
|
BUILD_SINGLE_SELECTOR(gradI->dataType(), clipByNormBp_,
|
|
(context, &casted, gradO, gradI, dimensions, clipNorm, useAverage), SD_FLOAT_TYPES);
|
|
}
|
|
|
|
template <typename T>
|
|
void clipByGlobalNorm_(LaunchContext* context, std::vector<NDArray*>& inputs, double clipNorm,
|
|
memory::Workspace* workspace, std::vector<NDArray*>& outputs, bool isInplace) {
|
|
T globalNorm = static_cast<T>(0.f);
|
|
|
|
for (auto i = 0; i < inputs.size(); i++) {
|
|
auto input = inputs[i];
|
|
auto l2norm = input->reduceNumber(reduce::Norm2);
|
|
globalNorm += l2norm.e<T>(0) * l2norm.e<T>(0);
|
|
}
|
|
|
|
globalNorm = math::sd_sqrt<T,T>(globalNorm);
|
|
outputs[inputs.size()]->p(0, globalNorm);
|
|
const T factor = static_cast<T>(clipNorm) / globalNorm;
|
|
|
|
for (size_t e = 0; e < inputs.size(); e++) {
|
|
// all-reduce
|
|
auto input = inputs[e];
|
|
auto output = outputs[e];
|
|
|
|
if (static_cast<double>(globalNorm) <= clipNorm) {
|
|
output->assign(input);
|
|
} else {
|
|
auto lambda = LAMBDA_T(_x, factor) { return _x * factor; });
|
|
input->applyLambda(lambda, output);
|
|
}
|
|
}
|
|
}
|
|
|
|
void clipByGlobalNorm(LaunchContext* context, std::vector<NDArray*>& inputs, double clipNorm,
|
|
memory::Workspace* workspace, std::vector<NDArray*>& outputs, bool isInplace) {
|
|
BUILD_SINGLE_SELECTOR(outputs[0]->dataType(), clipByGlobalNorm_,
|
|
(context, inputs, clipNorm, workspace, outputs, isInplace), SD_FLOAT_TYPES);
|
|
}
|
|
|
|
BUILD_SINGLE_TEMPLATE( void clipByGlobalNorm_,
|
|
(sd::LaunchContext * context, std::vector<NDArray*> & inputs, double clipNorm,
|
|
sd::memory::Workspace* workspace, std::vector<NDArray*>& outputs, bool isInplace),
|
|
SD_FLOAT_TYPES);
|
|
|
|
template <typename T>
|
|
static void SD_KERNEL clipByValueKernel(void* input, const LongType* inputShape, void* output,
|
|
const LongType* outputShape, double leftBound, double rightBound) {
|
|
__shared__ T* outputBuf;
|
|
__shared__ T* inputBuf;
|
|
__shared__ LongType length;
|
|
__shared__ LongType inputRank;
|
|
__shared__ LongType outputRank;
|
|
__shared__ LongType* inputShapePtr;
|
|
__shared__ LongType* outputShapePtr;
|
|
__shared__ LongType* inputStridePtr;
|
|
__shared__ LongType* outputStridePtr;
|
|
|
|
if (threadIdx.x == 0) {
|
|
outputBuf = reinterpret_cast<T*>(output);
|
|
inputBuf = reinterpret_cast<T*>(input);
|
|
length = shape::length(inputShape);
|
|
|
|
// Cache shape information
|
|
inputRank = shape::rank(inputShape);
|
|
outputRank = shape::rank(outputShape);
|
|
inputShapePtr = shape::shapeOf(inputShape);
|
|
outputShapePtr = shape::shapeOf(outputShape);
|
|
inputStridePtr = shape::stride(inputShape);
|
|
outputStridePtr = shape::stride(outputShape);
|
|
}
|
|
__syncthreads();
|
|
|
|
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
|
|
const auto step = gridDim.x * blockDim.x;
|
|
|
|
for (LongType e = tid; e < length; e += step) {
|
|
LongType inputCoords[SD_MAX_RANK];
|
|
LongType outputCoords[SD_MAX_RANK];
|
|
LongType inputOffset;
|
|
LongType outputOffset;
|
|
|
|
INDEX2COORDS(e, inputRank, inputShapePtr, inputCoords);
|
|
COORDS2INDEX(inputRank, inputStridePtr, inputCoords, inputOffset);
|
|
INDEX2COORDS(e, outputRank, outputShapePtr, outputCoords);
|
|
COORDS2INDEX(outputRank, outputStridePtr, outputCoords, outputOffset);
|
|
|
|
if (inputBuf[inputOffset] > rightBound)
|
|
outputBuf[outputOffset] = (T)rightBound;
|
|
else if (inputBuf[inputOffset] < leftBound)
|
|
outputBuf[outputOffset] = (T)leftBound;
|
|
else
|
|
outputBuf[outputOffset] = inputBuf[inputOffset];
|
|
}
|
|
}
|
|
|
|
template <typename T>
|
|
static void clipByValue_(LaunchContext* context, NDArray* input, double leftBound, double rightBound,
|
|
NDArray* output) {
|
|
auto stream = context->getCudaStream();
|
|
if (!input->isActualOnDeviceSide()) input->syncToDevice();
|
|
NDArray::prepareSpecialUse({output}, {input});
|
|
dim3 launchDims = getLaunchDims("clip");
|
|
clipByValueKernel<T><<<launchDims.x, launchDims.y, launchDims.z, *stream>>>(input->specialBuffer(), input->specialShapeInfo(),
|
|
output->specialBuffer(), output->specialShapeInfo(), leftBound,
|
|
rightBound);
|
|
sd::DebugHelper::checkGlobalErrorCode("clipByValue failed");
|
|
|
|
NDArray::registerSpecialUse({output}, {input});
|
|
}
|
|
|
|
void clipByValue(LaunchContext* context, NDArray* input, double leftBound, double rightBound, NDArray* output) {
|
|
BUILD_SINGLE_SELECTOR(input->dataType(), clipByValue_, (context, input, leftBound, rightBound, output),
|
|
SD_FLOAT_TYPES);
|
|
}
|
|
|
|
BUILD_SINGLE_TEMPLATE( void clipByValue_, (sd::LaunchContext * context, NDArray* input, double leftBound,
|
|
double rightBound, NDArray* output);
|
|
, SD_FLOAT_TYPES);
|
|
|
|
} // namespace helpers
|
|
} // namespace ops
|
|
} // namespace sd
|