SeNet（Squeeze-and-Excitation Networks）反向传播数学推导详解（感谢百度和deepseek）

cudaMalloc(X_tilde, N * C * H * W * sizeof(float));int total_elements N * C * H * W; scale_nchw_kernel(total_elements 255) / 256, 256 (U, s, X_tilde, N, C, H, W);// 4. 残差连接: X X_tilde Uadd_kernel(total_elements 255) / 256, 256(X_tilde, U, X, total_elements);cudaFree(z); cudaFree(s); cudaFree(X_tilde);}5.5.2 反向传播实现NCHW 反向传播需要计算损失对输入U和权重W1、W2的梯度。以下是关键步骤的CUDA实现// senet_backward_nchw.cu// Scale层反向传播 - NCHW布局__global__ void scale_backward_nchw_kernel(const float* dL_dX_tilde, const float* s, float* dL_dU_scale, float* dL_ds_partial, const float* U, int N, int C, int H, int W){int idx blockIdx.x * blockDim.x threadIdx.x;int total N * C * H * W;if (idx total) return;// 分解索引int n idx / (C * H * W);int c (idx % (C * H * W)) / (H * W);int spatial_idx idx % (H * W);// dL/dU from scale: dL/dX_tilde * s_c dL_dU_scale[idx] dL_dX_tilde[idx] * s[n * C c];// 部分梯度用于计算dL/ds: dL/dX_tilde * u_c(i,j)// 使用原子操作累加每个样本每个通道的梯度 atomicAdd(dL_ds_partial[n * C c], dL_dX_tilde[idx] * U[idx]);}// Excitation层反向传播批处理void excitation_backward_nchw(const float* dL_ds, const float* s, const float* z, const float* a, const float* b, const float* W1, const float* W2, float* dL_dW1, float* dL_dW2, float* dL_dz, int N, int C, int r, cublasHandle_t handle){int mid_dim C / r;// 1. dL/dc dL/ds ⊙ (s ⊙ (1-s))float* dL_dc;cudaMalloc(dL_dc, N * C * sizeof(float));int threads 256; int blocks (N * C threads - 1) / threads; dc_kernelblocks, threads(dL_ds, s, dL_dc, N * C);// 2. dL/dW2 sum_over_batch(dL/dc * b^T) / Nfloat* dL_dW2_temp;cudaMalloc(dL_dW2_temp, C * mid_dim * sizeof(float));float alpha 1.0f, beta 0.0f;// 批处理矩阵乘法dL_dc: (C, N), b: (mid_dim, N) → dL_dW2_temp: (C, mid_dim)cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, C, mid_dim, N, alpha, dL_dc, C, b, mid_dim, beta, dL_dW2_temp, C);// 平均梯度可选取决于优化器float scale 1.0f / N;cublasSscal(handle, C * mid_dim, scale, dL_dW2_temp, 1); cudaMemcpy(dL_dW2, dL_dW2_temp, C * mid_dim * sizeof(float), cudaMemcpyDeviceToDevice);// 3. dL/db W2^T * dL/dc float* dL_db;cudaMalloc(dL_db, N * mid_dim * sizeof(float)); cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, mid_dim, N, C, alpha, W2, C, dL_dc, C, beta, dL_db, mid_dim);// 4. dL/da dL/db ⊙ ReLU(a)float* dL_da;cudaMalloc(dL_da, N * mid_dim * sizeof(float));da_kernel(N * mid_dim threads - 1) / threads, threads (dL_db, a, dL_da, N * mid_dim);// 5. dL/dW1 sum_over_batch(dL/da * z^T) / Nfloat* dL_dW1_temp;cudaMalloc(dL_dW1_temp, mid_dim * C * sizeof(float)); cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, mid_dim, C, N, alpha, dL_da, mid_dim, z, C, beta, dL_dW1_temp, mid_dim);cublasSscal(handle, mid_dim * C, scale, dL_dW1_temp, 1); cudaMemcpy(dL_dW1, dL_dW1_temp, mid_dim * C * sizeof(float), cudaMemcpyDeviceToDevice);// 6. dL/dz W1^T * dL/dacublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, C, N, mid_dim, alpha, W1, mid_dim, dL_da, mid_dim, beta, dL_dz, C); cudaFree(dL_dc);cudaFree(dL_dW2_temp); cudaFree(dL_db); cudaFree(dL_da); cudaFree(dL_dW1_temp); }// Squeeze层反向传播全局平均池化的反向- NCHW布局 __global__ void squeeze_backward_nchw_kernel(const float* dL_dz, float* dL_dU_squeeze, int N, int C, int H, int W){int idx blockIdx.x * blockDim.x threadIdx.x; int total N * C * H * W; if (idx total) return; int n idx / (C * H * W); int c (idx % (C * H * W)) / (H * W);// 平均池化的反向梯度均匀分配到每个空间位置 dL_dU_squeeze[idx] dL_dz[n * C c] / (H * W);}// 反向传播主函数NCHWvoid senet_backward_nchw(const float* dL_dX_tilde, const float* U, const float* s, const float* z, const float* a, const float* b, const float* W1, const float* W2, float* dL_dU, float* dL_dW1, float* dL_dW2, int N, int C, int H, int W, int r, cublasHandle_t handle){int total_elements N * C * H * W;// 分配临时内存float *dL_dU_scale, *dL_dU_squeeze, *dL_ds, *dL_dz; cudaMalloc(dL_dU_scale, total_elements * sizeof(float)); cudaMalloc(dL_dU_squeeze, total_elements * sizeof(float)); cudaMalloc(dL_ds, N * C * sizeof(float)); cudaMalloc(dL_dz, N * C * sizeof(float));// 1. Scale层反向传播 float* dL_ds_partial; cudaMalloc(dL_ds_partial, N * C * sizeof(float)); cudaMemset(dL_ds_partial, 0, N * C * sizeof(float)); scale_backward_nchw_kernel(total_elements 255) / 256, 256 (dL_dX_tilde, s, dL_dU_scale, dL_ds_partial, U, N, C, H, W); cudaMemcpy(dL_ds, dL_ds_partial, N * C * sizeof(float), cudaMemcpyDeviceToDevice);// 2. Excitation层反向传播 excitation_backward_nchw(dL_ds, s, z, a, b, W1, W2, dL_dW1, dL_dW2, dL_dz, N, C, r, handle);// 3. Squeeze层反向传播 squeeze_backward_nchw_kernel(total_elements 255) / 256, 256 (dL_dz, dL_dU_squeeze, N, C, H, W);// 4. 合并梯度dL/dU dL/dU_scale dL/dU_squeeze add_kernel(total_elements 255) / 256, 256 (dL_dU_scale, dL_dU_squeeze, dL_dU, total_elements);// 清理内存 cudaFree(dL_dU_scale); cudaFree(dL_dU_squeeze); cudaFree(dL_ds); cudaFree(dL_dz); cudaFree(dL_ds_partial); }泄露的另一个宝贝class Optimizer {public: virtual ~Optimizer() default;virtual void update(float* param, const float* grad, int size, cudaStream_t stream 0) 0;virtual void reset() 0; };/ SGD优化器 class SGD : public Optimizer { private: float learning_rate_; float momentum_; float* velocity_;// 动量项int size_;public:SGD(float lr 0.01f, float momentum 0.9f) : learning_rate_(lr), momentum_(momentum), velocity_(nullptr), size_(0) {}~SGD() {if (velocity_) cudaFree(velocity_);}void update(float* param, const float* grad, int size, cudaStream_t stream 0) override{// 第一次调用时分配动量内存if (!velocity_ || size_ ! size) {if (velocity_) cudaFree(velocity_);cudaMalloc(velocity_, size * sizeof(float));cudaMemset(velocity_, 0, size * sizeof(float)); size_ size; }// 动量更新: v momentum * v - lr * grad // param v float beta -learning_rate_; cublasHandle_t handle;cublasCreate(handle);// v momentum * vcublasSscal(handle, size, momentum_, velocity_, 1);// v v - lr * gradcublasSaxpy(handle, size, beta, grad, 1, velocity_, 1);// param param vfloat alpha 1.0f;cublasSaxpy(handle, size, alpha, velocity_, 1, param, 1);cublasDestroy(handle);}void reset() override {if (velocity_) { cudaMemset(velocity_, 0, size_ * sizeof(float)); }}};class SENetOptimizer {private: Optimizer* w1_optimizer_;Optimizer* w2_optimizer_;int w1_size_; int w2_size_;public:SENetOptimizer(Optimizer* w1_opt, Optimizer* w2_opt, int channels, int reduction){ w1_optimizer_ w1_opt;w2_optimizer_ w2_opt;int mid_dim channels / reduction;w1_size_ mid_dim * channels;// W1: (mid_dim, channels)w2_size_ channels * mid_dim;// W2: (channels, mid_dim) }~SENetOptimizer() { delete w1_optimizer_; delete w2_optimizer_;}void update_weights(float* W1, float* W2, const float* dL_dW1, const float* dL_dW2, cudaStream_t stream 0) {// 更新W1权重w1_optimizer_-update(W1, dL_dW1, w1_size_, stream);// 更新W2权重w2_optimizer_-update(W2, dL_dW2, w2_size_, stream);}void reset() { w1_optimizer_-reset(); w2_optimizer_-reset(); }};