遗漏的部分
This commit is contained in:
qing committed
1 parent
84fb370da8
commit
0493383ee8
63 files changed
+28237
-4
No files matched your search
@@ -0,0 +1,85 @@
|
||||
#include "nn/ILayer.hpp"
|
||||
#include <stdexcept>
|
||||
|
||||
/* 获取函数对象 - 获取激活函数 */
|
||||
f_t ILayer::getf(ActivationFunc type) {
|
||||
switch (type) {
|
||||
case ActivationFunc::ReLU:
|
||||
return relu;
|
||||
case ActivationFunc::Leaky_ReLU:
|
||||
return leaky_relu;
|
||||
case ActivationFunc::Sigmoid:
|
||||
return sigmoid;
|
||||
case ActivationFunc::Tanh:
|
||||
return tanh;
|
||||
default:
|
||||
throw std::invalid_argument("ActivationFunc function does not support.");
|
||||
}
|
||||
}
|
||||
|
||||
/* 获取函数对象 - 求激活函数导数 */
|
||||
f_t ILayer::getfdot(ActivationFunc type) {
|
||||
switch (type) {
|
||||
case ActivationFunc::ReLU:
|
||||
return relu_derivative;
|
||||
case ActivationFunc::Leaky_ReLU:
|
||||
return leaky_relu_derivative;
|
||||
case ActivationFunc::Sigmoid:
|
||||
return sigmoid_derivative;
|
||||
case ActivationFunc::Tanh:
|
||||
return tanh_derivative;
|
||||
default:
|
||||
throw std::invalid_argument("ActivationFunc Derivative does not support.");
|
||||
}
|
||||
}
|
||||
const char* ILayer::get_f_type(ActivationFunc& af) {
|
||||
switch (af) {
|
||||
case ActivationFunc::Sigmoid:
|
||||
return "Sigmoid";
|
||||
case ActivationFunc::ReLU:
|
||||
return "ReLU";
|
||||
case ActivationFunc::Leaky_ReLU:
|
||||
return "Leaky_ReLU";
|
||||
}
|
||||
throw std::runtime_error("What activation function");
|
||||
}
|
||||
|
||||
ILayer::ActivationFunc ILayer::parse_ff_type(const std::string& name) {
|
||||
if(name == "Sigmoid") {
|
||||
return ActivationFunc::Sigmoid;
|
||||
} else if (name == "ReLU") {
|
||||
return ActivationFunc::ReLU;
|
||||
} else if (name == "Leaky_ReLU") {
|
||||
return ActivationFunc::Leaky_ReLU;
|
||||
} else {
|
||||
throw std::runtime_error("What activation function");
|
||||
}
|
||||
}
|
||||
|
||||
/* ReLU 激活函数 */
|
||||
float ILayer::relu(float x) { return (x > 0.0) ? x : 0.0; }
|
||||
|
||||
/* ReLU 导数 */
|
||||
float ILayer::relu_derivative(float x) { return (x > 0) ? 1: 0; }
|
||||
|
||||
/*
|
||||
* leaky ReLU 激活函数
|
||||
*
|
||||
* NOTE: alpha太小或者选择使用ReLU函数的情况下，会导致梯度消失
|
||||
* NOTE: Leaky ReLU 不会使梯度完全消失，但也会使它变得很小
|
||||
*/
|
||||
float ILayer::leaky_relu(float x) { return (x > 0) ? x : LEAKY_RELU_ALPHA * x; }
|
||||
/* leaky ReLU 导数 */
|
||||
float ILayer::leaky_relu_derivative(float x) { return (x > 0) ? 1.0f : LEAKY_RELU_ALPHA; }
|
||||
/* Sigmoid 激活函数 */
|
||||
float ILayer::sigmoid(float x) { return 1.0 / (1.0 + std::exp(-x)); }
|
||||
/* Sigmoid 导数 */
|
||||
float ILayer::sigmoid_derivative(float x)
|
||||
{
|
||||
double s = sigmoid(x);
|
||||
return s * (1.0 - s);
|
||||
}
|
||||
/* Tanh 激活函数 */
|
||||
float ILayer::tanh(float x) { return std::tanh(x); }
|
||||
/* Tanh 导数 */
|
||||
float ILayer::tanh_derivative(float x) { return 1.0f - x * x; }
|
||||
@@ -0,0 +1,144 @@
|
||||
#include "nn/Matrix.hh"
|
||||
namespace qing {
|
||||
|
||||
template <typename T>
|
||||
Matrx<T>::Matrx(): row(0), col(0) {}
|
||||
|
||||
template <typename T>
|
||||
Matrx<T>::Matrx(const long row, const long col)
|
||||
: row(row), col(col) { /* 该构造函数用于从头开始构造 */
|
||||
arr = std::vector<T>(col * row);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
Matrx<T>::Matrx(const long row, const long col, const std::vector<T>& src)
|
||||
: row(row), col(col) { /* 该构造函数用于通过已有的vector进行构造 */
|
||||
arr = src;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
const T& Matrx<T>::get(long row, long col) const { /* 访问器 - 访问矩阵元素 */
|
||||
return arr[col * this->row + row];
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
T& Matrx<T>::witch(long row, long col) { /* 访问器 - 访问矩阵元素 */
|
||||
return arr[col * this->row + row];
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
const std::vector<T>& Matrx<T>::get_data() const{ /* 访问器 - 用于获取原始数据 */
|
||||
return arr;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
const long Matrx<T>::get_col() const { /* 获取列数 */
|
||||
return col;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
const long Matrx<T>::get_row() const { /* 获取行数 */
|
||||
return row;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
bool Matrx<T>::has_same_shape(const Matrx<T>& other) { /* 两个矩阵是否具有相同形状 */
|
||||
return other.get_col() == this->get_col() || other.get_row() == this->get_row();
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
bool Matrx<T>::has_dot_shape(const Matrx<T>& other) { /* 具有矩阵乘法的形状 */
|
||||
return this->get_col() == other.get_row();
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void Matrx<T>::assert_same_shape(const Matrx<T>& other) { /* 断言两个矩阵具有相同形状 */
|
||||
if (!has_same_shape(other))
|
||||
throw std::invalid_argument("The shapes of matrix are not same.");
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void Matrx<T>::assert_dot_shape(const Matrx<T>& other) { /* 断言两个矩阵具有点积前置形状 */
|
||||
if (!has_dot_shape(other)) {
|
||||
throw std::invalid_argument("The shapes of matrix can not be dot.");
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
Matrx<T> Matrx<T>::operator+(const Matrx<T>& other) { /* 加法操作符重定义 - 矩阵加法 */
|
||||
assert_same_shape(other); /* 检查两个矩阵形状是否一致 */
|
||||
std::vector<T> v0 = this->get_data();
|
||||
std::vector<T> vec = other.get_data();
|
||||
for (auto it = v0.begin(), it_or = vec.begin(); it != v0.end(); ++it, ++it_or ) {
|
||||
*it += *it_or;
|
||||
}
|
||||
return Matrx<T>{this->get_row(), this->get_col(), v0};
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
Matrx<T> Matrx<T>::operator+(T& item) { /* 矩阵加法广播 */
|
||||
auto matrix = *this;
|
||||
for (auto it = arr.begin(); it != arr.end(); ++it) {
|
||||
*it += item;
|
||||
}
|
||||
return matrix;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
Matrx<T> Matrx<T>::operator-(const Matrx<T>& other) { /* 操作符重载 - 矩阵减法 */
|
||||
assert_same_shape(other); /* 检查两个矩阵形状是否一致 */
|
||||
std::vector<T> v0 = this->get_data();
|
||||
std::vector <T> vec = other.get_data();
|
||||
for (auto it = v0.begin(), it_or = vec.begin(); it != v0.end(); ++it, ++it_or ) {
|
||||
*it -= *it_or;
|
||||
}
|
||||
return Matrx{this->get_row(), this->get_col(), v0};
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
Matrx<T> Matrx<T>::operator-(T& item) { /* 操作符重载 - 矩阵减法广播 */
|
||||
auto matrix = *this;
|
||||
for (auto it = arr.begin(); it != arr.end(); ++it) {
|
||||
*it -= item;
|
||||
}
|
||||
return matrix;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
const Matrx<T> Matrx<T>::operator*(const Matrx<T>& other) { /* 操作符重载 - 矩阵乘法 */
|
||||
assert_dot_shape(other); /* 它们的形状是否能够点积 */
|
||||
auto matrix = Matrx<T>(this->get_row(), other.get_col());
|
||||
for (long i=0; i<this->get_row(); ++i) {/* 点积运算 */
|
||||
for (long j=0; j<other.get_col(); ++j) {
|
||||
T sum{0};
|
||||
for (long k=0; k<this->get_col(); ++k) {
|
||||
sum += this->get(i, k) * other.get(k, j);
|
||||
}
|
||||
matrix.witch(i, j) = sum;
|
||||
}
|
||||
}
|
||||
return matrix;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
Matrx<T> Matrx<T>::operator*(T& item) { /* 操作符重载 - 矩阵乘法广播 */
|
||||
std::vector<T> vec = this->get_data();
|
||||
for (auto it = vec.begin(); it != vec.end(); ++it) {
|
||||
*it *= item;
|
||||
}
|
||||
return Matrx<T>{this->get_row(), this->get_col(), vec};
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
Matrx<T> Matrx<T>::operator/(T& item) { /* 操作符重载 - 矩阵除法广播 */
|
||||
std::vector<T> vec = this->get_data();
|
||||
for (auto it = vec.begin(); it != vec.end(); ++it) {
|
||||
*it /= item;
|
||||
}
|
||||
return Matrx<T>(this->get_row(), this->get_col(), vec);
|
||||
}
|
||||
|
||||
template class Matrx<float>; /* 显式实例化 */
|
||||
template class Matrx<double>;
|
||||
|
||||
}
|
||||
@@ -0,0 +1,396 @@
|
||||
/*
|
||||
* MatrixCuda.cu - MatrixCuda类的CUDA实现文件
|
||||
*
|
||||
* 这个文件包含MatrixCuda类的具体实现和CUDA内核的显式实例化
|
||||
* 注意：这个文件需要使用nvcc编译器编译
|
||||
*/
|
||||
|
||||
#include "nn/MatrixCuda.cuh"
|
||||
#include "nn/Matrix.hh" // 用于与现有Matrx类交互
|
||||
|
||||
namespace qing {
|
||||
|
||||
// ============================================================================
|
||||
// 显式实例化CUDA内核函数
|
||||
// ============================================================================
|
||||
|
||||
namespace cuda_kernels {
|
||||
|
||||
// float类型内核显式实例化
|
||||
template __global__ void matrix_add_kernel<float>(
|
||||
const float* A, const float* B, float* C,
|
||||
long rows, long cols, long total_elements);
|
||||
|
||||
template __global__ void matrix_sub_kernel<float>(
|
||||
const float* A, const float* B, float* C,
|
||||
long rows, long cols, long total_elements);
|
||||
|
||||
template __global__ void matrix_scalar_mul_kernel<float>(
|
||||
const float* A, float scalar, float* C, long total_elements);
|
||||
|
||||
template __global__ void matrix_scalar_add_kernel<float>(
|
||||
const float* A, float scalar, float* C, long total_elements);
|
||||
|
||||
template __global__ void matrix_transpose_kernel<float>(
|
||||
const float* A, float* C, long rows, long cols);
|
||||
|
||||
template __global__ void matrix_hadamard_kernel<float>(
|
||||
const float* A, const float* B, float* C, long total_elements);
|
||||
|
||||
template __global__ void matrix_copy_kernel<float>(
|
||||
const float* src, float* dst, long total_elements);
|
||||
|
||||
// double类型内核显式实例化
|
||||
template __global__ void matrix_add_kernel<double>(
|
||||
const double* A, const double* B, double* C,
|
||||
long rows, long cols, long total_elements);
|
||||
|
||||
template __global__ void matrix_sub_kernel<double>(
|
||||
const double* A, const double* B, double* C,
|
||||
long rows, long cols, long total_elements);
|
||||
|
||||
template __global__ void matrix_scalar_mul_kernel<double>(
|
||||
const double* A, double scalar, double* C, long total_elements);
|
||||
|
||||
template __global__ void matrix_scalar_add_kernel<double>(
|
||||
const double* A, double scalar, double* C, long total_elements);
|
||||
|
||||
template __global__ void matrix_transpose_kernel<double>(
|
||||
const double* A, double* C, long rows, long cols);
|
||||
|
||||
template __global__ void matrix_hadamard_kernel<double>(
|
||||
const double* A, const double* B, double* C, long total_elements);
|
||||
|
||||
template __global__ void matrix_copy_kernel<double>(
|
||||
const double* src, double* dst, long total_elements);
|
||||
|
||||
} // namespace cuda_kernels
|
||||
|
||||
// ============================================================================
|
||||
// MatrixCuda类与现有Matrx类的交互实现
|
||||
// ============================================================================
|
||||
|
||||
// 从Matrx<float>转换到MatrixCuda<float>
|
||||
template <>
|
||||
CUDA_HOST MatrixCuda<float> MatrixCuda<float>::from_matrx(const Matrx<float>& mat) {
|
||||
long rows = mat.get_row();
|
||||
long cols = mat.get_col();
|
||||
const std::vector<float>& data = mat.get_data();
|
||||
|
||||
return MatrixCuda<float>(rows, cols, data);
|
||||
}
|
||||
|
||||
// 从MatrixCuda<float>转换到Matrx<float>
|
||||
template <>
|
||||
CUDA_HOST Matrx<float> MatrixCuda<float>::to_matrx() const {
|
||||
sync_device_to_host();
|
||||
return Matrx<float>(rows_, cols_, host_data_);
|
||||
}
|
||||
|
||||
// 从Matrx<double>转换到MatrixCuda<double>
|
||||
template <>
|
||||
CUDA_HOST MatrixCuda<double> MatrixCuda<double>::from_matrx(const Matrx<double>& mat) {
|
||||
long rows = mat.get_row();
|
||||
long cols = mat.get_col();
|
||||
const std::vector<double>& data = mat.get_data();
|
||||
|
||||
return MatrixCuda<double>(rows, cols, data);
|
||||
}
|
||||
|
||||
// 从MatrixCuda<double>转换到Matrx<double>
|
||||
template <>
|
||||
CUDA_HOST Matrx<double> MatrixCuda<double>::to_matrx() const {
|
||||
sync_device_to_host();
|
||||
return Matrx<double>(rows_, cols_, host_data_);
|
||||
}
|
||||
|
||||
// ============================================================================
|
||||
// 优化的矩阵乘法实现（使用共享内存）
|
||||
// ============================================================================
|
||||
|
||||
/**
|
||||
* @brief 使用共享内存优化的矩阵乘法内核
|
||||
*
|
||||
* 这个内核使用平铺（tiling）技术来提高矩阵乘法的性能
|
||||
* 每个线程块计算输出矩阵的一个子块
|
||||
*/
|
||||
template <typename T>
|
||||
__global__ void matrix_mul_optimized_kernel(
|
||||
const T* A, const T* B, T* C,
|
||||
long A_rows, long A_cols, long B_cols,
|
||||
long tile_size) {
|
||||
|
||||
// 共享内存声明
|
||||
extern __shared__ T shared_mem[];
|
||||
T* tile_A = shared_mem;
|
||||
T* tile_B = &shared_mem[tile_size * tile_size];
|
||||
|
||||
// 线程块索引
|
||||
long block_row = blockIdx.y;
|
||||
long block_col = blockIdx.x;
|
||||
|
||||
// 线程索引
|
||||
long thread_row = threadIdx.y;
|
||||
long thread_col = threadIdx.x;
|
||||
|
||||
// 输出矩阵中的位置
|
||||
long row = block_row * tile_size + thread_row;
|
||||
long col = block_col * tile_size + thread_col;
|
||||
|
||||
T sum = T(0);
|
||||
|
||||
// 遍历所有平铺
|
||||
for (long tile_idx = 0; tile_idx < (A_cols + tile_size - 1) / tile_size; ++tile_idx) {
|
||||
// 从全局内存加载到共享内存
|
||||
long A_col = tile_idx * tile_size + thread_col;
|
||||
long B_row = tile_idx * tile_size + thread_row;
|
||||
|
||||
if (row < A_rows && A_col < A_cols) {
|
||||
tile_A[thread_row * tile_size + thread_col] = A[row * A_cols + A_col];
|
||||
} else {
|
||||
tile_A[thread_row * tile_size + thread_col] = T(0);
|
||||
}
|
||||
|
||||
if (B_row < A_cols && col < B_cols) {
|
||||
tile_B[thread_row * tile_size + thread_col] = B[B_row * B_cols + col];
|
||||
} else {
|
||||
tile_B[thread_row * tile_size + thread_col] = T(0);
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
// 计算平铺内的点积
|
||||
for (long k = 0; k < tile_size; ++k) {
|
||||
sum += tile_A[thread_row * tile_size + k] * tile_B[k * tile_size + thread_col];
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
// 将结果写回全局内存
|
||||
if (row < A_rows && col < B_cols) {
|
||||
C[row * B_cols + col] = sum;
|
||||
}
|
||||
}
|
||||
|
||||
// ============================================================================
|
||||
// 优化的矩阵乘法方法
|
||||
// ============================================================================
|
||||
|
||||
/**
|
||||
* @brief 使用优化CUDA内核的矩阵乘法
|
||||
*
|
||||
* 这个方法使用共享内存和平铺技术来加速矩阵乘法
|
||||
* 对于大矩阵，性能比朴素算法有显著提升
|
||||
*/
|
||||
template <typename T>
|
||||
CUDA_HOST MatrixCuda<T> MatrixCuda<T>::matmul_optimized(const MatrixCuda<T>& other) const {
|
||||
assert_dot_shape(other);
|
||||
|
||||
MatrixCuda<T> result(rows_, other.cols_);
|
||||
|
||||
// 确保数据在设备上
|
||||
sync_host_to_device();
|
||||
other.sync_host_to_device();
|
||||
result.allocate_device_memory();
|
||||
|
||||
// 设置平铺大小（典型值为16或32）
|
||||
const long tile_size = 16;
|
||||
|
||||
// 计算线程块和网格大小
|
||||
dim3 block_size(tile_size, tile_size);
|
||||
dim3 grid_size((other.cols_ + tile_size - 1) / tile_size,
|
||||
(rows_ + tile_size - 1) / tile_size);
|
||||
|
||||
// 计算共享内存大小
|
||||
size_t shared_mem_size = 2 * tile_size * tile_size * sizeof(T);
|
||||
|
||||
// 启动优化内核
|
||||
matrix_mul_optimized_kernel<T><<<grid_size, block_size, shared_mem_size>>>(
|
||||
device_data_, other.device_data_, result.device_data_,
|
||||
rows_, cols_, other.cols_, tile_size
|
||||
);
|
||||
|
||||
CHECK_CUDA_ERROR(cudaGetLastError());
|
||||
CHECK_CUDA_ERROR(cudaDeviceSynchronize());
|
||||
|
||||
result.device_modified_ = true;
|
||||
result.sync_device_to_host();
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
// 显式实例化优化矩阵乘法
|
||||
template MatrixCuda<float> MatrixCuda<float>::matmul_optimized(const MatrixCuda<float>& other) const;
|
||||
template MatrixCuda<double> MatrixCuda<double>::matmul_optimized(const MatrixCuda<double>& other) const;
|
||||
|
||||
// ============================================================================
|
||||
// 使用cuBLAS的矩阵乘法（如果可用）
|
||||
// ============================================================================
|
||||
|
||||
#ifdef HAVE_CUBLAS
|
||||
#include <cublas_v2.h>
|
||||
|
||||
template <typename T>
|
||||
CUDA_HOST MatrixCuda<T> MatrixCuda<T>::matmul_cublas(const MatrixCuda<T>& other) const {
|
||||
assert_dot_shape(other);
|
||||
|
||||
MatrixCuda<T> result(rows_, other.cols_);
|
||||
|
||||
// 确保数据在设备上
|
||||
sync_host_to_device();
|
||||
other.sync_host_to_device();
|
||||
result.allocate_device_memory();
|
||||
|
||||
// 创建cuBLAS句柄
|
||||
cublasHandle_t handle;
|
||||
cublasCreate(&handle);
|
||||
|
||||
// 设置cuBLAS参数
|
||||
T alpha = T(1);
|
||||
T beta = T(0);
|
||||
|
||||
if constexpr (std::is_same_v<T, float>) {
|
||||
cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
|
||||
other.cols_, rows_, cols_,
|
||||
&alpha,
|
||||
other.device_data_, other.cols_,
|
||||
device_data_, cols_,
|
||||
&beta,
|
||||
result.device_data_, other.cols_);
|
||||
} else if constexpr (std::is_same_v<T, double>) {
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
|
||||
other.cols_, rows_, cols_,
|
||||
&alpha,
|
||||
other.device_data_, other.cols_,
|
||||
device_data_, cols_,
|
||||
&beta,
|
||||
result.device_data_, other.cols_);
|
||||
}
|
||||
|
||||
// 销毁cuBLAS句柄
|
||||
cublasDestroy(handle);
|
||||
|
||||
CHECK_CUDA_ERROR(cudaGetLastError());
|
||||
CHECK_CUDA_ERROR(cudaDeviceSynchronize());
|
||||
|
||||
result.device_modified_ = true;
|
||||
result.sync_device_to_host();
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
// 显式实例化cuBLAS矩阵乘法
|
||||
template MatrixCuda<float> MatrixCuda<float>::matmul_cublas(const MatrixCuda<float>& other) const;
|
||||
template MatrixCuda<double> MatrixCuda<double>::matmul_cublas(const MatrixCuda<double>& other) const;
|
||||
|
||||
#endif // HAVE_CUBLAS
|
||||
|
||||
// ============================================================================
|
||||
// 测试函数
|
||||
// ============================================================================
|
||||
|
||||
/**
|
||||
* @brief 测试MatrixCuda类的功能
|
||||
*
|
||||
* 这个函数用于验证MatrixCuda类的正确性
|
||||
* 可以单独编译和运行进行测试
|
||||
*/
|
||||
template <typename T>
|
||||
CUDA_HOST void test_matrix_cuda() {
|
||||
std::cout << "Testing MatrixCuda<" << typeid(T).name() << ">..." << std::endl;
|
||||
|
||||
try {
|
||||
// 测试1：创建矩阵
|
||||
std::cout << "Test 1: Creating matrices..." << std::endl;
|
||||
MatrixCuda<T> A = MatrixCuda<T>::zeros(3, 3);
|
||||
MatrixCuda<T> B = MatrixCuda<T>::ones(3, 3);
|
||||
MatrixCuda<T> C = MatrixCuda<T>::identity(3);
|
||||
|
||||
A.print("A (zeros)");
|
||||
B.print("B (ones)");
|
||||
C.print("C (identity)");
|
||||
|
||||
// 测试2：矩阵加法
|
||||
std::cout << "\nTest 2: Matrix addition..." << std::endl;
|
||||
MatrixCuda<T> D = A + B;
|
||||
D.print("D = A + B");
|
||||
|
||||
// 测试3：标量运算
|
||||
std::cout << "\nTest 3: Scalar operations..." << std::endl;
|
||||
MatrixCuda<T> E = B * T(2.5);
|
||||
E.print("E = B * 2.5");
|
||||
|
||||
// 测试4：矩阵转置
|
||||
std::cout << "\nTest 4: Matrix transpose..." << std::endl;
|
||||
MatrixCuda<T> F(2, 3);
|
||||
F.fill(T(1));
|
||||
F.print("F (2x3)");
|
||||
MatrixCuda<T> FT = F.transpose();
|
||||
FT.print("F^T (3x2)");
|
||||
|
||||
// 测试5：元素对应乘积
|
||||
std::cout << "\nTest 5: Hadamard product..." << std::endl;
|
||||
MatrixCuda<T> G = MatrixCuda<T>::random(2, 2, T(0), T(1));
|
||||
MatrixCuda<T> H = MatrixCuda<T>::random(2, 2, T(0), T(1));
|
||||
G.print("G");
|
||||
H.print("H");
|
||||
MatrixCuda<T> I = G.hadamard_product(H);
|
||||
I.print("G ⊙ H");
|
||||
|
||||
// 测试6：矩阵乘法
|
||||
std::cout << "\nTest 6: Matrix multiplication..." << std::endl;
|
||||
MatrixCuda<T> J(2, 3);
|
||||
MatrixCuda<T> K(3, 2);
|
||||
J.fill(T(1));
|
||||
K.fill(T(2));
|
||||
J.print("J (2x3)");
|
||||
K.print("K (3x2)");
|
||||
MatrixCuda<T> L = J * K;
|
||||
L.print("J * K (2x2)");
|
||||
|
||||
std::cout << "All tests passed!" << std::endl;
|
||||
|
||||
} catch (const std::exception& e) {
|
||||
std::cerr << "Test failed: " << e.what() << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
// 显式实例化测试函数
|
||||
template void test_matrix_cuda<float>();
|
||||
template void test_matrix_cuda<double>();
|
||||
|
||||
} // namespace qing
|
||||
|
||||
// ============================================================================
|
||||
// 主测试程序（可选）
|
||||
// ============================================================================
|
||||
|
||||
#ifdef TEST_MATRIX_CUDA
|
||||
|
||||
int main() {
|
||||
std::cout << "=== MatrixCuda Test Program ===" << std::endl;
|
||||
|
||||
// 初始化CUDA
|
||||
cudaError_t err = cudaSetDevice(0);
|
||||
if (err != cudaSuccess) {
|
||||
std::cerr << "Failed to set CUDA device: " << cudaGetErrorString(err) << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// 测试float版本
|
||||
qing::test_matrix_cuda<float>();
|
||||
|
||||
std::cout << "\n";
|
||||
|
||||
// 测试double版本
|
||||
qing::test_matrix_cuda<double>();
|
||||
|
||||
// 重置CUDA设备
|
||||
cudaDeviceReset();
|
||||
|
||||
std::cout << "\n=== Test Program Finished ===" << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
#endif // TEST_MATRIX_CUDA
|
||||
@@ -0,0 +1,48 @@
|
||||
#include "nn/NNBuilder.hpp"
|
||||
namespace qing {
|
||||
|
||||
void NNBuilder::add(NeuralNetwork &layer) { nn.push_back(layer); }
|
||||
|
||||
std::vector<std::pair<long, long>> NNBuilder::get_shape() const {
|
||||
std::vector<std::pair<long, long>> shape;
|
||||
for (long j = 0; j < nn.size(); ++j) {
|
||||
shape.emplace_back(nn[j].get_inputno(), nn[j].get_outputno());
|
||||
}
|
||||
return shape;
|
||||
}
|
||||
|
||||
void NNBuilder::print_shape() const {
|
||||
auto shape = get_shape();
|
||||
for (const auto& p : shape) {
|
||||
std::cout << p.first << " -> " << p.second << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
/* 前向反馈 */
|
||||
std::vector<float> NNBuilder::forward(std::vector<float>& r) {
|
||||
auto x = r;
|
||||
for (long j=0; j< nn.size(); ++j) {
|
||||
x = nn[j].forward(x);
|
||||
}
|
||||
return x;
|
||||
}
|
||||
|
||||
/* 反向传播 */
|
||||
std::vector<float> NNBuilder::backward(std::vector<float>& errs, float discount) {
|
||||
auto e = errs;
|
||||
for (long j=nn.size()-1; j>=0;--j) {
|
||||
e = nn[j].backward(e);
|
||||
nn[j].update(discount);
|
||||
}
|
||||
return e;
|
||||
}
|
||||
|
||||
/* 计算误差 */
|
||||
std::vector<float> NNBuilder::cal_err(std::vector<float>& r, std::vector<float>& t) {
|
||||
auto res = r;
|
||||
for (long j=0; j<r.size(); ++j) {
|
||||
res[j] = t[j] - r[j];
|
||||
}
|
||||
return res;
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,231 @@
|
||||
#include "nn/NeuralNetwork.hpp"
|
||||
namespace qing{
|
||||
/* 构造函数 使用基本参数构造 */
|
||||
NeuralNetwork::NeuralNetwork(
|
||||
const long inputno, const long outputno,
|
||||
const std::vector<float> &weights,
|
||||
const std::vector<float> &bias,
|
||||
const float learning_rate,
|
||||
const ActivationFunc func_type
|
||||
): inputno(inputno), outputno(outputno),
|
||||
weights(inputno, outputno, weights), bias( 1, outputno, bias),
|
||||
learning_rate(learning_rate),
|
||||
f_type(func_type) {
|
||||
;/* 进行形状检查 */
|
||||
}
|
||||
|
||||
/* 获取输入层大小 */
|
||||
long NeuralNetwork::get_inputno() const {
|
||||
return this->inputno;
|
||||
}
|
||||
|
||||
/* 获取输出层大小 */
|
||||
long NeuralNetwork::get_outputno() const {
|
||||
return this->outputno;
|
||||
}
|
||||
|
||||
/* 获取权重 */
|
||||
Matrx<float>& NeuralNetwork::get_weights() {
|
||||
return this->weights;
|
||||
}
|
||||
|
||||
/* 获取偏置 */
|
||||
Matrx<float>& NeuralNetwork::get_bias() {
|
||||
return this->bias;
|
||||
}
|
||||
|
||||
/* 正态分布Xavier初始化 */
|
||||
void NeuralNetwork::xavierNormalInit(std::vector<float> &weights, long fan_in, long fan_out)
|
||||
{
|
||||
|
||||
/* 高质量的随机数生成器 */
|
||||
static std::mt19937 generator(std::random_device{}());
|
||||
|
||||
/* 计算标准差 */
|
||||
float sigma = std::sqrt(2.0f / (fan_in + fan_out));
|
||||
|
||||
/* 创建正态分布对象，均值为0，标准差为sigma */
|
||||
std::normal_distribution<float> distribution(0.0f, sigma);
|
||||
|
||||
/* 遍历权重矩阵，进行赋值 */
|
||||
long k=0;
|
||||
for (long i=0; i<fan_in; ++i)
|
||||
for (long j=0; j<fan_out; ++j)
|
||||
weights[k++] = distribution(generator);
|
||||
|
||||
}
|
||||
|
||||
/* 均匀分布Xavier初始化函数 */
|
||||
void NeuralNetwork::xavierUniformInit(std::vector<float> &weights, int fan_in, int fan_out)
|
||||
{
|
||||
/* 使用高质量的随机数生成器 */
|
||||
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);
|
||||
|
||||
/* 遍历权重矩阵，进行赋值 */
|
||||
long k=0;
|
||||
for (long i=0; i<fan_in; ++i)
|
||||
for (long j=0; j<fan_out; ++j)
|
||||
weights[k++] = distribution(generator);
|
||||
}
|
||||
|
||||
/* 工厂模式 通过形状构建全连接层 */
|
||||
NeuralNetwork NeuralNetwork::Create_in_Factory(const long inputno, const long outputno, float learning_rate, ActivationFunc func_type)
|
||||
{
|
||||
std::vector<float> weights(inputno * outputno);
|
||||
xavierUniformInit(weights, inputno, outputno);
|
||||
std::vector<float> bias(outputno);
|
||||
NeuralNetwork layer(inputno, outputno, weights, bias, learning_rate, func_type);
|
||||
return layer; /* FIXME: 使用移动语义 */
|
||||
}
|
||||
|
||||
|
||||
/* 前向反馈 */
|
||||
std::vector<float> NeuralNetwork::forward(const std::vector<float>& x)
|
||||
|
||||
{
|
||||
|
||||
|
||||
/*
|
||||
* 矩阵乘法（点积）
|
||||
* wi => weights inputs(column)
|
||||
* wo => weights outputs(row)
|
||||
*
|
||||
* NOTE: 通常矩阵乘法是三层循环，这里我们用两层
|
||||
*/
|
||||
//std::vector<float> res(this->outputno); /* (1, outputs) */
|
||||
//for (long i=0; i<this->outputno; ++i) {
|
||||
// long wo = i;
|
||||
// float sum=0.0;
|
||||
// for (long j=0, wi=wo; j<this->inputno; ++j, wi+=this->outputno) {
|
||||
// sum += x[j] * this->weights[wi];
|
||||
// }
|
||||
// res[i]= sum;
|
||||
//}
|
||||
Matrx<float>mx{1, this->inputno, x};
|
||||
|
||||
|
||||
/*
|
||||
* 为输入进行备份
|
||||
*
|
||||
* FIXME: 使用移动语义传入参数，并且储存
|
||||
*/
|
||||
this->prev_inputs = mx; /* (1, inputs) */
|
||||
|
||||
|
||||
Matrx<float>r = mx * this->weights; /* (1, outputs) */
|
||||
|
||||
/* 偏置操作 */
|
||||
//for (long i=0; i<this->outputno; ++i) {
|
||||
// res[i]+=this->bias[i];
|
||||
//}
|
||||
r = r + this->bias;
|
||||
|
||||
/*
|
||||
* 为原始计算结果进行备份
|
||||
*/
|
||||
this->prev_raw_inputs = r;
|
||||
|
||||
/*
|
||||
* 使用ReLU激活函数
|
||||
*
|
||||
* FIXME: 建立一个算法类
|
||||
*/
|
||||
f_t f = getf(f_type);
|
||||
|
||||
std::vector<float> vec =r.get_data();
|
||||
for (long i=0; i<this->outputno; ++i) {
|
||||
float& v = vec[i];
|
||||
v = f(v);
|
||||
}
|
||||
|
||||
r = Matrx<float>(1, outputno, vec);
|
||||
|
||||
/*
|
||||
* 为输出进行备份
|
||||
*/
|
||||
this->prev_outputs = r;
|
||||
|
||||
|
||||
/* 返回前向反馈的结果。 (1, outputs) */
|
||||
return r.get_data();
|
||||
}
|
||||
|
||||
|
||||
/* 反向传播 */
|
||||
std::vector<float> NeuralNetwork::backward(const std::vector<float>& errors) /* errors => (outputs, 1) */
|
||||
|
||||
{
|
||||
|
||||
/* 匹配激活导数 */
|
||||
f_t fdot = getfdot(f_type);
|
||||
|
||||
std::vector<float> vec = this->prev_raw_inputs.get_data();
|
||||
/* 将原始输出转化为导数 */
|
||||
for (long i=0; i<this->outputno; ++i) {
|
||||
float& v = vec[i];
|
||||
v = errors[i] * fdot(v);
|
||||
}
|
||||
|
||||
/* 误差乘以导数 */
|
||||
//std::vector<float> first(this->outputno);
|
||||
//for (long i=0; i<this->outputno; ++i) {
|
||||
// first[i] = errors[i] * this->prev_raw_inputs[i];
|
||||
//}
|
||||
Matrx<float> first{1, this->outputno, vec};
|
||||
|
||||
/* 更新权重梯度 => (errors * 导数) @ this->inputs => (1, outputs ) @ (inputs, 1) = (outputs, inputs) */
|
||||
//this->weights_grad = std::vector<float>(this->outputno * this->inputno);
|
||||
//long k = 0;
|
||||
//for (long i=0; i<this->inputno; ++i) {
|
||||
// for (long j=0; j<this->outputno; ++j) {
|
||||
// this->weights_grad[k++] = first[j] * this->prev_inputs[i];
|
||||
// }
|
||||
//}
|
||||
Matrx<float> mx {this->inputno, 1, this->prev_inputs.get_data()}; /* (inputs, 1) */
|
||||
this->weights_grad = mx * first; /* (inputs, 1) @ (1, outputs) => (inputs, outputs) */
|
||||
|
||||
/*
|
||||
* 更新偏置梯度
|
||||
*
|
||||
* NOTE: 在多批次情况下，应该沿着每一样本纵向求平均数
|
||||
*/
|
||||
//this->bias_grad = std::vector<float>(this->outputno);
|
||||
//for (long i=0; i<this->outputno; ++i) {
|
||||
// this->bias_grad[i] = first[i];
|
||||
//}
|
||||
this->bias_grad = first;
|
||||
|
||||
|
||||
/*
|
||||
* 计算上一层的误差责任分担，然后返回
|
||||
*
|
||||
* this->weights @ first => (outputs, inputs) @ (1, outputs) = (inputs, 1)
|
||||
*/
|
||||
//std::vector<float> uplevel_errors(this->inputno);
|
||||
//for (long i=0; i<this->inputno; ++i) {
|
||||
// long wii = i * this->outputno; /* 每一列的起始点 */
|
||||
// float sum = 0.0f;
|
||||
// for (long j=0, wi=wii; j < this->outputno; ++j, ++wi) {
|
||||
// sum += this->weights[wi] * first[j];
|
||||
// }
|
||||
// uplevel_errors[i] = sum;
|
||||
//}
|
||||
Matrx<float> first1{this->outputno, 1, first.get_data()};
|
||||
Matrx<float> uplevel_errs = this->weights * first1; /* (inputs, outputs) @ (outputs, 1) */
|
||||
return uplevel_errs.get_data();
|
||||
|
||||
/* 使用移动语义返回上一层误差 */
|
||||
//return std::move(uplevel_errors);
|
||||
}
|
||||
|
||||
void NeuralNetwork::update(float discount) {
|
||||
this->weights = this->weights + (this->weights_grad * this->learning_rate * discount);
|
||||
this->bias = this->bias + (this->bias_grad * this->learning_rate * discount);
|
||||
}
|
||||
}
|
||||
Reference in new issue
Block a user