template backward.cl

cl/backward.cl

@@ -0,0 +1,11 @@
+__kernel void backward(__global float* activate,
+__global float* output,
+__global float* delta,
+__global float* input,
+__global float* mul,
+__global float* add,
+int input_width,
+int param_width)
+{
+
+}

src/layer.rs

@@ -1,11 +1,11 @@
 use ocl::{
     builders::{BufferBuilder, KernelBuilder},
-    MemFlags, SpatialDims,
+    Buffer, SpatialDims,
 };
 use serde::{Deserialize, Serialize};
 
 use crate::{
-    state::{Context, PROGRAM_FORWARD, PROGRAM_BACKWARD},
+    state::{Context, PROGRAM_BACKWARD, PROGRAM_FORWARD},
     Error,
 };
 
@@ -96,11 +96,11 @@ impl Layer {
     /// FIXME: we should use host memory instead device memory (EG. GPU)
     ///
     /// MEM_USE_HOST_PTR: use host memory, cache by device memory
-    pub fn forward(&mut self, state: &Context, activation: &ocl::Buffer<f32>) -> Result<(), Error> {
+    pub fn forward(&mut self, ctx: &Context, activation: &ocl::Buffer<f32>) -> Result<(), Error> {
         let kernel = KernelBuilder::new()
-            .queue(state.queue.clone())
+            .queue(ctx.queue.clone())
             .global_work_size(SpatialDims::One(self.inter))
-            .program(&state.program[PROGRAM_FORWARD])
+            .program(&ctx.program[PROGRAM_FORWARD])
             .arg(&self.activate)
             .arg(&self.output)
             .arg(activation)
@@ -118,15 +118,23 @@ impl Layer {
     }
     /// forward pagination
     /// 
+    /// delta: da superscript [L], kernel is require to rewrite it to da superscript [L-1   ]
+    ///
     /// FIXME: we should use host memory instead device memory (EG. GPU)
     ///
     /// MEM_USE_HOST_PTR: use host memory, cache by device memory
-    pub fn backward(&mut self, state: &Context) -> Result<(), Error> {
-        todo!();
+    pub fn backward(&mut self, ctx: &Context, delta: &mut Buffer<f32>) -> Result<(), Error> {
         let kernel = KernelBuilder::new()
-            .queue(state.queue.clone())
+            .queue(ctx.queue.clone())
             .global_work_size(SpatialDims::One(self.inter))
-            .program(&state.program[PROGRAM_BACKWARD])
+            .program(&ctx.program[PROGRAM_BACKWARD])
+            .arg(&self.activate)
+            .arg(&self.output)
+            .arg(delta)
+            .arg(&self.mul_conn)
+            .arg(&self.offset_conn)
+            .arg(self.input)
+            .arg(self.inter)
             .build()?;
 
         unsafe {

src/state.rs

@@ -83,11 +83,18 @@ impl Layers {
             input = &layer.activate;
         }
 
+        let mut delta = BufferBuilder::<f32>::new()
+            .context(&ctx.context)
+            .len(1024)
+            .fill_val(1.0)
+            .build()?;
+
         for layer in self.0.iter_mut().rev() {
-            layer.backward(ctx)?;
+            layer.backward(ctx, &mut delta)?;
         }
 
         drop(data);
+        
         Ok(())
     }
 }