From b8ac3b860d9765d7aeec5fa0a29769c6da597aa2 Mon Sep 17 00:00:00 2001 From: "Randall C. O'Reilly" Date: Tue, 26 Nov 2024 13:15:47 -0800 Subject: [PATCH] add missing WtFromDWtLayer --- axon/act-net.go | 3 +- axon/act-net.goal | 3 +- axon/gosl.go | 43 + axon/learn-layer.go | 41 +- axon/learn-layer.goal | 41 +- axon/learn-net.go | 25 +- axon/learn-net.goal | 25 +- axon/shaders/ApplyExtsNeuron.wgsl | 2 + axon/shaders/BetweenGi.wgsl | 2 + axon/shaders/CycleInc.wgsl | 2 + axon/shaders/CycleNeuron.wgsl | 2 + axon/shaders/CyclePost.wgsl | 2 + axon/shaders/DWtFromDiSyn.wgsl | 2 + axon/shaders/DWtSubMeanPath.wgsl | 2 + axon/shaders/DWtSyn.wgsl | 2 + axon/shaders/GPUTestWrite.wgsl | 2 + axon/shaders/GatherSpikes.wgsl | 2 + axon/shaders/InitGBuffsPath.wgsl | 2 + axon/shaders/LayerGi.wgsl | 2 + axon/shaders/MinusPhaseNeuron.wgsl | 2 + axon/shaders/MinusPhasePool.wgsl | 2 + axon/shaders/MinusPhasePost.wgsl | 2 + axon/shaders/NewStateLayer.wgsl | 2 + axon/shaders/NewStateNeuron.wgsl | 2 + axon/shaders/PlusPhaseNeuron.wgsl | 2 + axon/shaders/PlusPhasePool.wgsl | 2 + axon/shaders/PlusPhasePost.wgsl | 2 + axon/shaders/PlusPhaseStartNeuron.wgsl | 2 + axon/shaders/PoolGi.wgsl | 2 + axon/shaders/SendSpike.wgsl | 2 + axon/shaders/WtFromDWtLayer.wgsl | 1487 ++++++++++++++++++++++++ axon/shaders/WtFromDWtSyn.wgsl | 2 + 32 files changed, 1650 insertions(+), 66 deletions(-) create mode 100644 axon/shaders/WtFromDWtLayer.wgsl diff --git a/axon/act-net.go b/axon/act-net.go index 4c3b1330..875acdd1 100644 --- a/axon/act-net.go +++ b/axon/act-net.go @@ -124,7 +124,8 @@ func (nt *Network) PlusPhaseStart() { RunPlusPhaseStartNeuron(nd) } -// PlusPhase does updating after end of plus phase +// PlusPhase does updating after end of plus phase. +// On GPU this is when we finally sync back Layers and Neurons. func (nt *Network) PlusPhase() { nix := nt.NetIxs() ctx := nt.Context() diff --git a/axon/act-net.goal b/axon/act-net.goal index 27623805..e5167a45 100644 --- a/axon/act-net.goal +++ b/axon/act-net.goal @@ -118,7 +118,8 @@ func (nt *Network) PlusPhaseStart() { RunPlusPhaseStartNeuron(nd) } -// PlusPhase does updating after end of plus phase +// PlusPhase does updating after end of plus phase. +// On GPU this is when we finally sync back Layers and Neurons. func (nt *Network) PlusPhase() { nix := nt.NetIxs() ctx := nt.Context() diff --git a/axon/gosl.go b/axon/gosl.go index 0078450e..27aac156 100644 --- a/axon/gosl.go +++ b/axon/gosl.go @@ -84,6 +84,7 @@ func GPUInit() { gpu.NewComputePipelineShaderFS(shaders, "shaders/PlusPhaseStartNeuron.wgsl", sy) gpu.NewComputePipelineShaderFS(shaders, "shaders/PoolGi.wgsl", sy) gpu.NewComputePipelineShaderFS(shaders, "shaders/SendSpike.wgsl", sy) + gpu.NewComputePipelineShaderFS(shaders, "shaders/WtFromDWtLayer.wgsl", sy) gpu.NewComputePipelineShaderFS(shaders, "shaders/WtFromDWtSyn.wgsl", sy) vars := sy.Vars() { @@ -1125,6 +1126,48 @@ func RunOneSendSpike(n int, syncVars ...GPUVars) { RunSendSpikeCPU(n) } } +// RunWtFromDWtLayer runs the WtFromDWtLayer kernel with given number of elements, +// on either the CPU or GPU depending on the UseGPU variable. +// Can call multiple Run* kernels in a row, which are then all launched +// in the same command submission on the GPU, which is by far the most efficient. +// MUST call RunDone (with optional vars to sync) after all Run calls. +// Alternatively, a single-shot RunOneWtFromDWtLayer call does Run and Done for a +// single run-and-sync case. +func RunWtFromDWtLayer(n int) { + if UseGPU { + RunWtFromDWtLayerGPU(n) + } else { + RunWtFromDWtLayerCPU(n) + } +} + +// RunWtFromDWtLayerGPU runs the WtFromDWtLayer kernel on the GPU. See [RunWtFromDWtLayer] for more info. +func RunWtFromDWtLayerGPU(n int) { + sy := GPUSystem + pl := sy.ComputePipelines["WtFromDWtLayer"] + ce, _ := sy.BeginComputePass() + pl.Dispatch1D(ce, n, 64) +} + +// RunWtFromDWtLayerCPU runs the WtFromDWtLayer kernel on the CPU. +func RunWtFromDWtLayerCPU(n int) { + gpu.VectorizeFunc(0, n, WtFromDWtLayer) +} + +// RunOneWtFromDWtLayer runs the WtFromDWtLayer kernel with given number of elements, +// on either the CPU or GPU depending on the UseGPU variable. +// This version then calls RunDone with the given variables to sync +// after the Run, for a single-shot Run-and-Done call. If multiple kernels +// can be run in sequence, it is much more efficient to do multiple Run* +// calls followed by a RunDone call. +func RunOneWtFromDWtLayer(n int, syncVars ...GPUVars) { + if UseGPU { + RunWtFromDWtLayerGPU(n) + RunDone(syncVars...) + } else { + RunWtFromDWtLayerCPU(n) + } +} // RunWtFromDWtSyn runs the WtFromDWtSyn kernel with given number of elements, // on either the CPU or GPU depending on the UseGPU variable. // Can call multiple Run* kernels in a row, which are then all launched diff --git a/axon/learn-layer.go b/axon/learn-layer.go index 5df23833..7e3ae3c6 100644 --- a/axon/learn-layer.go +++ b/axon/learn-layer.go @@ -8,23 +8,25 @@ package axon import "cogentcore.org/core/math32" +//gosl:start + // DTrgSubMean subtracts the mean from DTrgAvg values // Called by TrgAvgFromD -func (ly *Layer) DTrgSubMean(ctx *Context) { - submean := ly.Params.Learn.TrgAvgAct.SubMean +func (ly *LayerParams) DTrgSubMean(ctx *Context) { + submean := ly.Learn.TrgAvgAct.SubMean if submean == 0 { return } - if ly.Params.HasPoolInhib() && ly.Params.Learn.TrgAvgAct.Pool.IsTrue() { - np := ly.NPools + if ly.HasPoolInhib() && ly.Learn.TrgAvgAct.Pool.IsTrue() { + np := ly.Indexes.NPools for spi := uint32(1); spi < np; spi++ { - pi := ly.Params.PoolIndex(spi) // only for idxs + pi := ly.PoolIndex(spi) nsi := PoolsInt.Value(int(pi), int(0), int(PoolNeurSt)) nei := PoolsInt.Value(int(pi), int(0), int(PoolNeurEd)) nn := 0 avg := float32(0) for lni := nsi; lni < nei; lni++ { - ni := ly.NeurStIndex + uint32(lni) + ni := ly.Indexes.NeurSt + uint32(lni) if NeuronIsOff(ni) { continue } @@ -37,7 +39,7 @@ func (ly *Layer) DTrgSubMean(ctx *Context) { avg /= float32(nn) avg *= submean for lni := nsi; lni < nei; lni++ { - ni := ly.NeurStIndex + uint32(lni) + ni := ly.Indexes.NeurSt + uint32(lni) if NeuronIsOff(ni) { continue } @@ -47,8 +49,9 @@ func (ly *Layer) DTrgSubMean(ctx *Context) { } else { nn := 0 avg := float32(0) - for lni := uint32(0); lni < ly.NNeurons; lni++ { - ni := ly.NeurStIndex + lni + tn := ly.Indexes.NNeurons + for lni := uint32(0); lni < tn; lni++ { + ni := ly.Indexes.NeurSt + lni if NeuronIsOff(ni) { continue } @@ -60,8 +63,8 @@ func (ly *Layer) DTrgSubMean(ctx *Context) { } avg /= float32(nn) avg *= submean - for lni := uint32(0); lni < ly.NNeurons; lni++ { - ni := ly.NeurStIndex + lni + for lni := uint32(0); lni < tn; lni++ { + ni := ly.Indexes.NeurSt + lni if NeuronIsOff(ni) { continue } @@ -71,32 +74,34 @@ func (ly *Layer) DTrgSubMean(ctx *Context) { } // TrgAvgFromD updates TrgAvg from DTrgAvg -- called in PlusPhasePost -func (ly *Layer) TrgAvgFromD(ctx *Context) { - lr := ly.Params.LearnTrgAvgErrLRate() +func (ly *LayerParams) TrgAvgFromD(ctx *Context) { + lr := ly.LearnTrgAvgErrLRate() if lr == 0 { return } ly.DTrgSubMean(ctx) - nn := ly.NNeurons + nn := ly.Indexes.NNeurons for lni := uint32(0); lni < nn; lni++ { - ni := ly.NeurStIndex + lni + ni := ly.Indexes.NeurSt + lni if NeuronIsOff(ni) { continue } ntrg := NeuronAvgs.Value(int(ni), int(TrgAvg)) + NeuronAvgs.Value(int(ni), int(DTrgAvg)) - ntrg = ly.Params.Learn.TrgAvgAct.TrgRange.ClipValue(ntrg) + ntrg = ly.Learn.TrgAvgAct.TrgRange.ClipValue(ntrg) NeuronAvgs.Set(ntrg, int(ni), int(TrgAvg)) - NeuronAvgs.Set(0, int(ni), int(DTrgAvg)) + NeuronAvgs.Set(0.0, int(ni), int(DTrgAvg)) } } // WtFromDWtLayer does weight update at the layer level. // does NOT call main pathway-level WtFromDWt method. // in base, only calls TrgAvgFromD -func (ly *Layer) WtFromDWtLayer(ctx *Context) { +func (ly *LayerParams) WtFromDWtLayer(ctx *Context) { ly.TrgAvgFromD(ctx) } +//gosl:end + // SlowAdapt is the layer-level slow adaptation functions. // Calls AdaptInhib and AvgDifFromTrgAvg for Synaptic Scaling. // Does NOT call pathway-level methods. diff --git a/axon/learn-layer.goal b/axon/learn-layer.goal index f64e6a0d..7847629a 100644 --- a/axon/learn-layer.goal +++ b/axon/learn-layer.goal @@ -6,23 +6,25 @@ package axon import "cogentcore.org/core/math32" +//gosl:start + // DTrgSubMean subtracts the mean from DTrgAvg values // Called by TrgAvgFromD -func (ly *Layer) DTrgSubMean(ctx *Context) { - submean := ly.Params.Learn.TrgAvgAct.SubMean +func (ly *LayerParams) DTrgSubMean(ctx *Context) { + submean := ly.Learn.TrgAvgAct.SubMean if submean == 0 { return } - if ly.Params.HasPoolInhib() && ly.Params.Learn.TrgAvgAct.Pool.IsTrue() { - np := ly.NPools + if ly.HasPoolInhib() && ly.Learn.TrgAvgAct.Pool.IsTrue() { + np := ly.Indexes.NPools for spi := uint32(1); spi < np; spi++ { - pi := ly.Params.PoolIndex(spi) // only for idxs + pi := ly.PoolIndex(spi) nsi := PoolsInt[pi, 0, PoolNeurSt] nei := PoolsInt[pi, 0, PoolNeurEd] nn := 0 avg := float32(0) for lni := nsi; lni < nei; lni++ { - ni := ly.NeurStIndex + uint32(lni) + ni := ly.Indexes.NeurSt + uint32(lni) if NeuronIsOff(ni) { continue } @@ -35,7 +37,7 @@ func (ly *Layer) DTrgSubMean(ctx *Context) { avg /= float32(nn) avg *= submean for lni := nsi; lni < nei; lni++ { - ni := ly.NeurStIndex + uint32(lni) + ni := ly.Indexes.NeurSt + uint32(lni) if NeuronIsOff(ni) { continue } @@ -45,8 +47,9 @@ func (ly *Layer) DTrgSubMean(ctx *Context) { } else { nn := 0 avg := float32(0) - for lni := uint32(0); lni < ly.NNeurons; lni++ { - ni := ly.NeurStIndex + lni + tn := ly.Indexes.NNeurons + for lni := uint32(0); lni < tn; lni++ { + ni := ly.Indexes.NeurSt + lni if NeuronIsOff(ni) { continue } @@ -58,8 +61,8 @@ func (ly *Layer) DTrgSubMean(ctx *Context) { } avg /= float32(nn) avg *= submean - for lni := uint32(0); lni < ly.NNeurons; lni++ { - ni := ly.NeurStIndex + lni + for lni := uint32(0); lni < tn; lni++ { + ni := ly.Indexes.NeurSt + lni if NeuronIsOff(ni) { continue } @@ -69,32 +72,34 @@ func (ly *Layer) DTrgSubMean(ctx *Context) { } // TrgAvgFromD updates TrgAvg from DTrgAvg -- called in PlusPhasePost -func (ly *Layer) TrgAvgFromD(ctx *Context) { - lr := ly.Params.LearnTrgAvgErrLRate() +func (ly *LayerParams) TrgAvgFromD(ctx *Context) { + lr := ly.LearnTrgAvgErrLRate() if lr == 0 { return } ly.DTrgSubMean(ctx) - nn := ly.NNeurons + nn := ly.Indexes.NNeurons for lni := uint32(0); lni < nn; lni++ { - ni := ly.NeurStIndex + lni + ni := ly.Indexes.NeurSt + lni if NeuronIsOff(ni) { continue } ntrg := NeuronAvgs[ni, TrgAvg] + NeuronAvgs[ni, DTrgAvg] - ntrg = ly.Params.Learn.TrgAvgAct.TrgRange.ClipValue(ntrg) + ntrg = ly.Learn.TrgAvgAct.TrgRange.ClipValue(ntrg) NeuronAvgs[ni, TrgAvg] = ntrg - NeuronAvgs[ni, DTrgAvg] = 0 + NeuronAvgs[ni, DTrgAvg] = 0.0 } } // WtFromDWtLayer does weight update at the layer level. // does NOT call main pathway-level WtFromDWt method. // in base, only calls TrgAvgFromD -func (ly *Layer) WtFromDWtLayer(ctx *Context) { +func (ly *LayerParams) WtFromDWtLayer(ctx *Context) { ly.TrgAvgFromD(ctx) } +//gosl:end + // SlowAdapt is the layer-level slow adaptation functions. // Calls AdaptInhib and AvgDifFromTrgAvg for Synaptic Scaling. // Does NOT call pathway-level methods. diff --git a/axon/learn-net.go b/axon/learn-net.go index cdbd9340..453b31f3 100644 --- a/axon/learn-net.go +++ b/axon/learn-net.go @@ -24,6 +24,7 @@ func (nt *Network) DWt() { func (nt *Network) WtFromDWt() { nix := nt.NetIxs() ctx := nt.Context() + RunWtFromDWtLayer(int(nix.NLayers)) RunDWtSubMeanPath(int(nix.NPaths)) RunWtFromDWtSyn(int(nix.NSyns)) RunDoneSynapses() @@ -33,8 +34,8 @@ func (nt *Network) WtFromDWt() { } // DWtToWt computes the weight change (learning) based on current -// running-average activation values, and then WtFromDWt, syncing -// back only the synapses (not SynapseTraces). +// running-average activation values, and then WtFromDWt, +// without syncing any synapse-level state. // This should be used when not viewing the weights. func (nt *Network) DWtToWt() { nix := nt.NetIxs() @@ -53,14 +54,8 @@ func (nt *Network) DWtToWt() { } // SlowAdapt is the layer-level slow adaptation functions: Synaptic scaling, -// and adapting inhibition +// and adapting inhibition. func (nt *Network) SlowAdapt() { - // note: for now doing all this slow stuff CPU-side - // These Sync calls always check if GPU is On - // nt.GPU.SyncAllFromGPU() // todo: - - // todo: convert this to GPU mode - ctx := nt.Context() for _, ly := range nt.Layers { if ly.Off { @@ -71,11 +66,6 @@ func (nt *Network) SlowAdapt() { for _, pt := range nt.Paths { pt.SlowAdapt(ctx) } - // nt.LayerMapSeq(func(ly *Layer) { ly.SlowAdapt(ctx) }, "SlowAdapt") - // nt.PathMapSeq(func(pj *Path) { pj.SlowAdapt(ctx) }, "SlowAdapt") - - // nt.GPU.SyncAllToGPU() - // nt.GPU.SyncSynCaToGPU() // was cleared } // LRateMod sets the LRate modulation parameter for Paths, which is @@ -249,6 +239,13 @@ func DWtFromDiSyn(syni uint32) { //gosl:kernel Paths[pti].DWtFromDi(ctx, syni) } +// WtFromDWtLayer is the kernel over Layers for layer-level Wt update. +// Does TrgAvg updating. +func WtFromDWtLayer(li uint32) { //gosl:kernel + ctx := GetCtx(0) + Layers[li].WtFromDWtLayer(ctx) +} + // DWtSubMeanPath is the kernel over Paths to // compute DWt - mean(DWt). func DWtSubMeanPath(pti uint32) { //gosl:kernel diff --git a/axon/learn-net.goal b/axon/learn-net.goal index e8fcad2a..0d72827d 100644 --- a/axon/learn-net.goal +++ b/axon/learn-net.goal @@ -22,6 +22,7 @@ func (nt *Network) DWt() { func (nt *Network) WtFromDWt() { nix := nt.NetIxs() ctx := nt.Context() + RunWtFromDWtLayer(int(nix.NLayers)) RunDWtSubMeanPath(int(nix.NPaths)) RunWtFromDWtSyn(int(nix.NSyns)) RunDoneSynapses() @@ -32,8 +33,8 @@ func (nt *Network) WtFromDWt() { // DWtToWt computes the weight change (learning) based on current -// running-average activation values, and then WtFromDWt, syncing -// back only the synapses (not SynapseTraces). +// running-average activation values, and then WtFromDWt, +// without syncing any synapse-level state. // This should be used when not viewing the weights. func (nt *Network) DWtToWt() { nix := nt.NetIxs() @@ -52,14 +53,8 @@ func (nt *Network) DWtToWt() { } // SlowAdapt is the layer-level slow adaptation functions: Synaptic scaling, -// and adapting inhibition +// and adapting inhibition. func (nt *Network) SlowAdapt() { - // note: for now doing all this slow stuff CPU-side - // These Sync calls always check if GPU is On - // nt.GPU.SyncAllFromGPU() // todo: - - // todo: convert this to GPU mode - ctx := nt.Context() for _, ly := range nt.Layers { if ly.Off { @@ -70,11 +65,6 @@ func (nt *Network) SlowAdapt() { for _, pt := range nt.Paths { pt.SlowAdapt(ctx) } - // nt.LayerMapSeq(func(ly *Layer) { ly.SlowAdapt(ctx) }, "SlowAdapt") - // nt.PathMapSeq(func(pj *Path) { pj.SlowAdapt(ctx) }, "SlowAdapt") - - // nt.GPU.SyncAllToGPU() - // nt.GPU.SyncSynCaToGPU() // was cleared } // LRateMod sets the LRate modulation parameter for Paths, which is @@ -245,6 +235,13 @@ func DWtFromDiSyn(syni uint32) { //gosl:kernel Paths[pti].DWtFromDi(ctx, syni) } +// WtFromDWtLayer is the kernel over Layers for layer-level Wt update. +// Does TrgAvg updating. +func WtFromDWtLayer(li uint32) { //gosl:kernel + ctx := GetCtx(0) + Layers[li].WtFromDWtLayer(ctx) +} + // DWtSubMeanPath is the kernel over Paths to // compute DWt - mean(DWt). func DWtSubMeanPath(pti uint32) { //gosl:kernel diff --git a/axon/shaders/ApplyExtsNeuron.wgsl b/axon/shaders/ApplyExtsNeuron.wgsl index e1d13bd5..e8dbb646 100644 --- a/axon/shaders/ApplyExtsNeuron.wgsl +++ b/axon/shaders/ApplyExtsNeuron.wgsl @@ -775,6 +775,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/BetweenGi.wgsl b/axon/shaders/BetweenGi.wgsl index a47be7b6..fc1c07e3 100644 --- a/axon/shaders/BetweenGi.wgsl +++ b/axon/shaders/BetweenGi.wgsl @@ -740,6 +740,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/CycleInc.wgsl b/axon/shaders/CycleInc.wgsl index 70c08928..f6fc7f4a 100644 --- a/axon/shaders/CycleInc.wgsl +++ b/axon/shaders/CycleInc.wgsl @@ -719,6 +719,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/CycleNeuron.wgsl b/axon/shaders/CycleNeuron.wgsl index 13d19f96..7438a5e3 100644 --- a/axon/shaders/CycleNeuron.wgsl +++ b/axon/shaders/CycleNeuron.wgsl @@ -1565,6 +1565,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/CyclePost.wgsl b/axon/shaders/CyclePost.wgsl index b9690386..806a29fa 100644 --- a/axon/shaders/CyclePost.wgsl +++ b/axon/shaders/CyclePost.wgsl @@ -853,6 +853,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/DWtFromDiSyn.wgsl b/axon/shaders/DWtFromDiSyn.wgsl index f8337917..f66ab3c7 100644 --- a/axon/shaders/DWtFromDiSyn.wgsl +++ b/axon/shaders/DWtFromDiSyn.wgsl @@ -704,6 +704,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" fn DWtFromDiSyn(syni: u32) { //gosl:kernel var ctx = Ctx[0]; diff --git a/axon/shaders/DWtSubMeanPath.wgsl b/axon/shaders/DWtSubMeanPath.wgsl index ed61ea75..fadc1834 100644 --- a/axon/shaders/DWtSubMeanPath.wgsl +++ b/axon/shaders/DWtSubMeanPath.wgsl @@ -704,6 +704,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" fn DWtSubMeanPath(pti: u32) { //gosl:kernel var ctx = Ctx[0]; diff --git a/axon/shaders/DWtSyn.wgsl b/axon/shaders/DWtSyn.wgsl index 2ed93231..712a7971 100644 --- a/axon/shaders/DWtSyn.wgsl +++ b/axon/shaders/DWtSyn.wgsl @@ -720,6 +720,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" fn DWtSyn(i: u32) { //gosl:kernel var ctx = Ctx[0]; diff --git a/axon/shaders/GPUTestWrite.wgsl b/axon/shaders/GPUTestWrite.wgsl index ee34c3b0..3ced44df 100644 --- a/axon/shaders/GPUTestWrite.wgsl +++ b/axon/shaders/GPUTestWrite.wgsl @@ -719,6 +719,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/GatherSpikes.wgsl b/axon/shaders/GatherSpikes.wgsl index 4922f701..1d58b3b6 100644 --- a/axon/shaders/GatherSpikes.wgsl +++ b/axon/shaders/GatherSpikes.wgsl @@ -820,6 +820,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/InitGBuffsPath.wgsl b/axon/shaders/InitGBuffsPath.wgsl index a0b56c21..bafb76cb 100644 --- a/axon/shaders/InitGBuffsPath.wgsl +++ b/axon/shaders/InitGBuffsPath.wgsl @@ -724,6 +724,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/LayerGi.wgsl b/axon/shaders/LayerGi.wgsl index 8e4638f9..f0fb6bb7 100644 --- a/axon/shaders/LayerGi.wgsl +++ b/axon/shaders/LayerGi.wgsl @@ -813,6 +813,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/MinusPhaseNeuron.wgsl b/axon/shaders/MinusPhaseNeuron.wgsl index 8f1f3af3..583af42b 100644 --- a/axon/shaders/MinusPhaseNeuron.wgsl +++ b/axon/shaders/MinusPhaseNeuron.wgsl @@ -722,6 +722,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/MinusPhasePool.wgsl b/axon/shaders/MinusPhasePool.wgsl index 738c995b..0db71405 100644 --- a/axon/shaders/MinusPhasePool.wgsl +++ b/axon/shaders/MinusPhasePool.wgsl @@ -738,6 +738,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/MinusPhasePost.wgsl b/axon/shaders/MinusPhasePost.wgsl index 15e27313..561cafbb 100644 --- a/axon/shaders/MinusPhasePost.wgsl +++ b/axon/shaders/MinusPhasePost.wgsl @@ -822,6 +822,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/NewStateLayer.wgsl b/axon/shaders/NewStateLayer.wgsl index 2c16d85c..646d9245 100644 --- a/axon/shaders/NewStateLayer.wgsl +++ b/axon/shaders/NewStateLayer.wgsl @@ -794,6 +794,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/NewStateNeuron.wgsl b/axon/shaders/NewStateNeuron.wgsl index d899672e..48f5ea19 100644 --- a/axon/shaders/NewStateNeuron.wgsl +++ b/axon/shaders/NewStateNeuron.wgsl @@ -811,6 +811,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/PlusPhaseNeuron.wgsl b/axon/shaders/PlusPhaseNeuron.wgsl index d79d0a69..7e51f495 100644 --- a/axon/shaders/PlusPhaseNeuron.wgsl +++ b/axon/shaders/PlusPhaseNeuron.wgsl @@ -788,6 +788,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/PlusPhasePool.wgsl b/axon/shaders/PlusPhasePool.wgsl index 30a2b1de..a65ab461 100644 --- a/axon/shaders/PlusPhasePool.wgsl +++ b/axon/shaders/PlusPhasePool.wgsl @@ -721,6 +721,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/PlusPhasePost.wgsl b/axon/shaders/PlusPhasePost.wgsl index a1963652..7e34edcf 100644 --- a/axon/shaders/PlusPhasePost.wgsl +++ b/axon/shaders/PlusPhasePost.wgsl @@ -909,6 +909,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/PlusPhaseStartNeuron.wgsl b/axon/shaders/PlusPhaseStartNeuron.wgsl index 19de8405..f0cc2b40 100644 --- a/axon/shaders/PlusPhaseStartNeuron.wgsl +++ b/axon/shaders/PlusPhaseStartNeuron.wgsl @@ -734,6 +734,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/PoolGi.wgsl b/axon/shaders/PoolGi.wgsl index f6a07432..f46af327 100644 --- a/axon/shaders/PoolGi.wgsl +++ b/axon/shaders/PoolGi.wgsl @@ -819,6 +819,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/SendSpike.wgsl b/axon/shaders/SendSpike.wgsl index 6c087a15..4e110d68 100644 --- a/axon/shaders/SendSpike.wgsl +++ b/axon/shaders/SendSpike.wgsl @@ -951,6 +951,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" ///////////// import: "learn-path.go" diff --git a/axon/shaders/WtFromDWtLayer.wgsl b/axon/shaders/WtFromDWtLayer.wgsl new file mode 100644 index 00000000..276dea7d --- /dev/null +++ b/axon/shaders/WtFromDWtLayer.wgsl @@ -0,0 +1,1487 @@ +// Code generated by "gosl"; DO NOT EDIT +// kernel: WtFromDWtLayer + +// // Layers are all the layer parameters. +@group(0) @binding(0) +var Layers: array; +@group(0) @binding(1) +var Paths: array; +// // NetworkIxs have indexes and sizes for entire network (one only). +@group(1) @binding(0) +var NetworkIxs: array; +@group(1) @binding(1) +var NeuronIxs: array; +@group(1) @binding(2) +var SynapseIxs: array; +@group(1) @binding(3) +var PathSendCon: array; +@group(1) @binding(4) +var RecvPathIxs: array; +@group(1) @binding(5) +var PathRecvCon: array; +@group(1) @binding(6) +var RecvSynIxs: array; +// // Ctx is the current context state (one only). +@group(2) @binding(0) +var Ctx: array; +@group(2) @binding(1) +var Neurons: array; +@group(2) @binding(2) +var NeuronAvgs: array; +@group(2) @binding(3) +var LayerStates: array; +@group(2) @binding(4) +var GlobalScalars: array; +@group(2) @binding(5) +var GlobalVectors: array; +@group(2) @binding(6) +var Exts: array; +// // Pools are the [PoolVars] float32 state values for layer and sub-pool inhibition, // Including the float32 AvgMax values by Phase and variable: use [AvgMaxVarIndex]. // [Layer * Pools][PoolVars+AvgMax][Data] +@group(3) @binding(0) +var Pools: array; +@group(3) @binding(1) +var PoolsInt: array; +@group(3) @binding(2) +var PathGBuf: array; +@group(3) @binding(3) +var PathGSyns: array; +@group(3) @binding(4) +var Synapses: array; +@group(3) @binding(5) +var SynapseTraces: array; + +alias GPUVars = i32; + +@compute @workgroup_size(64, 1, 1) +fn main(@builtin(global_invocation_id) idx: vec3) { + WtFromDWtLayer(idx.x); +} + +fn IndexU322D(s0: u32, s1: u32, i0: u32, i1: u32) -> u32 { + return u32(2) + s0 * i0 + s1 * i1; +} + +fn IndexU321D(s0: u32, i0: u32) -> u32 { + return u32(1) + s0 * i0; +} + +fn IndexF323D(s0: f32, s1: f32, s2: f32, i0: u32, i1: u32, i2: u32) -> u32 { + return u32(3) + bitcast(s0) * i0 + bitcast(s1) * i1 + bitcast(s2) * i2; +} + +fn IndexF322D(s0: f32, s1: f32, i0: u32, i1: u32) -> u32 { + return u32(2) + bitcast(s0) * i0 + bitcast(s1) * i1; +} + +fn IndexI323D(s0: i32, s1: i32, s2: i32, i0: u32, i1: u32, i2: u32) -> u32 { + return u32(3) + u32(s0) * i0 + u32(s1) * i1 + u32(s2) * i2; +} + + +///////////// import: "vars.go" + +///////////// import: "act-layer.go" +fn LayerParams_IsLearnTrgAvg(ly: ptr) -> bool { + if ((*ly).Acts.Clamp.IsInput == 1 || (*ly).Acts.Clamp.IsTarget == 1 || (*ly).Learn.TrgAvgAct.RescaleOn == 0) { + return false; + }return true; +} +fn LayerParams_LearnTrgAvgErrLRate(ly: ptr) -> f32 { + if (!LayerParams_IsLearnTrgAvg(ly)) { + return f32(0); + }return (*ly).Learn.TrgAvgAct.ErrLRate; +} + +///////////// import: "act-net.go" + +///////////// import: "act-path.go" +alias PathGTypes = i32; //enums:enum +const ExcitatoryG: PathGTypes = 0; +const InhibitoryG: PathGTypes = 1; +const ModulatoryG: PathGTypes = 2; +const MaintG: PathGTypes = 3; +const ContextG: PathGTypes = 4; +struct SynComParams { + GType: PathGTypes, + Delay: u32, + MaxDelay: u32, + DelLen: u32, +} +struct PathScaleParams { + Rel: f32, + Abs: f32, + pad: f32, + pad1: f32, +} + +///////////// import: "act.go" +fn NeuronHasFlag(flag: NeuronFlags, ni: u32,di: u32) -> bool { + return (NeuronFlags(bitcast(Neurons[IndexF323D(Neurons[0], Neurons[1], Neurons[2], u32(ni),u32(di),u32(NeurFlags))])) & flag) > 0; // weird: != 0 does NOT work on GPU +} +fn NeuronIsOff(ni: u32) -> bool { + return NeuronHasFlag(NeuronOff, ni, u32(u32(0))); +} +struct SpikeParams { + Thr: f32, + VmR: f32, + Tr: i32, + RTau: f32, + Exp: i32, + ExpSlope: f32, + ExpThr: f32, + MaxHz: f32, + ISITau: f32, + ISIDt: f32, + RDt: f32, + pad: i32, +} +struct DendParams { + GbarExp: f32, + GbarR: f32, + SSGi: f32, + HasMod: i32, + ModGain: f32, + ModACh: i32, + ModBase: f32, + pad: i32, +} +struct ActInitParams { + Vm: f32, + Act: f32, + GeBase: f32, + GiBase: f32, + GeVar: f32, + GiVar: f32, + pad: i32, + pad1: i32, +} +struct DecayParams { + Act: f32, + Glong: f32, + AHP: f32, + LearnCa: f32, + OnRew: i32, + pad: f32, + pad1: f32, + pad2: f32, +} +struct DtParams { + Integ: f32, + VmTau: f32, + VmDendTau: f32, + VmSteps: i32, + GeTau: f32, + GiTau: f32, + IntTau: f32, + LongAvgTau: f32, + MaxCycStart: i32, + VmDt: f32, + VmDendDt: f32, + DtStep: f32, + GeDt: f32, + GiDt: f32, + IntDt: f32, + LongAvgDt: f32, +} +struct SpikeNoiseParams { + On: i32, + GeHz: f32, + Ge: f32, + GiHz: f32, + Gi: f32, + MaintGe: i32, + GeExpInt: f32, + GiExpInt: f32, +} +struct ClampParams { + IsInput: i32, + IsTarget: i32, + Ge: f32, + Add: i32, + ErrThr: f32, + pad: f32, + pad1: f32, + pad2: f32, +} +struct SMaintParams { + On: i32, + NNeurons: f32, + Gbar: f32, + Inhib: f32, + ISI: F32, +} +struct PopCodeParams { + On: i32, + Ge: f32, + Min: f32, + Max: f32, + MinAct: f32, + MinSigma: f32, + MaxSigma: f32, + Clip: i32, +} +struct ActParams { + Spikes: SpikeParams, + Dend: DendParams, + Init: ActInitParams, + Decay: DecayParams, + Dt: DtParams, + Gbar: Chans, + Erev: Chans, + Clamp: ClampParams, + Noise: SpikeNoiseParams, + VmRange: F32, + Mahp: MahpParams, + Sahp: SahpParams, + KNa: KNaMedSlow, + Kir: KirParams, + NMDA: NMDAParams, + MaintNMDA: NMDAParams, + GabaB: GABABParams, + VGCC: VGCCParams, + AK: AKsParams, + SKCa: SKCaParams, + SMaint: SMaintParams, + PopCode: PopCodeParams, +} + +///////////// import: "chans-ak.go" +struct AKsParams { + Gbar: f32, + Hf: f32, + Mf: f32, + Voff: f32, + Vmax: f32, + pad: i32, + pad1: i32, + pad2: i32, +} + +///////////// import: "chans-chans.go" +struct Chans { + E: f32, + L: f32, + I: f32, + K: f32, +} + +///////////// import: "chans-gabab.go" +struct GABABParams { + Gbar: f32, + RiseTau: f32, + DecayTau: f32, + Gbase: f32, + GiSpike: f32, + MaxTime: f32, + TauFact: f32, + RiseDt: f32, + DecayDt: f32, + pad: f32, + pad1: f32, + pad2: f32, +} + +///////////// import: "chans-kir.go" +struct KirParams { + Gbar: f32, + MinfOff: f32, + MinfTau: f32, + RiseOff: f32, + RiseTau: f32, + DecayOff: f32, + DecayTau: f32, + Mrest: f32, +} + +///////////// import: "chans-kna.go" +struct KNaParams { + On: i32, + Rise: f32, + Max: f32, + Tau: f32, + Dt: f32, + pad: i32, + pad1: i32, + pad2: i32, +} +struct KNaMedSlow { + On: i32, + TrialSlow: i32, + pad: i32, + pad1: i32, + Med: KNaParams, + Slow: KNaParams, +} + +///////////// import: "chans-mahp.go" +struct MahpParams { + Gbar: f32, + Voff: f32, + Vslope: f32, + TauMax: f32, + Tadj: f32, + DtMax: f32, + pad: i32, + pad2: i32, +} + +///////////// import: "chans-nmda.go" +struct NMDAParams { + Gbar: f32, + Tau: f32, + ITau: f32, + MgC: f32, + Voff: f32, + Dt: f32, + IDt: f32, + MgFact: f32, +} + +///////////// import: "chans-sahp.go" +struct SahpParams { + Gbar: f32, + CaTau: f32, + Off: f32, + Slope: f32, + TauMax: f32, + CaDt: f32, + DtMax: f32, + pad: i32, +} + +///////////// import: "chans-skca.go" +struct SKCaParams { + Gbar: f32, + C50: f32, + ActTau: f32, + DeTau: f32, + KCaR: f32, + CaRDecayTau: f32, + CaInThr: f32, + CaInTau: f32, + ActDt: f32, + DeDt: f32, + CaRDecayDt: f32, + CaInDt: f32, +} + +///////////// import: "chans-vgcc.go" +struct VGCCParams { + Gbar: f32, + Ca: f32, + pad: i32, + pad1: i32, +} + +///////////// import: "context.go" +struct Context { + NData: u32, + Mode: i32, + Testing: i32, + Phase: i32, + PlusPhase: i32, + PhaseCycle: i32, + Cycle: i32, + ThetaCycles: i32, + CyclesTotal: i32, + Time: f32, + TrialsTotal: i32, + TimePerCycle: f32, + SlowInterval: i32, + SlowCounter: i32, + pad: i32, + pad1: i32, + RandCounter: RandCounter, +} + +///////////// import: "deep-layer.go" +struct BurstParams { + ThrRel: f32, + ThrAbs: f32, + pad: f32, + pad1: f32, +} +struct CTParams { + GeGain: f32, + DecayTau: f32, + OFCposPT: i32, + DecayDt: f32, +} +struct PulvParams { + DriveScale: f32, + FullDriveAct: f32, + DriveLayIndex: i32, + pad: f32, +} + +///////////// import: "deep-path.go" + +///////////// import: "enumgen.go" +const PathGTypesN: PathGTypes = 5; +const GlobalScalarVarsN: GlobalScalarVars = 57; +const GlobalVectorVarsN: GlobalVectorVars = 10; +const GPUVarsN: GPUVars = 22; +const LayerTypesN: LayerTypes = 30; +const LayerVarsN: LayerVars = 11; +const ViewTimesN: ViewTimes = 7; +const DAModTypesN: DAModTypes = 4; +const ValenceTypesN: ValenceTypes = 3; +const NeuronFlagsN: NeuronFlags = 9; +const NeuronVarsN: NeuronVars = 90; +const NeuronAvgVarsN: NeuronAvgVars = 7; +const NeuronIndexVarsN: NeuronIndexVars = 3; +const PathTypesN: PathTypes = 12; +const GPLayerTypesN: GPLayerTypes = 3; +const PoolIntVarsN: PoolIntVars = 10; +const AvgMaxN: AvgMax = 2; +const AvgMaxPhasesN: AvgMaxPhases = 4; +const AvgMaxVarsN: AvgMaxVars = 7; +const SynapseVarsN: SynapseVars = 5; +const SynapseTraceVarsN: SynapseTraceVars = 3; +const SynapseIndexVarsN: SynapseIndexVars = 3; + +///////////// import: "fsfffb-enumgen.go" +const InhibVarsN: InhibVars = 16; + +///////////// import: "fsfffb-fsfffb.go" +struct GiParams { + On: i32, + Gi: f32, + FB: f32, + FSTau: f32, + SS: f32, + SSfTau: f32, + SSiTau: f32, + FS0: f32, + FFAvgTau: f32, + FFPrv: f32, + ClampExtMin: f32, + FSDt: f32, + SSfDt: f32, + SSiDt: f32, + FFAvgDt: f32, + pad: f32, +} + +///////////// import: "fsfffb-inhib.go" +alias InhibVars = i32; //enums:enum +const FFsRaw: InhibVars = 0; +const FBsRaw: InhibVars = 1; +const GeExtRaw: InhibVars = 2; +const FFs: InhibVars = 3; +const FBs: InhibVars = 4; +const GeExts: InhibVars = 5; +const FSi: InhibVars = 6; +const SSi: InhibVars = 7; +const SSf: InhibVars = 8; +const FSGi: InhibVars = 9; +const SSGi: InhibVars = 10; +const TotalGi: InhibVars = 11; +const GiOrig: InhibVars = 12; +const LayGi: InhibVars = 13; +const FFAvg: InhibVars = 14; +const FFAvgPrv: InhibVars = 15; + +///////////// import: "globals.go" +alias GlobalScalarVars = i32; //enums:enum +const GvRew: GlobalScalarVars = 0; +const GvHasRew: GlobalScalarVars = 1; +const GvRewPred: GlobalScalarVars = 2; +const GvPrevPred: GlobalScalarVars = 3; +const GvHadRew: GlobalScalarVars = 4; +const GvDA: GlobalScalarVars = 5; +const GvDAtonic: GlobalScalarVars = 6; +const GvACh: GlobalScalarVars = 7; +const GvNE: GlobalScalarVars = 8; +const GvSer: GlobalScalarVars = 9; +const GvAChRaw: GlobalScalarVars = 10; +const GvGoalMaint: GlobalScalarVars = 11; +const GvVSMatrixJustGated: GlobalScalarVars = 12; +const GvVSMatrixHasGated: GlobalScalarVars = 13; +const GvCuriosityPoolGated: GlobalScalarVars = 14; +const GvTime: GlobalScalarVars = 15; +const GvEffort: GlobalScalarVars = 16; +const GvUrgencyRaw: GlobalScalarVars = 17; +const GvUrgency: GlobalScalarVars = 18; +const GvHasPosUS: GlobalScalarVars = 19; +const GvHadPosUS: GlobalScalarVars = 20; +const GvNegUSOutcome: GlobalScalarVars = 21; +const GvHadNegUSOutcome: GlobalScalarVars = 22; +const GvPVposSum: GlobalScalarVars = 23; +const GvPVpos: GlobalScalarVars = 24; +const GvPVnegSum: GlobalScalarVars = 25; +const GvPVneg: GlobalScalarVars = 26; +const GvPVposEst: GlobalScalarVars = 27; +const GvPVposVar: GlobalScalarVars = 28; +const GvPVnegEst: GlobalScalarVars = 29; +const GvPVnegVar: GlobalScalarVars = 30; +const GvGoalDistEst: GlobalScalarVars = 31; +const GvGoalDistPrev: GlobalScalarVars = 32; +const GvProgressRate: GlobalScalarVars = 33; +const GvGiveUpUtility: GlobalScalarVars = 34; +const GvContUtility: GlobalScalarVars = 35; +const GvGiveUpTiming: GlobalScalarVars = 36; +const GvContTiming: GlobalScalarVars = 37; +const GvGiveUpProgress: GlobalScalarVars = 38; +const GvContProgress: GlobalScalarVars = 39; +const GvGiveUpSum: GlobalScalarVars = 40; +const GvContSum: GlobalScalarVars = 41; +const GvGiveUpProb: GlobalScalarVars = 42; +const GvGiveUp: GlobalScalarVars = 43; +const GvGaveUp: GlobalScalarVars = 44; +const GvVSPatchPos: GlobalScalarVars = 45; +const GvVSPatchPosThr: GlobalScalarVars = 46; +const GvVSPatchPosRPE: GlobalScalarVars = 47; +const GvVSPatchPosSum: GlobalScalarVars = 48; +const GvVSPatchPosPrev: GlobalScalarVars = 49; +const GvVSPatchPosVar: GlobalScalarVars = 50; +const GvLHbDip: GlobalScalarVars = 51; +const GvLHbBurst: GlobalScalarVars = 52; +const GvLHbPVDA: GlobalScalarVars = 53; +const GvCeMpos: GlobalScalarVars = 54; +const GvCeMneg: GlobalScalarVars = 55; +const GvVtaDA: GlobalScalarVars = 56; +const MaxGlobalVecN = 16; +alias GlobalVectorVars = i32; //enums:enum +const GvCost: GlobalVectorVars = 0; +const GvCostRaw: GlobalVectorVars = 1; +const GvUSneg: GlobalVectorVars = 2; +const GvUSnegRaw: GlobalVectorVars = 3; +const GvDrives: GlobalVectorVars = 4; +const GvUSpos: GlobalVectorVars = 5; +const GvVSPatchD1: GlobalVectorVars = 6; +const GvVSPatchD2: GlobalVectorVars = 7; +const GvOFCposPTMaint: GlobalVectorVars = 8; +const GvVSMatrixPoolGated: GlobalVectorVars = 9; + +///////////// import: "hip_paths.go" +struct HipPathParams { + Hebb: f32, + Err: f32, + SAvgCor: f32, + SAvgThr: f32, + SNominal: f32, + pad: f32, + pad1: f32, + pad2: f32, +} + +///////////// import: "inhib.go" +struct ActAvgParams { + Nominal: f32, + AdaptGi: i32, + Offset: f32, + HiTol: f32, + LoTol: f32, + AdaptRate: f32, + pad: f32, + pad1: f32, +} +struct InhibParams { + ActAvg: ActAvgParams, + Layer: GiParams, + Pool: GiParams, +} + +///////////// import: "init-layer.go" + +///////////// import: "kinase-params.go" +struct CaDtParams { //types:add + MTau: f32, + PTau: f32, + DTau: f32, + MDt: f32, + PDt: f32, + DDt: f32, + pad: i32, + pad1: i32, +} +struct NeurCaParams { + SpikeG: f32, + SynTau: f32, + SynDt: f32, + pad: i32, + Dt: CaDtParams, +} +struct SynCaParams { //types:add + CaScale: f32, + pad: i32, + pad1: i32, + pad2: i32, + Dt: CaDtParams, +} +struct BinWeights { //types:add + Bin0: f32, + Bin1: f32, + Bin2: f32, + Bin3: f32, + Bin4: f32, + Bin5: f32, + Bin6: f32, + Bin7: f32, +} +struct SynCaLinear { //types:add + CaP: BinWeights, + CaD: BinWeights, + CaGain: f32, + pad: f32, + pad1: f32, + pad2: f32, +} + +///////////// import: "layerparams.go" +struct LayerIndexes { + NPools: u32, + NeurSt: u32, + NNeurons: u32, + RecvSt: u32, + RecvN: u32, + SendSt: u32, + SendN: u32, + ExtsSt: u32, + ShpPlY: i32, + ShpPlX: i32, + ShpUnY: i32, + ShpUnX: i32, +} +struct LayerInhibIndexes { + Index1: i32, + Index2: i32, + Index3: i32, + Index4: i32, +} +struct LayerParams { + Type: LayerTypes, + Index: u32, + MaxData: u32, + PoolSt: u32, + Acts: ActParams, + Inhib: InhibParams, + LayInhib: LayerInhibIndexes, + Learn: LearnNeurParams, + Bursts: BurstParams, + CT: CTParams, + Pulv: PulvParams, + Matrix: MatrixParams, + GP: GPParams, + LDT: LDTParams, + VTA: VTAParams, + RWPred: RWPredParams, + RWDa: RWDaParams, + TDInteg: TDIntegParams, + TDDa: TDDaParams, + Indexes: LayerIndexes, +} +fn LayerParams_PoolIndex(ly: ptr, pi: u32) -> u32 { + return (*ly).PoolSt + pi; +} +fn LayerParams_HasPoolInhib(ly: ptr) -> bool { + return (*ly).Inhib.Pool.On == 1; +} + +///////////// import: "layertypes.go" +alias LayerTypes = i32; //enums:enum +const SuperLayer: LayerTypes = 0; +const InputLayer: LayerTypes = 1; +const TargetLayer: LayerTypes = 2; +const CompareLayer: LayerTypes = 3; +const CTLayer: LayerTypes = 4; +const PulvinarLayer: LayerTypes = 5; +const TRNLayer: LayerTypes = 6; +const PTMaintLayer: LayerTypes = 7; +const PTPredLayer: LayerTypes = 8; +const MatrixLayer: LayerTypes = 9; +const STNLayer: LayerTypes = 10; +const GPLayer: LayerTypes = 11; +const BGThalLayer: LayerTypes = 12; +const VSGatedLayer: LayerTypes = 13; +const BLALayer: LayerTypes = 14; +const CeMLayer: LayerTypes = 15; +const VSPatchLayer: LayerTypes = 16; +const LHbLayer: LayerTypes = 17; +const DrivesLayer: LayerTypes = 18; +const UrgencyLayer: LayerTypes = 19; +const USLayer: LayerTypes = 20; +const PVLayer: LayerTypes = 21; +const LDTLayer: LayerTypes = 22; +const VTALayer: LayerTypes = 23; +const RewLayer: LayerTypes = 24; +const RWPredLayer: LayerTypes = 25; +const RWDaLayer: LayerTypes = 26; +const TDPredLayer: LayerTypes = 27; +const TDIntegLayer: LayerTypes = 28; +const TDDaLayer: LayerTypes = 29; + +///////////// import: "layervars.go" +alias LayerVars = i32; //enums:enum +const LayerActMAvg: LayerVars = 0; +const LayerActPAvg: LayerVars = 1; +const LayerAvgMaxGeM: LayerVars = 2; +const LayerAvgMaxGiM: LayerVars = 3; +const LayerGiMult: LayerVars = 4; +const LayerPhaseDiff: LayerVars = 5; +const LayerPhaseDiffAvg: LayerVars = 6; +const LayerPhaseDiffVar: LayerVars = 7; +const LayerRT: LayerVars = 8; +const LayerRewPredPos: LayerVars = 9; +const LayerRewPredNeg: LayerVars = 10; + +///////////// import: "learn-layer.go" +fn LayerParams_DTrgSubMean(ly: ptr, ctx: ptr) { + var submean = (*ly).Learn.TrgAvgAct.SubMean; + if (submean == 0) { + return; + } + if (LayerParams_HasPoolInhib(ly) && (*ly).Learn.TrgAvgAct.Pool == 1) { + var np = (*ly).Indexes.NPools; + for (var spi = u32(1); spi < np; spi++) { + var pi = LayerParams_PoolIndex(ly, spi); + var nsi = PoolsInt[IndexI323D(PoolsInt[0], PoolsInt[1], PoolsInt[2], u32(pi),u32(0),u32(PoolNeurSt))]; + var nei = PoolsInt[IndexI323D(PoolsInt[0], PoolsInt[1], PoolsInt[2], u32(pi),u32(0),u32(PoolNeurEd))]; + var nn = 0; + var avg = f32(0); + for (var lni = nsi; lni < nei; lni++) { + var ni = (*ly).Indexes.NeurSt + u32(lni); + if (NeuronIsOff(ni)) { + continue; + } + avg += NeuronAvgs[IndexF322D(NeuronAvgs[0], NeuronAvgs[1], u32(ni),u32(DTrgAvg))]; + nn++; + } + if (nn == 0) { + continue; + } + avg /= f32(nn); + avg *= submean; + for (var lni = nsi; lni < nei; lni++) { + var ni = (*ly).Indexes.NeurSt + u32(lni); + if (NeuronIsOff(ni)) { + continue; + } + NeuronAvgs[IndexF322D(NeuronAvgs[0], NeuronAvgs[1], u32(ni),u32(DTrgAvg))] -= avg; + } + } + } else { + var nn = 0; + var avg = f32(0); + var tn = (*ly).Indexes.NNeurons; + for (var lni = u32(0); lni < tn; lni++) { + var ni = (*ly).Indexes.NeurSt + lni; + if (NeuronIsOff(ni)) { + continue; + } + avg += NeuronAvgs[IndexF322D(NeuronAvgs[0], NeuronAvgs[1], u32(ni),u32(DTrgAvg))]; + nn++; + } + if (nn == 0) { + return; + } + avg /= f32(nn); + avg *= submean; + for (var lni = u32(0); lni < tn; lni++) { + var ni = (*ly).Indexes.NeurSt + lni; + if (NeuronIsOff(ni)) { + continue; + } + NeuronAvgs[IndexF322D(NeuronAvgs[0], NeuronAvgs[1], u32(ni),u32(DTrgAvg))] -= avg; + } + } +} +fn LayerParams_TrgAvgFromD(ly: ptr, ctx: ptr) { + var lr = LayerParams_LearnTrgAvgErrLRate(ly); + if (lr == 0) { + return; + } + LayerParams_DTrgSubMean(ly, ctx); + var nn = (*ly).Indexes.NNeurons; + for (var lni = u32(0); lni < nn; lni++) { + var ni = (*ly).Indexes.NeurSt + lni; + if (NeuronIsOff(ni)) { + continue; + } + var ntrg = NeuronAvgs[IndexF322D(NeuronAvgs[0], NeuronAvgs[1], u32(ni),u32(TrgAvg))] + NeuronAvgs[IndexF322D(NeuronAvgs[0], NeuronAvgs[1], u32(ni),u32(DTrgAvg))]; + ntrg = F32_ClipValue(&(*ly).Learn.TrgAvgAct.TrgRange, ntrg); + NeuronAvgs[IndexF322D(NeuronAvgs[0], NeuronAvgs[1], u32(ni),u32(TrgAvg))] = ntrg; + NeuronAvgs[IndexF322D(NeuronAvgs[0], NeuronAvgs[1], + u32(ni),u32(DTrgAvg))] = 0.0; + } +} +fn LayerParams_WtFromDWtLayer(ly: ptr, ctx: ptr) { + LayerParams_TrgAvgFromD(ly, ctx); +} + +///////////// import: "learn-net.go" +fn WtFromDWtLayer(li: u32) { //gosl:kernel + var ctx = Ctx[0]; + var layers=Layers[li]; LayerParams_WtFromDWtLayer(&layers, &ctx); + Ctx[0] = ctx; +} + +///////////// import: "learn-path.go" + +///////////// import: "learn.go" +struct CaLrnParams { + Norm: f32, + SpkVGCC: i32, + SpkVgccCa: f32, + VgccTau: f32, + Dt: CaDtParams, + UpdateThr: f32, + VgccDt: f32, + NormInv: f32, + pad: i32, +} +struct TrgAvgActParams { + GiBaseInit: f32, + RescaleOn: i32, + ErrLRate: f32, + SynScaleRate: f32, + SubMean: f32, + Permute: i32, + Pool: i32, + pad: i32, + TrgRange: F32, +} +struct RLRateParams { + On: i32, + SigmoidLinear: i32, + SigmoidMin: f32, + Diff: i32, + SpkThr: f32, + DiffThr: f32, + Min: f32, + pad: i32, +} +struct LearnNeurParams { + CaLearn: CaLrnParams, + CaSpk: NeurCaParams, + LrnNMDA: NMDAParams, + TrgAvgAct: TrgAvgActParams, + RLRate: RLRateParams, + NeuroMod: NeuroModParams, +} +struct SWtInitParams { + SPct: f32, + Mean: f32, + Var: f32, + Sym: i32, +} +struct SWtAdaptParams { + On: i32, + LRate: f32, + SubMean: f32, + SigGain: f32, +} +struct SWtParams { + Init: SWtInitParams, + Adapt: SWtAdaptParams, + Limit: F32, +} +struct LRateParams { + Base: f32, + Sched: f32, + Mod: f32, + Eff: f32, +} +struct TraceParams { + Tau: f32, + SubMean: f32, + LearnThr: f32, + Dt: f32, +} +struct LRateMod { + On: i32, + Base: f32, + pad: i32, + pad1: i32, + Range: F32, +} +struct HebbParams { + On: i32, + Up: f32, + Down: f32, + pad: f32, +} +struct LearnSynParams { + Learn: i32, + pad: i32, + pad1: i32, + pad2: i32, + LRate: LRateParams, + Trace: TraceParams, + KinaseCa: SynCaLinear, + Hebb: HebbParams, +} + +///////////// import: "looper.go" +alias ViewTimes = i32; //enums:enum +const Cycle: ViewTimes = 0; +const FastSpike: ViewTimes = 1; +const Gamma: ViewTimes = 2; +const Beta: ViewTimes = 3; +const Alpha: ViewTimes = 4; +const Phase: ViewTimes = 5; +const Theta: ViewTimes = 6; + +///////////// import: "math32-fastexp.go" + +///////////// import: "minmax-avgmax.go" +const MaxFloat32: f32 = 3.402823466e+38; +const MinFloat32: f32 = 1.175494351e-38; +struct AvgMax32 { + Avg: f32, + Max: f32, + Sum: f32, + MaxIndex: i32, + N: i32, + pad: i32, + pad1: i32, + pad2: i32, +} + +///////////// import: "minmax-minmax32.go" +struct F32 { + Min: f32, + Max: f32, + pad: i32, + pad1: i32, // for gpu use +} +fn F32_ClipValue(mr: ptr, val: f32) -> f32 { + if (val < (*mr).Min) { + return (*mr).Min; + } + if (val > (*mr).Max) { + return (*mr).Max; + }return val; +} + +///////////// import: "network.go" +struct NetworkIndexes { + MaxData: u32, + MaxDelay: u32, + NLayers: u32, + NNeurons: u32, + NPools: u32, + NPaths: u32, + NSyns: u32, + RubiconNPosUSs: u32, + RubiconNCosts: u32, + RubiconNNegUSs: u32, + GPUMaxBuffFloats: u32, + GPUSynCaBanks: u32, +} + +///////////// import: "neuromod.go" +alias DAModTypes = i32; //enums:enum +const NoDAMod: DAModTypes = 0; +const D1Mod: DAModTypes = 1; +const D2Mod: DAModTypes = 2; +const D1AbsMod: DAModTypes = 3; +alias ValenceTypes = i32; //enums:enum +const Positive: ValenceTypes = 0; +const Negative: ValenceTypes = 1; +const Cost: ValenceTypes = 2; +struct NeuroModParams { + DAMod: DAModTypes, + Valence: ValenceTypes, + DAModGain: f32, + DALRateSign: i32, + DALRateMod: f32, + AChLRateMod: f32, + AChDisInhib: f32, + BurstGain: f32, + DipGain: f32, + pad: f32, + pad1: f32, + pad2: f32, +} + +///////////// import: "neuron.go" +alias NeuronFlags = i32; //enums:enum +const NeuronOff: NeuronFlags = 1; +const NeuronHasExt: NeuronFlags = 2; +const NeuronHasTarg: NeuronFlags = 4; +const NeuronHasCmpr: NeuronFlags = 8; +alias NeuronVars = i32; //enums:enum +const Spike: NeuronVars = 0; +const Spiked: NeuronVars = 1; +const Act: NeuronVars = 2; +const ActInt: NeuronVars = 3; +const Ge: NeuronVars = 4; +const Gi: NeuronVars = 5; +const Gk: NeuronVars = 6; +const Inet: NeuronVars = 7; +const Vm: NeuronVars = 8; +const VmDend: NeuronVars = 9; +const ISI: NeuronVars = 10; +const ISIAvg: NeuronVars = 11; +const Ext: NeuronVars = 12; +const Target: NeuronVars = 13; +const CaSpkM: NeuronVars = 14; +const CaSpkP: NeuronVars = 15; +const CaSpkD: NeuronVars = 16; +const CaSpkPM: NeuronVars = 17; +const CaLrn: NeuronVars = 18; +const NrnCaM: NeuronVars = 19; +const NrnCaP: NeuronVars = 20; +const NrnCaD: NeuronVars = 21; +const CaDiff: NeuronVars = 22; +const RLRate: NeuronVars = 23; +const GnmdaSyn: NeuronVars = 24; +const Gnmda: NeuronVars = 25; +const GnmdaLrn: NeuronVars = 26; +const GnmdaMaint: NeuronVars = 27; +const NmdaCa: NeuronVars = 28; +const Gvgcc: NeuronVars = 29; +const VgccM: NeuronVars = 30; +const VgccH: NeuronVars = 31; +const VgccCa: NeuronVars = 32; +const VgccCaInt: NeuronVars = 33; +const Burst: NeuronVars = 34; +const BurstPrv: NeuronVars = 35; +const CtxtGe: NeuronVars = 36; +const CtxtGeRaw: NeuronVars = 37; +const CtxtGeOrig: NeuronVars = 38; +const GgabaB: NeuronVars = 39; +const GABAB: NeuronVars = 40; +const GABABx: NeuronVars = 41; +const Gak: NeuronVars = 42; +const SSGiDend: NeuronVars = 43; +const GknaMed: NeuronVars = 44; +const GknaSlow: NeuronVars = 45; +const Gkir: NeuronVars = 46; +const KirM: NeuronVars = 47; +const Gsk: NeuronVars = 48; +const SKCaIn: NeuronVars = 49; +const SKCaR: NeuronVars = 50; +const SKCaM: NeuronVars = 51; +const Gmahp: NeuronVars = 52; +const MahpN: NeuronVars = 53; +const Gsahp: NeuronVars = 54; +const SahpCa: NeuronVars = 55; +const SahpN: NeuronVars = 56; +const ActM: NeuronVars = 57; +const ActP: NeuronVars = 58; +const SpkSt1: NeuronVars = 59; +const SpkSt2: NeuronVars = 60; +const SpkMax: NeuronVars = 61; +const SpkMaxCa: NeuronVars = 62; +const SpkBin0: NeuronVars = 63; +const SpkBin1: NeuronVars = 64; +const SpkBin2: NeuronVars = 65; +const SpkBin3: NeuronVars = 66; +const SpkBin4: NeuronVars = 67; +const SpkBin5: NeuronVars = 68; +const SpkBin6: NeuronVars = 69; +const SpkBin7: NeuronVars = 70; +const SpkPrv: NeuronVars = 71; +const GeNoise: NeuronVars = 72; +const GeNoiseP: NeuronVars = 73; +const GiNoise: NeuronVars = 74; +const GiNoiseP: NeuronVars = 75; +const GeExt: NeuronVars = 76; +const GeRaw: NeuronVars = 77; +const GeSyn: NeuronVars = 78; +const GiRaw: NeuronVars = 79; +const GiSyn: NeuronVars = 80; +const GeInt: NeuronVars = 81; +const GeIntNorm: NeuronVars = 82; +const GiInt: NeuronVars = 83; +const GModRaw: NeuronVars = 84; +const GModSyn: NeuronVars = 85; +const SMaintP: NeuronVars = 86; +const GMaintRaw: NeuronVars = 87; +const GMaintSyn: NeuronVars = 88; +const NeurFlags: NeuronVars = 89; +alias NeuronAvgVars = i32; //enums:enum +const ActAvg: NeuronAvgVars = 0; +const AvgPct: NeuronAvgVars = 1; +const TrgAvg: NeuronAvgVars = 2; +const DTrgAvg: NeuronAvgVars = 3; +const AvgDif: NeuronAvgVars = 4; +const GeBase: NeuronAvgVars = 5; +const GiBase: NeuronAvgVars = 6; +alias NeuronIndexVars = i32; //enums:enum +const NrnNeurIndex: NeuronIndexVars = 0; +const NrnLayIndex: NeuronIndexVars = 1; +const NrnSubPool: NeuronIndexVars = 2; + +///////////// import: "pathparams.go" +const StartOff: i32 = 0; +const Nitems: i32 = 1; +const StartNN: i32 = 2; +struct StartN { + Start: u32, + N: u32, + pad: u32, + pad1: u32, // todo: see if we can do without these? +} +struct PathIndexes { + RecvLayer: u32, + RecvNeurSt: u32, + RecvNeurN: u32, + SendLayer: u32, + SendNeurSt: u32, + SendNeurN: u32, + SynapseSt: u32, + SendConSt: u32, + RecvConSt: u32, + RecvSynSt: u32, + NPathNeurSt: u32, + pad: u32, +} +struct GScaleValues { + Scale: f32, + Rel: f32, + pad: f32, + pad1: f32, +} +struct PathParams { + Type: PathTypes, + Index: u32, + pad: i32, + pad1: i32, + Indexes: PathIndexes, + Com: SynComParams, + PathScale: PathScaleParams, + SWts: SWtParams, + Learn: LearnSynParams, + GScale: GScaleValues, + RLPred: RLPredPathParams, + Matrix: MatrixPathParams, + BLA: BLAPathParams, + Hip: HipPathParams, +} + +///////////// import: "pathtypes.go" +alias PathTypes = i32; //enums:enum +const ForwardPath: PathTypes = 0; +const BackPath: PathTypes = 1; +const LateralPath: PathTypes = 2; +const InhibPath: PathTypes = 3; +const CTCtxtPath: PathTypes = 4; +const RWPath: PathTypes = 5; +const TDPredPath: PathTypes = 6; +const BLAPath: PathTypes = 7; +const HipPath: PathTypes = 8; +const VSPatchPath: PathTypes = 9; +const VSMatrixPath: PathTypes = 10; +const DSMatrixPath: PathTypes = 11; + +///////////// import: "pcore-layer.go" +struct MatrixParams { + GateThr: f32, + IsVS: i32, + OtherMatrixIndex: i32, + ThalLay1Index: i32, + ThalLay2Index: i32, + ThalLay3Index: i32, + ThalLay4Index: i32, + ThalLay5Index: i32, + ThalLay6Index: i32, + pad: i32, + pad1: i32, + pad2: i32, +} +alias GPLayerTypes = i32; //enums:enum +const GPePr: GPLayerTypes = 0; +const GPeAk: GPLayerTypes = 1; +const GPi: GPLayerTypes = 2; +struct GPParams { + GPType: GPLayerTypes, + pad: u32, + pad1: u32, + pad2: u32, +} + +///////////// import: "pcore-path.go" +struct MatrixPathParams { + Credit: f32, + BasePF: f32, + Delta: f32, + VSRewLearn: i32, +} + +///////////// import: "pool.go" +alias PoolIntVars = i32; //enums:enum +const PoolNeurSt: PoolIntVars = 0; +const PoolNeurEd: PoolIntVars = 1; +const PoolLayerIdx: PoolIntVars = 2; +const PoolIsLayer: PoolIntVars = 3; +const Clamped: PoolIntVars = 4; +const PoolGated: PoolIntVars = 5; +const FFsRawInt: PoolIntVars = 6; +const FBsRawInt: PoolIntVars = 7; +const GeExtRawInt: PoolIntVars = 8; +const PoolIntAvgMaxStart: PoolIntVars = 9; +alias AvgMax = i32; //enums:enum +const Avg: AvgMax = 0; +const Max: AvgMax = 1; +alias AvgMaxPhases = i32; //enums:enum -trim-prefix AM +const AMCycle: AvgMaxPhases = 0; +const AMMinus: AvgMaxPhases = 1; +const AMPlus: AvgMaxPhases = 2; +const AMPrev: AvgMaxPhases = 3; +alias AvgMaxVars = i32; //enums:enum -trim-prefix AM +const AMCaSpkP: AvgMaxVars = 0; +const AMCaSpkD: AvgMaxVars = 1; +const AMSpkMax: AvgMaxVars = 2; +const AMAct: AvgMaxVars = 3; +const AMGeInt: AvgMaxVars = 4; +const AMGiInt: AvgMaxVars = 5; +const AMAvgDif: AvgMaxVars = 6; +const poolFloatAvgMaxStart = InhibVarsN; +const PoolVarsN = poolFloatAvgMaxStart + InhibVars(i32(AvgMaxVarsN)*i32(AvgMaxN)*i32(AvgMaxPhasesN)); +const PoolIntVarsTot = PoolIntAvgMaxStart + PoolIntVars(i32(AvgMaxVarsN)*i32(AvgMaxN)); +const avgMaxToNeuron = array(CaSpkP, CaSpkD, SpkMax, Act, GeInt, GiInt); + +///////////// import: "rand.go" +alias RandFunIndex = u32; +const RandFunActPGe: RandFunIndex = 0; +const RandFunActPGi: RandFunIndex = 1; +const RandFunActSMaintP: RandFunIndex = 2; +const RandFunIndexN: RandFunIndex = 3; + +///////////// import: "rl-layer.go" +struct RWPredParams { + PredRange: F32, +} +struct RWDaParams { + TonicGe: f32, + RWPredLayIndex: i32, + pad: u32, + pad1: u32, +} +struct TDIntegParams { + Discount: f32, + PredGain: f32, + TDPredLayIndex: i32, + pad: u32, +} +struct TDDaParams { + TonicGe: f32, + TDIntegLayIndex: i32, + pad: u32, + pad1: u32, +} + +///////////// import: "rl-path.go" +struct RLPredPathParams { + OppSignLRate: f32, + DaTol: f32, + pad: f32, + pad1: f32, +} + +///////////// import: "rubicon-layer.go" +struct LDTParams { + SrcThr: f32, + Rew: i32, + MaintInhib: f32, + SrcLay1Index: i32, + SrcLay2Index: i32, + SrcLay3Index: i32, + SrcLay4Index: i32, + pad: f32, +} +struct VTAParams { + CeMGain: f32, + LHbGain: f32, + AChThr: f32, + pad: f32, +} + +///////////// import: "rubicon-path.go" +struct BLAPathParams { + NegDeltaLRate: f32, + AChThr: f32, + USTrace: f32, + pad: f32, +} + +///////////// import: "rubicon.go" + +///////////// import: "stats.go" + +///////////// import: "synapse.go" +alias SynapseVars = i32; //enums:enum +const Wt: SynapseVars = 0; +const LWt: SynapseVars = 1; +const SWt: SynapseVars = 2; +const DWt: SynapseVars = 3; +const DSWt: SynapseVars = 4; +alias SynapseTraceVars = i32; //enums:enum +const Tr: SynapseTraceVars = 0; +const DTr: SynapseTraceVars = 1; +const DiDWt: SynapseTraceVars = 2; +alias SynapseIndexVars = i32; //enums:enum +const SynRecvIndex: SynapseIndexVars = 0; +const SynSendIndex: SynapseIndexVars = 1; +const SynPathIndex: SynapseIndexVars = 2; + +///////////// import: "slrand.wgsl" +fn Philox2x32round(counter: su64, key: u32) -> su64 { + let mul = Uint32Mul64(u32(0xD256D193), counter.x); + var ctr: su64; + ctr.x = mul.y ^ key ^ counter.y; + ctr.y = mul.x; + return ctr; +} +fn Philox2x32bumpkey(key: u32) -> u32 { + return key + u32(0x9E3779B9); +} +fn Philox2x32(counter: su64, key: u32) -> vec2 { + var ctr = Philox2x32round(counter, key); // 1 + var ky = Philox2x32bumpkey(key); + ctr = Philox2x32round(ctr, ky); // 2 + ky = Philox2x32bumpkey(ky); + ctr = Philox2x32round(ctr, ky); // 3 + ky = Philox2x32bumpkey(ky); + ctr = Philox2x32round(ctr, ky); // 4 + ky = Philox2x32bumpkey(ky); + ctr = Philox2x32round(ctr, ky); // 5 + ky = Philox2x32bumpkey(ky); + ctr = Philox2x32round(ctr, ky); // 6 + ky = Philox2x32bumpkey(ky); + ctr = Philox2x32round(ctr, ky); // 7 + ky = Philox2x32bumpkey(ky); + ctr = Philox2x32round(ctr, ky); // 8 + ky = Philox2x32bumpkey(ky); + ctr = Philox2x32round(ctr, ky); // 9 + ky = Philox2x32bumpkey(ky); + return Philox2x32round(ctr, ky); // 10 +} +fn RandUint32Vec2(counter: su64, funcIndex: u32, key: u32) -> vec2 { + return Philox2x32(Uint64Add32(counter, funcIndex), key); +} +fn RandUint32(counter: su64, funcIndex: u32, key: u32) -> u32 { + return Philox2x32(Uint64Add32(counter, funcIndex), key).x; +} +fn RandFloat32Vec2(counter: su64, funcIndex: u32, key: u32) -> vec2 { + return Uint32ToFloat32Vec2(RandUint32Vec2(counter, funcIndex, key)); +} +fn RandFloat32(counter: su64, funcIndex: u32, key: u32) -> f32 { + return Uint32ToFloat32(RandUint32(counter, funcIndex, key)); +} +fn RandFloat32Range11Vec2(counter: su64, funcIndex: u32, key: u32) -> vec2 { + return Uint32ToFloat32Vec2(RandUint32Vec2(counter, funcIndex, key)); +} +fn RandFloat32Range11(counter: su64, funcIndex: u32, key: u32) -> f32 { + return Uint32ToFloat32Range11(RandUint32(counter, funcIndex, key)); +} +fn RandBoolP(counter: su64, funcIndex: u32, key: u32, p: f32) -> bool { + return (RandFloat32(counter, funcIndex, key) < p); +} +fn sincospi(x: f32) -> vec2 { + let PIf = 3.1415926535897932; + var r: vec2; + r.x = cos(PIf*x); + r.y = sin(PIf*x); + return r; +} +fn RandFloat32NormVec2(counter: su64, funcIndex: u32, key: u32) -> vec2 { + let ur = RandUint32Vec2(counter, funcIndex, key); + var f = sincospi(Uint32ToFloat32Range11(ur.x)); + let r = sqrt(-2.0 * log(Uint32ToFloat32(ur.y))); // guaranteed to avoid 0. + return f * r; +} +fn RandFloat32Norm(counter: su64, funcIndex: u32, key: u32) -> f32 { + return RandFloat32Vec2(counter, funcIndex, key).x; +} +fn RandUint32N(counter: su64, funcIndex: u32, key: u32, n: u32) -> u32 { + let v = RandFloat32(counter, funcIndex, key); + return u32(v * f32(n)); +} +struct RandCounter { + Counter: su64, + HiSeed: u32, + pad: u32, +} +fn RandCounter_Reset(ct: ptr) { + (*ct).Counter.x = u32(0); + (*ct).Counter.y = (*ct).HiSeed; +} +fn RandCounter_Seed(ct: ptr, seed: u32) { + (*ct).HiSeed = seed; + RandCounter_Reset(ct); +} +fn RandCounter_Add(ct: ptr, inc: u32) { + (*ct).Counter = Uint64Add32((*ct).Counter, inc); +} + +///////////// import: "sltype.wgsl" +alias su64 = vec2; +fn Uint32Mul64(a: u32, b: u32) -> su64 { + let LOMASK = (((u32(1))<<16)-1); + var r: su64; + r.x = a * b; /* full low multiply */ + let ahi = a >> 16; + let alo = a & LOMASK; + let bhi = b >> 16; + let blo = b & LOMASK; + let ahbl = ahi * blo; + let albh = alo * bhi; + let ahbl_albh = ((ahbl&LOMASK) + (albh&LOMASK)); + var hit = ahi*bhi + (ahbl>>16) + (albh>>16); + hit += ahbl_albh >> 16; /* carry from the sum of lo(ahbl) + lo(albh) ) */ + /* carry from the sum with alo*blo */ + if ((r.x >> u32(16)) < (ahbl_albh&LOMASK)) { + hit += u32(1); + } + r.y = hit; + return r; +} +/* +fn Uint32Mul64(a: u32, b: u32) -> su64 { + return su64(a) * su64(b); +} +*/ +fn Uint64Add32(a: su64, b: u32) -> su64 { + if (b == 0) { + return a; + } + var s = a; + if (s.x > u32(0xffffffff) - b) { + s.y++; + s.x = (b - 1) - (u32(0xffffffff) - s.x); + } else { + s.x += b; + } + return s; +} +fn Uint64Incr(a: su64) -> su64 { + var s = a; + if(s.x == 0xffffffff) { + s.y++; + s.x = u32(0); + } else { + s.x++; + } + return s; +} +fn Uint32ToFloat32(val: u32) -> f32 { + let factor = f32(1.0) / (f32(u32(0xffffffff)) + f32(1.0)); + let halffactor = f32(0.5) * factor; + var f = f32(val) * factor + halffactor; + if (f == 1.0) { // exclude 1 + return bitcast(0x3F7FFFFF); + } + return f; +} +fn Uint32ToFloat32Vec2(val: vec2) -> vec2 { + var r: vec2; + r.x = Uint32ToFloat32(val.x); + r.y = Uint32ToFloat32(val.y); + return r; +} +fn Uint32ToFloat32Range11(val: u32) -> f32 { + let factor = f32(1.0) / (f32(i32(0x7fffffff)) + f32(1.0)); + let halffactor = f32(0.5) * factor; + return (f32(val) * factor + halffactor); +} +fn Uint32ToFloat32Range11Vec2(val: vec2) -> vec2 { + var r: vec2; + r.x = Uint32ToFloat32Range11(val.x); + r.y = Uint32ToFloat32Range11(val.y); + return r; +} \ No newline at end of file diff --git a/axon/shaders/WtFromDWtSyn.wgsl b/axon/shaders/WtFromDWtSyn.wgsl index 2eee2baf..1034ba05 100644 --- a/axon/shaders/WtFromDWtSyn.wgsl +++ b/axon/shaders/WtFromDWtSyn.wgsl @@ -704,6 +704,8 @@ const LayerRT: LayerVars = 8; const LayerRewPredPos: LayerVars = 9; const LayerRewPredNeg: LayerVars = 10; +///////////// import: "learn-layer.go" + ///////////// import: "learn-net.go" fn WtFromDWtSyn(syni: u32) { //gosl:kernel var ctx = Ctx[0];