872 lines
22 KiB
Plaintext
872 lines
22 KiB
Plaintext
#include <cstdint>
|
|
#include <stdint.h>
|
|
#include <time.h>
|
|
#include <random>
|
|
#include "error.cuh"
|
|
extern "C" {
|
|
/* 创建一个设备数组 */
|
|
static void* _crtDevArr(const size_t m)
|
|
{
|
|
void *d;
|
|
CHECK(cudaMalloc(&d, m));
|
|
return d;
|
|
}
|
|
|
|
/* 释放一个设备数组 */
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_Free(void* d)
|
|
#else
|
|
void Cu4cj_Free(void* d)
|
|
#endif
|
|
{
|
|
CHECK(cudaFree(d));
|
|
}
|
|
|
|
|
|
/* 创建一个单精度浮点数组 输入数组的长度（不初始化） */
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* Cu4cj_Orig(uint64_t n)
|
|
#else
|
|
void* Cu4cj_Orig(int64_t n)
|
|
#endif
|
|
{
|
|
return _crtDevArr(n * sizeof(float));
|
|
}
|
|
|
|
|
|
/*
|
|
* 从主机内存拷贝数据到设备内存
|
|
* FIXME: 是否需要做成输入目的地址
|
|
*/
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) float* Cu4cj_ToDev(float* h, uint64_t n)
|
|
#else
|
|
float* Cu4cj_ToDev(float* h, int64_t n)
|
|
#endif
|
|
{
|
|
float* d = (float*)Cu4cj_Orig(n);
|
|
CHECK(cudaMemcpy(d, h, n * sizeof(float), cudaMemcpyHostToDevice));
|
|
return d;
|
|
}
|
|
|
|
|
|
/* 从设备内存拷贝数据到主机内存 */
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_ToHost(float* d, float* h, const uint64_t n)
|
|
#else
|
|
void Cu4cj_ToHost(float* d, float* h, int64_t n)
|
|
#endif
|
|
{
|
|
CHECK(cudaMemcpy(h, d, n * sizeof(float), cudaMemcpyDeviceToHost));
|
|
}
|
|
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_Copy(float* d, float* h, const uint64_t n)
|
|
#else
|
|
void Cu4cj_Copy(float* d, float* h, int64_t n)
|
|
#endif
|
|
{
|
|
CHECK(cudaMemcpy(h, d, n * sizeof(float), cudaMemcpyDeviceToHost));
|
|
}
|
|
|
|
/* 创建一个单精度浮点数组并将其初始化为零
|
|
* FIXME: 当n == 10000 时进程被杀死 */
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) float* Cu4cj_Zero(const uint64_t n)
|
|
#else
|
|
float* Cu4cj_Zero(int64_t n)
|
|
#endif
|
|
{
|
|
float* d = (float*)Cu4cj_Orig(n);
|
|
CHECK(cudaMemset(d, 0, n*sizeof(float)));
|
|
return d;
|
|
}
|
|
|
|
|
|
/* 正态分布初始化（适用于神经网络权重） */
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) float* Cu4cj_Normal(uint64_t n, uint64_t fan_in, uint64_t fan_out)
|
|
#else
|
|
float* Cu4cj_Normal(int64_t n, int64_t fan_in, int64_t fan_out)
|
|
#endif
|
|
{
|
|
// 计算 Xavier/Glorot 初始化的标准差
|
|
float std = sqrt(2.0f / (fan_in + fan_out));
|
|
|
|
// 先在主机上分配内存并生成正态分布随机数
|
|
float* h = (float*)malloc(n * sizeof(float));
|
|
|
|
/* 初始化随机数种子 */
|
|
srand(time(NULL));
|
|
|
|
for (size_t i = 0; i < n; i++) {
|
|
// 使用 Box-Muller 变换生成正态分布随机数
|
|
float u1 = (float)rand() / RAND_MAX;
|
|
float u2 = (float)rand() / RAND_MAX;
|
|
float z = sqrt(-2.0f * log(u1)) * cos(2.0f * 3.1415926535f * u2);
|
|
h[i] = z * std;
|
|
}
|
|
|
|
// 在设备上分配内存
|
|
float* d = (float*)Cu4cj_Orig(n);
|
|
|
|
// 将主机内存中的数据拷贝到设备内存
|
|
CHECK(cudaMemcpy(d, h, n * sizeof(float), cudaMemcpyHostToDevice));
|
|
|
|
// 释放主机内存
|
|
free(h);
|
|
|
|
return d;
|
|
}
|
|
|
|
|
|
|
|
|
|
// 向量加法
|
|
static void __global__ _devAdd(float* d1, float* d2, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] + d2[i];
|
|
}
|
|
}
|
|
|
|
// 向量加法
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_Add(float* d1, float* d2, float* d_r, uint64_t n)
|
|
#else
|
|
void Cu4cj_Add(float* d1, float* d2, float* d_r, int64_t n)
|
|
#endif
|
|
{
|
|
_devAdd<<<(n + 255) / 256, 256>>>(d1, d2, d_r, n);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
}
|
|
|
|
// 向量减法
|
|
static void __global__ _devSub(float* d1, float* d2, float* r, int64_t n)
|
|
{
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] - d2[i];
|
|
}
|
|
}
|
|
|
|
// 向量减法
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_Sub(float* d1, float* d2, float* d_r, uint64_t n)
|
|
#else
|
|
void Cu4cj_Sub(float* d1, float* d2, float* d_r, int64_t n)
|
|
#endif
|
|
{
|
|
_devSub<<<(n + 255) / 256, 256>>>(d1, d2, d_r, n);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
}
|
|
|
|
// 向量依次相乘
|
|
static void __global__ _devMul(float* d1, float* d2, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] * d2[i];
|
|
}
|
|
}
|
|
|
|
// 向量依次相乘
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_Mul(float* d1, float* d2, float* d_r, uint64_t n)
|
|
#else
|
|
void Cu4cj_Mul(float* d1, float* d2, float* d_r, int64_t n)
|
|
#endif
|
|
{
|
|
_devMul<<<(n + 255) / 256, 256>>>(d1, d2, d_r, n);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
}
|
|
|
|
// 向量依次相除
|
|
static void __global__ _devDiv(float* d1, float* d2, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] / d2[i];
|
|
}
|
|
}
|
|
|
|
// 向量依次相除
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_Div(float* d1, float* d2, float* d_r, uint64_t n)
|
|
#else
|
|
void Cu4cj_Div(float* d1, float* d2, float* d_r, int64_t n)
|
|
#endif
|
|
{
|
|
_devDiv<<<(n + 255) / 256, 256>>>(d1, d2, d_r, n);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
}
|
|
|
|
// 向量加法广播
|
|
static void __global__ _devAddBroadcast(float* d1, float a, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] + a;
|
|
}
|
|
}
|
|
|
|
// 加法广播
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_AddBroadcast(float* d1, float a, float* d_r, uint64_t n)
|
|
#else
|
|
void Cu4cj_AddBroadcast(float* d1, float a, float* d_r, int64_t n)
|
|
#endif
|
|
{
|
|
_devAddBroadcast<<<(n + 255) / 256, 255>>>(d1, a, d_r, n);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
}
|
|
|
|
// 减法广播
|
|
static void __global__ _devSubBroadcast(float* d1, float a, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] - a;
|
|
}
|
|
}
|
|
|
|
// 减法广播
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_SubBroadcast(float* d1, float a, float* d_r, uint64_t n)
|
|
#else
|
|
void Cu4cj_SubBroadcast(float* d1, float a, float* d_r, int64_t n)
|
|
#endif
|
|
{
|
|
_devSubBroadcast<<<(n + 255) / 256, 256>>>(d1, a, d_r, n);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
}
|
|
|
|
// 乘法广播
|
|
static void __global__ _devMulBroadcast(float* d1, float a, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] * a;
|
|
}
|
|
}
|
|
|
|
// 乘法广播
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_MulBroadcast(float* d1, float a, float* d_r, uint64_t n)
|
|
#else
|
|
void Cu4cj_MulBroadcast(float* d1, float a, float* d_r, int64_t n)
|
|
#endif
|
|
{
|
|
_devMulBroadcast<<<(n + 255) / 256, 256>>>(d1, a, d_r, n);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
}
|
|
|
|
// 除法广播
|
|
static void __global__ _devDivBroadcast(float* d1, float a, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] / a;
|
|
}
|
|
}
|
|
|
|
// 除法广播
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_DivBroadcast(float* d1, float a, float* d_r, uint64_t n)
|
|
#else
|
|
void Cu4cj_DivBroadcast(float* d1, float a, float* d_r, int64_t n)
|
|
#endif
|
|
{
|
|
_devDivBroadcast<<<(n + 255) / 256, 256>>>(d1, a, d_r, n);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
}
|
|
|
|
// 矩阵乘法
|
|
static void __global__ _devMatDot(float* d1, float* d2, float* r, int64_t m, int64_t n, int64_t k) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
int j = blockIdx.y * blockDim.y + threadIdx.y;
|
|
if (i < m && j < k) {
|
|
float sum = 0.0f;
|
|
for (int p = 0; p < n; p++) {
|
|
sum += d1[i * n + p] * d2[p * k + j];
|
|
}
|
|
r[i * k + j] = sum;
|
|
}
|
|
}
|
|
|
|
// 矩阵乘法
|
|
// m、n、k 分别为矩阵 A、B、C 的行数、列数
|
|
// 其中 n 为矩阵 A 的列数，也为矩阵 B 的行数
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void Cu4cj_MatDot(float* d1, float* d2, float* d_r, uint64_t m, uint64_t n, uint64_t k)
|
|
#else
|
|
void Cu4cj_MatDot(float* d1, float* d2, float* d_r, int64_t m, int64_t n, int64_t k)
|
|
#endif
|
|
{
|
|
dim3 blockDim(16, 16);
|
|
dim3 gridDim((m + blockDim.x - 1) / blockDim.x, (k + blockDim.y - 1) / blockDim.y);
|
|
_devMatDot<<<gridDim, blockDim>>>(d1, d2, d_r, m, n, k);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
}
|
|
|
|
// 矩阵转置
|
|
static void __global__ _devTransport(float* d1, float* d2, int64_t m, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
int j = blockIdx.y * blockDim.y + threadIdx.y;
|
|
if (i < m && j < n) {
|
|
d2[j * m + i] = d1[i * n + j];
|
|
}
|
|
}
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) float* Cu4cj_Transport(float *m, uint64_t r, uint64_t c)
|
|
#else
|
|
float* Cu4cj_Transport(float *m, uint64_t r, uint64_t c)
|
|
#endif
|
|
{
|
|
float *d_r = (float*)Cu4cj_Zero(c * r * sizeof(float));
|
|
_devTransport<<<1, c * r>>>(m, d_r, c, r);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
return d_r;
|
|
}
|
|
|
|
|
|
// Sigmoid 激活函数，输入为矩阵，输出为矩阵
|
|
static void __global__ _devSigmoid(float* d1, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = 1.0f / (1.0f + exp(-d1[i]));
|
|
}
|
|
}
|
|
|
|
// Sigmoid 导数
|
|
static void __global__ _devSigmoidDerivative(float* d1, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] * (1.0f - d1[i]);
|
|
}
|
|
}
|
|
|
|
// ReLU
|
|
static void __global__ _devReLU(float* d1, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] > 0.0f ? d1[i] : 0.0f;
|
|
}
|
|
}
|
|
|
|
// ReLU 导数
|
|
static void __global__ _devReLUDerivative(float* d1, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] > 0.0f ? 1.0f : 0.0f;
|
|
}
|
|
}
|
|
|
|
// Leaky_ReLU
|
|
static void __global__ _devLeakyReLU(float* d1, float* r, int64_t n, float a) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] > 0.0f ? d1[i] : a * d1[i];
|
|
}
|
|
}
|
|
|
|
// Leaky_ReLU 导数
|
|
static void __global__ _devLeakyReLU_Derivative(float* d1, float* r, int64_t n, float a) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = d1[i] > 0.0f ? 1.0f : a;
|
|
}
|
|
}
|
|
|
|
// Tanh 激活函数，输入为矩阵，输出为矩阵
|
|
static void __global__ _devTanh(float* d1, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = (exp(d1[i]) - exp(-d1[i])) / (exp(d1[i]) + exp(-d1[i]));
|
|
}
|
|
}
|
|
|
|
// Tanh 导数
|
|
static void __global__ _devTanhDerivative(float* d1, float* r, int64_t n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) {
|
|
r[i] = 1.0f - d1[i] * d1[i];
|
|
}
|
|
}
|
|
|
|
}
|
|
|
|
namespace qing {
|
|
class Matrix {
|
|
public:
|
|
Matrix(uint64_t r, uint64_t c) {
|
|
this->r = r;
|
|
this->c = c;
|
|
this->d = (float*)Cu4cj_Zero(r * c * sizeof(float));
|
|
}
|
|
Matrix(uint64_t r, uint64_t c, float* d) {
|
|
this->r = r;
|
|
this->c = c;
|
|
this->d = d;
|
|
}
|
|
Matrix(const Matrix& m) {
|
|
this->r = m.getR();
|
|
this->c = m.getC();
|
|
this->d = (float*)Cu4cj_Zero(m.getR() * m.getC() * sizeof(float));
|
|
Cu4cj_Copy(m.getD(), d, m.getR() * m.getC() * sizeof(float));
|
|
}
|
|
~Matrix() {
|
|
Cu4cj_Free(d);
|
|
}
|
|
uint64_t getR() const {
|
|
return r;
|
|
}
|
|
uint64_t getC() const {
|
|
return c;
|
|
}
|
|
float* getD() const {
|
|
return d;
|
|
}
|
|
float get(uint64_t r, uint64_t c) const {
|
|
float h;
|
|
CHECK(cudaMemcpy(&h, &d[c * this->r + r], sizeof(float), cudaMemcpyDeviceToHost));
|
|
return h;
|
|
}
|
|
void set(uint64_t r, uint64_t c, float v) {
|
|
CHECK(cudaMemcpy(&d[c * this->r + r], &v, sizeof(float), cudaMemcpyHostToDevice));
|
|
}
|
|
bool hasSameShape(const Matrix& m) {
|
|
return r == m.getR() && c == m.getC();
|
|
}
|
|
bool hasDotShape(const Matrix& m) {
|
|
return c == m.getR();
|
|
}
|
|
Matrix operator+(const Matrix& m) {
|
|
float *d_r = (float*)Cu4cj_Zero(r * c * sizeof(float));
|
|
Cu4cj_Add(d, m.getD(), d_r, r * c);
|
|
return Matrix(r, c, d_r);
|
|
}
|
|
Matrix operator+(float a) {
|
|
float *d_r = (float*)Cu4cj_Zero(r * c * sizeof(float));
|
|
Cu4cj_AddBroadcast(d, a, d_r, r * c);
|
|
return Matrix(r, c, d_r);
|
|
}
|
|
Matrix operator-(const Matrix& m) {
|
|
float *d_r = (float*)Cu4cj_Zero(r * c * sizeof(float));
|
|
Cu4cj_Sub(d, m.getD(), d_r, r * c);
|
|
return Matrix(r, c, d_r);
|
|
}
|
|
Matrix operator-(float a) {
|
|
float *d_r = (float*)Cu4cj_Zero(r * c * sizeof(float));
|
|
Cu4cj_SubBroadcast(d, a, d_r, r * c);
|
|
return Matrix(r, c, d_r);
|
|
}
|
|
Matrix operator*(const Matrix& m) {
|
|
float *d_r = (float*)Cu4cj_Zero(r * c * sizeof(float));
|
|
Cu4cj_Mul(d, m.getD(), d_r, r * c);
|
|
return Matrix(r, c, d_r);
|
|
}
|
|
Matrix operator*(float a) {
|
|
float *d_r = (float*)Cu4cj_Zero(r * c * sizeof(float));
|
|
Cu4cj_MulBroadcast(d, a, d_r, r * c);
|
|
return Matrix(r, c, d_r);
|
|
}
|
|
Matrix operator/(const Matrix& m) {
|
|
float *d_r = (float*)Cu4cj_Zero(r * c * sizeof(float));
|
|
Cu4cj_Div(d, m.getD(), d_r, r * c);
|
|
return Matrix(r, c, d_r);
|
|
}
|
|
Matrix operator/(float a) {
|
|
float *d_r = (float*)Cu4cj_Zero(r * c * sizeof(float));
|
|
Cu4cj_DivBroadcast(d, a, d_r, r * c);
|
|
return Matrix(r, c, d_r);
|
|
}
|
|
Matrix dot(const Matrix& m) {
|
|
//if (!hasDotShape(m)) {
|
|
// throw std::invalid_argument("Dot product: matrices must have compatible shapes");
|
|
//}
|
|
float *d_r = (float*)Cu4cj_Zero(r * m.getC() * sizeof(float));
|
|
Cu4cj_MatDot(d, m.getD(), d_r, r, c, m.getC());
|
|
return Matrix(r, m.getC(), d_r);
|
|
}
|
|
Matrix T() {
|
|
float *d_r = (float*)Cu4cj_Transport(d, r, c);
|
|
return Matrix(c, r, d_r);
|
|
}
|
|
|
|
private:
|
|
uint64_t r, c;
|
|
float* d;
|
|
};
|
|
}
|
|
|
|
using namespace qing;
|
|
extern "C" {
|
|
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* createMatrix(const uint64_t r, const uint64_t c)
|
|
#else
|
|
void* createMatrix(int64_t r, int64_t c)
|
|
#endif
|
|
{
|
|
return(void*) new Matrix(r, c);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void freeMatrix(void* m)
|
|
#else
|
|
void freeMatrix(void* m)
|
|
#endif
|
|
{
|
|
delete (Matrix*)m;
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* xavierNormalInit(uint64_t fan_in, uint64_t fan_out)
|
|
#else
|
|
void* xavierNormalInit(uint64_t fan_in, uint64_t fan_out)
|
|
#endif
|
|
{
|
|
float* d = Cu4cj_Normal(fan_in * fan_out, fan_in, fan_out);
|
|
return (void *)new Matrix(fan_in, fan_out, d);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* xavierUniformInit(uint64_t fan_in, uint64_t fan_out)
|
|
#else
|
|
void* xavierUniformInit(uint64_t fan_in, uint64_t fan_out)
|
|
#endif
|
|
{
|
|
uint64_t n = fan_in * fan_out;
|
|
|
|
// 先在主机上分配内存并生成正态分布随机数
|
|
float* h = (float*)malloc(n * sizeof(float));
|
|
|
|
/* 使用高质量的随机数生成器 */
|
|
static std::mt19937 generator(std::random_device{}());
|
|
|
|
/* 计算区间范围：sqrt(2 / (fan_in + fan_out)) */
|
|
float scale = std::sqrt(2.0f / (fan_in + fan_out));
|
|
|
|
/* 创建均匀分布对象，范围为[-scale, scale] */
|
|
std::uniform_real_distribution<float> distribution(-scale, scale);
|
|
|
|
/* 遍历权重矩阵，进行赋值 */
|
|
uint64_t k=0;
|
|
for (uint64_t i=0; i<fan_in; ++i)
|
|
for (uint64_t j=0; j<fan_out; ++j)
|
|
h[k++] = distribution(generator);
|
|
|
|
|
|
|
|
// 在设备上分配内存
|
|
float* d = (float*)Cu4cj_Orig(n*sizeof(float));
|
|
|
|
// 将主机内存中的数据拷贝到设备内存
|
|
CHECK(cudaMemcpy(d, h, n * sizeof(float), cudaMemcpyHostToDevice));
|
|
|
|
// 释放主机内存
|
|
free(h);
|
|
|
|
return (void *)new Matrix(fan_in, fan_out, d);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) uint64_t getrowMatrix(void *m)
|
|
#else
|
|
uint64_t getrowMatrix(void *m)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
return m1->getR();
|
|
}
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) uint64_t getcolMatrix(void *m)
|
|
#else
|
|
uint64_t getcolMatrix(void *m)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
return m1->getC();
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) float getMatrix(void *m, uint64_t r, uint64_t c)
|
|
#else
|
|
float getMatrix(void *m, uint64_t r, uint64_t c)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
return m1->get(r, c);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void setMatrix(void *m, uint64_t r, uint64_t c, float v)
|
|
#else
|
|
void setMatrix(void *m, uint64_t r, uint64_t c, float v)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
m1->set(r, c, v);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) bool sameShapeMatrix(void *m, void *s)
|
|
#else
|
|
bool sameShapeMatrix(void *m, void *s)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix* s1 = (Matrix*)s;
|
|
return (*m1).hasSameShape(*s1);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) bool dotShapeMatrix(void *m, void *s)
|
|
#else
|
|
bool dotShapeMatrix(void *m, void *s)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix* s1 = (Matrix*)s;
|
|
return (*m1).hasDotShape(*s1);
|
|
}
|
|
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* plusMatrix(void *m, void *s)
|
|
#endif
|
|
#ifdef __linux__
|
|
void* plusMatrix(void *m, void *s)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix* s1 = (Matrix*)s;
|
|
Matrix r = *m1 + *s1;
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* plusBroadcastMatrix(void *m, float a)
|
|
#else
|
|
void* plusBroadcastMatrix(void *m, float a)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix r = *m1 + a;
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* subMatrix(void *m, void *s)
|
|
#else
|
|
void* subMatrix(void *m, void *s)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix* s1 = (Matrix*)s;
|
|
Matrix r = *m1 - *s1;
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* subBroadcastMatrix(void *m, float a)
|
|
#else
|
|
void* subBroadcastMatrix(void *m, float a)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix r = *m1 - a;
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* multiMatrix(void *m, void *s)
|
|
#else
|
|
void* multiMatrix(void *m, void *s)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix* s1 = (Matrix*)s;
|
|
Matrix r = (*m1) * (*s1);
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* multiBroadcastMatrix(void *m, float a)
|
|
#else
|
|
void* multiBroadcastMatrix(void *m, float a)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix r = (*m1) * a;
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* divMatrix(void *m, void *s)
|
|
#else
|
|
void* divMatrix(void *m, void *s)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix* s1 = (Matrix*)s;
|
|
Matrix r = (*m1) / (*s1);
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* divBroadcastMatrix(void *m, float a)
|
|
#else
|
|
void* divBroadcastMatrix(void *m, float a)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix r = (*m1) / a;
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* dotMatrix(void *m, void *s)
|
|
#else
|
|
void* dotMatrix(void *m, void *s)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix* s1 = (Matrix*)s;
|
|
Matrix r = (*m1).dot(*s1);
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* transposeMatrix(void *m)
|
|
#endif
|
|
#ifdef __linux__
|
|
void* transposeMatrix(void *m)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix r = m1->T();
|
|
return (void *)new Matrix(r);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* Sigmoid_do(void *m)
|
|
#else
|
|
void* Sigmoid_do(void *m)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix s = Matrix(*m1);
|
|
_devSigmoid<<<1, s.getR()* s.getC()>>>(s.getD(), s.getD(), s.getR()* s.getC());
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
return (void *)new Matrix(s);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* Sigmoid_dot(void *m)
|
|
#else
|
|
void* Sigmoid_dot(void *m)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix s = Matrix(*m1);
|
|
_devSigmoidDerivative<<<1, s.getR()* s.getC()>>>(s.getD(), s.getD(), s.getR()* s.getC());
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
return (void *)new Matrix(s);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* ReLU_do(void *m)
|
|
#else
|
|
void* ReLU_do(void *m)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix s = Matrix(*m1);
|
|
_devReLU<<<1, s.getR()* s.getC()>>>(s.getD(), s.getD(), s.getR()* s.getC());
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
return (void *)new Matrix(s);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* ReLU_dot(void *m)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix s = Matrix(*m1);
|
|
_devReLUDerivative<<<1, s.getR()* s.getC()>>>(s.getD(), s.getD(), s.getR()* s.getC());
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
return (void *)new Matrix(s);
|
|
}
|
|
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* Leaky_ReLU_do(void *m, float a)
|
|
#else
|
|
void* LeakyReLU_do(void *m, float a)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix s = Matrix(*m1);
|
|
_devLeakyReLU<<<1, s.getR()* s.getC()>>>(s.getD(), s.getD(), s.getR()* s.getC(), a);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
return (void *)new Matrix(s);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* Leaky_ReLU_dot(void *m, float a)
|
|
#else
|
|
void* LeakyReLU_dot(void *m, float a)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix s = Matrix(*m1);
|
|
_devLeakyReLU_Derivative<<<1, s.getR()* s.getC()>>>(s.getD(), s.getD(), s.getR()* s.getC(), a);
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
return (void *)new Matrix(s);
|
|
}
|
|
|
|
// Tanh
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* Tanh_do(void *m)
|
|
#else
|
|
void* Tanh_do(void *m)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix s = Matrix(*m1);
|
|
_devTanh<<<1, s.getR()* s.getC()>>>(s.getD(), s.getD(), s.getR()* s.getC());
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
return (void *)new Matrix(s);
|
|
}
|
|
|
|
#ifdef _WIN64
|
|
__declspec(dllexport) void* Tanh_dot(void *m)
|
|
#else
|
|
void* Tanh_dot(void *m)
|
|
#endif
|
|
{
|
|
Matrix* m1 = (Matrix*)m;
|
|
Matrix s = Matrix(*m1);
|
|
_devTanhDerivative<<<1, s.getR()* s.getC()>>>(s.getD(), s.getD(), s.getR()* s.getC());
|
|
CHECK(cudaGetLastError());
|
|
CHECK(cudaDeviceSynchronize());
|
|
return (void *)new Matrix(s);
|
|
}
|
|
|
|
|
|
} |