def adagradKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *param, const float *grad, float *h, float learnRate, float epsilon",
"h[i] += grad[i] * grad[i];"
"param[i] += learnRate * grad[i] / (sqrt(h[i]) + epsilon)",
"adagradKer")
def classicMomSGDKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *param, const float *grad, float *mom, float learnRate, float momRate",
"""
mom[i] = momRate * mom[i] + learnRate * grad[i];
param[i] += mom[i];
""", "classicMomSGDKer")
def rmspropKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *param, const float *grad, float *ms, float learnRate, float factor, float epsilon",
"""
ms[i] = factor * ms[i] + (1.0f - factor) * grad[i] * grad[i];
param[i] += learnRate * grad[i] / (sqrt(ms[i]) + epsilon);
""", "rmspropKer")
def nesterovMomSGDKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *param, const float *grad, float *mom, float learnRate, float momRate",
"""
float m = mom[i];
mom[i] = momRate * m + learnRate * grad[i];
param[i] += momRate * momRate * m + (1.0f + momRate) * learnRate * grad[i];
""", "nesterovMomSGDKer")
def adamKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *param, const float *grad, float *mg, float *ms, float learnRate, float fix1, float fix2, float epsilon",
"""
mg[i] += fix1 * (grad[i] - mg[i]);
ms[i] += fix2 * (grad[i] * grad[i] - ms[i]);
param[i] += learnRate * mg[i] / (sqrt(ms[i]) + epsilon);
""", "adamKer")
def adadeltaKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *param, const float *grad, float *msg, float *msdx, float rho, float epsilon",
"""
msg[i] += (1.0f - rho) * (grad[i] * grad[i] - msg[i]);
float dx = sqrt((msdx[i] + epsilon) / (msg[i] + epsilon)) * grad[i];
msdx[i] += (1.0f - rho) * (dx * dx - msdx[i]);
param[i] += dx;
""", "adadeltaKer")
def rmspropGravesKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"""
float *param, const float *grad, float *mg, float *ms, float *delta,
float learnRate, float alpha, float momRate, float epsilon
""", """
ms[i] = alpha * ms[i] + (1.0f - alpha) * grad[i] * grad[i];
mg[i] = alpha * mg[i] + (1.0f - alpha) * grad[i];
delta[i] = momRate * delta[i] + learnRate * grad[i] / sqrt(ms[i] - mg[i] * mg[i] + epsilon);
param[i] += delta[i];
""", "rmspropGravesKer")
def smorms3Ker(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *param, const float *grad, float *mem, float *g, float *g2, float learnRate, float epsilon",
"""
float memi = mem[i], r = 1.0f / (memi + 1.0f), gi = g[i], g2i = g2[i];
gi = (1.0f - r) * gi + r * grad[i];
g2i = (1.0f - r) * g2i + r * grad[i] * grad[i];
float x = gi * gi / (g2i + epsilon);
param[i] += grad[i] * min(learnRate, x) / (sqrt(g2i) + epsilon);
mem[i] = 1.0f + memi * (1.0f - x);
g[i] = gi;
g2[i] = g2i;
""", "smorms3Ker")
def linearKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *outdata, const float *indata, float a, float b",
"outdata[i] = a * indata[i] + b", "linearKer")
def mulKer(dtype):
assert dtype == np.float32
return ElementwiseKernel("float *outdata, const float *a, const float *b",
"outdata[i] = a[i] * b[i]", "mulKer")
def addKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *outdata, const float *a, float alpha, const float *b, float beta",
"outdata[i] = alpha * a[i] + beta * b[i]", "addKer")
def dropout2dKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *outdata, const float *indata, const float *uni, float p, int mapsize",
"outdata[i] = indata[i] * (uni[i / mapsize] < p) / p", "dropout2dKer")
def reluKer(dtype):
assert dtype == np.float32
return ElementwiseKernel("float *outdata, const float *indata",
"outdata[i] = indata[i] * (indata[i] > 0.0f)",
"reluKer")
def leakyReluKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *outdata, const float *indata, float a",
"outdata[i] = indata[i] * ((indata[i] > 0.0f) + a * (indata[i] <= 0.0f))",
"leakyReluKer")
def softPlusDerKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *ingrad, const float *outgrad, const float *outdata",
"ingrad[i] = outgrad[i] * (1.0f - exp(-outdata[i]))", "softPlusDerKer")
def clipKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *outdata, const float *indata, float a, float b",
"outdata[i] = min(b, max(a, indata[i]))", "clipKer")
def toVectorAddVectorKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *outdata, const float *indata, float alpha",
"outdata[i] += indata[i] * alpha", "toVectorAddVectorKer")
def tanhKer(dtype):
assert dtype == np.float32
return ElementwiseKernel("float *outdata, const float *indata",
"outdata[i] = tanh(indata[i])", "tanhKer")
def eluDerKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *ingrad, const float *outgrad, const float *outdata, float a",
"ingrad[i] = outgrad[i] * ((outdata[i] > 0.0f) + (outdata[i] + a) * (outdata[i] <= 0.0f))",
"eluDerKer")
def tanhDerKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *ingrad, const float *outgrad, const float *outdata",
"ingrad[i] = outgrad[i] * (1.0f - outdata[i] * outdata[i])",
"tanhDerKer")
def softPlusKer(dtype):
assert dtype == np.float32
return ElementwiseKernel("float *outdata, const float *indata",
"outdata[i] = log(1.0f + exp(indata[i]))",
"softPlusKer")
def reluDerKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *ingrad, const float *outgrad, const float *outdata",
"ingrad[i] = outgrad[i] * (outdata[i] > 0.0f)", "reluDerKer")
def clipDerKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *ingrad, const float *outgrad, const float *outdata, float a, float b",
"ingrad[i] = outgrad[i] * (outdata[i] > a && outdata[i] < b);",
"clipDerKer")
def eluKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *outdata, const float *indata, float a",
"outdata[i] = indata[i] * (indata[i] > 0.0f) + a * (exp(indata[i]) - 1.0f) * (indata[i] <= 0.0f)",
"eluKer")
def sigmoidKer(dtype):
assert dtype == np.float32
return ElementwiseKernel("float *outdata, const float *indata",
"outdata[i] = 1.0f / (1.0f + exp(-indata[i]))",
"sigmoidKer")
from string import Template
import numpy as np
from PuzzleLib.OpenCL.Driver import Driver
from PuzzleLib.OpenCL.Kernels.Templates import ElementwiseKernel, ReductionKernel
from PuzzleLib.OpenCL.Kernels.Utils import nthreads, roundUp, atomicAddTmpl
from PuzzleLib.OpenCL.Utils import context, queue, memoryPool as memPool
from PuzzleLib.OpenCL.Wrappers.MIOpen import softmax2d
bceKer = ElementwiseKernel(
"const float *scores, const int *labels, float *totalError, float *grad, int numsamples, int spatialDim",
"""
float prob = 1.0f / (1.0f + exp(-scores[i]));
float error = labels[i] == 1 ? -log(prob) : -log(1.0f - prob);
atomicAddCAS(totalError, error / spatialDim);
grad[i] = ((labels[i] == 1) - prob) / numsamples / spatialDim;
""",
"bceKer",
preamble=atomicAddTmpl)
hingeKer = ElementwiseKernel(
"const float *scores, const int *labels, float *totalError, float *grad, int numsamples, int numcases",
"""
float score = scores[i];
int label = labels[i];
float error = max(0.0f, 1.0f - score * label) / numcases;
atomicAddCAS(totalError, error);
grad[i] = score * label < 1.0f ? (float)label / numsamples / numcases : 0.0f;
""",
"hingeKer",
def mulKer(dtype):
assert dtype == np.float32
return ElementwiseKernel("float *outdata, const float *a, const float *b",
"outdata[i] = a[i] * b[i]", "mulKer")
@memoize
def linearKer(dtype):
assert dtype == np.float32
return ElementwiseKernel(
"float *outdata, const float *indata, float a, float b",
"outdata[i] = a * indata[i] + b", "linearKer")
rbmKer = ElementwiseKernel(
"float *outdata, const float *indata, const float *uni",
"float act = 1.0f / (1.0f + exp(-indata[i]));"
"outdata[i] = (uni[i] < act)", "rbmKer")
absKer = ElementwiseKernel("float *outdata, const float *indata",
"outdata[i] = fabs(indata[i])", "absKer")
weightDecayKer = ElementwiseKernel(
"float *grad, const float *param, float rate",
"grad[i] -= rate * param[i]", "weightDecayKer")
l1penaltyKer = ElementwiseKernel(
"float *outgrad, const float *ingrad, const float *data, float a",
"outgrad[i] = ingrad[i] - a * ((0.0f <= data[i]) - (data[i] < 0.0f))",
"l1penaltyKer")
l1gradKer = ElementwiseKernel(
