From ec6abacdb887f68f9cbd74de9a0a24bb5a1cbb76 Mon Sep 17 00:00:00 2001 From: Alex Black Date: Tue, 24 Mar 2020 20:33:43 +1100 Subject: [PATCH 1/6] Fix limits on flaky test to avoid spurious failure (#344) Signed-off-by: Alex Black --- .../org/nd4j/autodiff/opvalidation/RandomOpValidation.java | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/autodiff/opvalidation/RandomOpValidation.java b/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/autodiff/opvalidation/RandomOpValidation.java index 4f228717a..4585b4a15 100644 --- a/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/autodiff/opvalidation/RandomOpValidation.java +++ b/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/autodiff/opvalidation/RandomOpValidation.java @@ -76,7 +76,7 @@ public class RandomOpValidation extends BaseOpValidation { double min = in.minNumber().doubleValue(); double max = in.maxNumber().doubleValue(); double mean = in.meanNumber().doubleValue(); - if (min >= 1 && max <= 2 && (in.length() == 1 || Math.abs(mean - 1.5) < 0.1)) + if (min >= 1 && max <= 2 && (in.length() == 1 || Math.abs(mean - 1.5) < 0.2)) return null; return "Failed: min = " + min + ", max = " + max + ", mean = " + mean; }; @@ -87,7 +87,7 @@ public class RandomOpValidation extends BaseOpValidation { checkFn = in -> { double mean = in.meanNumber().doubleValue(); double stdev = in.std(true).getDouble(0); - if (in.length() == 1 || (Math.abs(mean - 1) < 0.1 && Math.abs(stdev - 1) < 0.1)) + if (in.length() == 1 || (Math.abs(mean - 1) < 0.2 && Math.abs(stdev - 1) < 0.2)) return null; return "Failed: mean = " + mean + ", stdev = " + stdev; }; From b1bc7df160128d3def31e4884c7a32a12868caa1 Mon Sep 17 00:00:00 2001 From: Fariz Rahman Date: Tue, 24 Mar 2020 13:37:27 +0400 Subject: [PATCH 2/6] tf.keras model import (#258) * tf op initial * .. * protobuf parsing working * model build working * test passing * headers * conffix * service loader + tests * revert cuda version * msg * override * refacc * pom * rem bad import * dtype fix + const cast caaching * rem unnecessary fields * rem println * rem dep * refacc * rem redundant arg * Ignore TFOpLayer in DTypeTests Signed-off-by: Alex Black Co-authored-by: Alex Black --- .../deeplearning4j/nn/dtypes/DTypeTests.java | 3 +- .../deeplearning4j-modelimport/pom.xml | 8 + .../config/Keras2LayerConfiguration.java | 2 + .../keras/layers/KerasTFOpLayer.java | 74 ++++++++ .../modelimport/keras/layers/TFOpLayer.java | 106 +++++++++++ .../keras/layers/TFOpLayerImpl.java | 169 ++++++++++++++++++ .../keras/utils/KerasLayerUtils.java | 17 ++ .../nn/modelimport/keras/TFKerasTests.java | 50 ++++++ deeplearning4j/deeplearning4j-nn/pom.xml | 6 +- .../nn/layers/AbstractLayer.java | 1 + .../java/org/nd4j/TFGraphRunnerService.java | 37 ++++ .../conversion/graphrunner/GraphRunner.java | 14 +- .../GraphRunnerServiceProvider.java | 52 ++++++ .../services/org.nd4j.TFGraphRunnerService | 17 ++ 14 files changed, 550 insertions(+), 6 deletions(-) create mode 100644 deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/KerasTFOpLayer.java create mode 100644 deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/TFOpLayer.java create mode 100644 deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/TFOpLayerImpl.java create mode 100644 deeplearning4j/deeplearning4j-modelimport/src/test/java/org/deeplearning4j/nn/modelimport/keras/TFKerasTests.java create mode 100644 nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/TFGraphRunnerService.java create mode 100644 nd4j/nd4j-tensorflow/src/main/java/org/nd4j/tensorflow/conversion/graphrunner/GraphRunnerServiceProvider.java create mode 100644 nd4j/nd4j-tensorflow/src/main/resources/META-INF/services/org.nd4j.TFGraphRunnerService diff --git a/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/nn/dtypes/DTypeTests.java b/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/nn/dtypes/DTypeTests.java index d9da12b62..6831af10b 100644 --- a/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/nn/dtypes/DTypeTests.java +++ b/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/nn/dtypes/DTypeTests.java @@ -17,6 +17,7 @@ package org.deeplearning4j.nn.dtypes; import org.deeplearning4j.nn.conf.layers.recurrent.TimeDistributed; +import org.deeplearning4j.nn.modelimport.keras.layers.TFOpLayer; import org.nd4j.shade.guava.collect.ImmutableSet; import org.nd4j.shade.guava.reflect.ClassPath; import lombok.extern.slf4j.Slf4j; @@ -128,7 +129,7 @@ public class DTypeTests extends BaseDL4JTest { throw new RuntimeException(e); } - if (Modifier.isAbstract(clazz.getModifiers()) || clazz.isInterface()) { + if (Modifier.isAbstract(clazz.getModifiers()) || clazz.isInterface() || TFOpLayer.class == clazz) { //Skip TFOpLayer here - dtype depends on imported model dtype continue; } diff --git a/deeplearning4j/deeplearning4j-modelimport/pom.xml b/deeplearning4j/deeplearning4j-modelimport/pom.xml index 566bf6012..6d71c394e 100644 --- a/deeplearning4j/deeplearning4j-modelimport/pom.xml +++ b/deeplearning4j/deeplearning4j-modelimport/pom.xml @@ -105,6 +105,14 @@ ${project.version} test + + + org.nd4j + nd4j-tensorflow + ${nd4j.version} + test + + diff --git a/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/config/Keras2LayerConfiguration.java b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/config/Keras2LayerConfiguration.java index 430b7407a..9b91d10cc 100644 --- a/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/config/Keras2LayerConfiguration.java +++ b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/config/Keras2LayerConfiguration.java @@ -103,4 +103,6 @@ public class Keras2LayerConfiguration extends KerasLayerConfiguration { /* Keras weight initializers. */ private final String LAYER_FIELD_INIT = "kernel_initializer"; + + private final String TENSORFLOW_OP_LAYER = "TensorFlowOpLayer"; } \ No newline at end of file diff --git a/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/KerasTFOpLayer.java b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/KerasTFOpLayer.java new file mode 100644 index 000000000..2dd95338a --- /dev/null +++ b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/KerasTFOpLayer.java @@ -0,0 +1,74 @@ +/******************************************************************************* + * Copyright (c) 2020 Konduit K.K. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +package org.deeplearning4j.nn.modelimport.keras.layers; + +import org.deeplearning4j.nn.conf.inputs.InputType; +import org.deeplearning4j.nn.modelimport.keras.KerasLayer; +import org.deeplearning4j.nn.modelimport.keras.exceptions.InvalidKerasConfigurationException; +import org.deeplearning4j.nn.modelimport.keras.exceptions.UnsupportedKerasConfigurationException; + +import java.util.Map; + + +public class KerasTFOpLayer extends KerasLayer { + + public KerasTFOpLayer(Integer kerasVersion) throws UnsupportedKerasConfigurationException { + super(kerasVersion); + if (kerasVersion != 2){ + throw new UnsupportedKerasConfigurationException("KerasTFOpLayer expects Keras version 2"); + } + } + + /** + * Constructor from parsed Keras layer configuration dictionary. + * + * @param layerConfig dictionary containing Keras layer configuration + * @throws InvalidKerasConfigurationException Invalid Keras config + * @throws UnsupportedKerasConfigurationException Unsupported Keras config + */ + public KerasTFOpLayer(Map layerConfig) + throws InvalidKerasConfigurationException, UnsupportedKerasConfigurationException { + this(layerConfig, true); + } + + /** + * Constructor from parsed Keras layer configuration dictionary. + * + * @param layerConfig dictionary containing Keras layer configuration + * @param enforceTrainingConfig whether to enforce training-related configuration options + * @throws InvalidKerasConfigurationException Invalid Keras config + * @throws UnsupportedKerasConfigurationException Unsupported Keras config + */ + public KerasTFOpLayer(Map layerConfig, boolean enforceTrainingConfig) throws UnsupportedKerasConfigurationException, InvalidKerasConfigurationException{ + super(layerConfig, enforceTrainingConfig); + this.layer = new TFOpLayer((Map)((Map)layerConfig.get("config")).get("node_def"), (Map)((Map)layerConfig.get("config")).get("constants")); + } + + /** + * Get layer output type. + * + * @param inputType Array of InputTypes + * @return output type as InputType + * @throws InvalidKerasConfigurationException Invalid Keras configuration + */ + public InputType getOutputType(InputType... inputType){ + return this.layer.getOutputType(0, inputType[0]); + } + + + +} diff --git a/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/TFOpLayer.java b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/TFOpLayer.java new file mode 100644 index 000000000..ecf64e8c0 --- /dev/null +++ b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/TFOpLayer.java @@ -0,0 +1,106 @@ +/******************************************************************************* + * Copyright (c) 2020 Konduit K.K. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +package org.deeplearning4j.nn.modelimport.keras.layers; + +import org.deeplearning4j.nn.api.ParamInitializer; +import org.deeplearning4j.nn.conf.GradientNormalization; +import org.deeplearning4j.nn.conf.InputPreProcessor; +import org.deeplearning4j.nn.conf.NeuralNetConfiguration; +import org.deeplearning4j.nn.conf.inputs.InputType; +import org.deeplearning4j.nn.conf.layers.Layer; +import org.deeplearning4j.nn.conf.memory.LayerMemoryReport; +import org.deeplearning4j.nn.modelimport.keras.layers.TFOpLayerImpl; +import org.deeplearning4j.nn.params.EmptyParamInitializer; +import org.deeplearning4j.optimize.api.TrainingListener; +import org.nd4j.linalg.api.buffer.DataType; +import org.nd4j.linalg.api.ndarray.INDArray; +import org.nd4j.linalg.factory.Nd4j; +import org.nd4j.linalg.learning.regularization.Regularization; + +import java.util.Collection; +import java.util.List; +import java.util.Map; + + +public class TFOpLayer extends Layer { + + private Map nodeDef; + private Map constants; + public TFOpLayer(Map nodeDef, Map constants){ + super(); + this.nodeDef = nodeDef; + this.constants = constants; + } + + @Override + public ParamInitializer initializer() { + return EmptyParamInitializer.getInstance(); + } + @Override + public InputPreProcessor getPreProcessorForInputType(InputType inputType) { + return null; + } + + @Override + public boolean isPretrainParam(String param){ + return false; + } + + @Override + public InputType getOutputType(int idx, InputType inputType){ + long[] shape = inputType.getShape(true); + TFOpLayerImpl tempLayer = new TFOpLayerImpl(nodeDef, constants, null, null); + long[] outputShape = tempLayer.getOutputShape(shape); + return InputType.inferInputType(Nd4j.create(outputShape)); + + } + + @Override + public void setNIn(InputType inputType, boolean override){} + + + @Override + public GradientNormalization getGradientNormalization(){return null;} + + + @Override + public org.deeplearning4j.nn.api.Layer instantiate(NeuralNetConfiguration conf, + Collection trainingListeners, int layerIndex, INDArray layerParamsView, + boolean initializeParams, DataType networkDataType) { + + TFOpLayerImpl tfOpLayerImpl = new TFOpLayerImpl(nodeDef, constants, conf, networkDataType); + tfOpLayerImpl.setListeners(trainingListeners); + tfOpLayerImpl.setIndex(layerIndex); + return tfOpLayerImpl; + } + + @Override + public double getGradientNormalizationThreshold(){return 0.;} + + @Override + public List getRegularizationByParam(String paramName){return null;} + + @Override + public LayerMemoryReport getMemoryReport(InputType inputType) { + return new LayerMemoryReport(); //TODO + } + + + + + +} diff --git a/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/TFOpLayerImpl.java b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/TFOpLayerImpl.java new file mode 100644 index 000000000..d7b0b3b56 --- /dev/null +++ b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/TFOpLayerImpl.java @@ -0,0 +1,169 @@ +/******************************************************************************* + * Copyright (c) 2020 Konduit K.K. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +package org.deeplearning4j.nn.modelimport.keras.layers; + +import lombok.Data; +import lombok.extern.slf4j.Slf4j; +import org.apache.commons.lang3.ArrayUtils; +import org.deeplearning4j.nn.conf.NeuralNetConfiguration; +import org.deeplearning4j.nn.gradient.Gradient; +import org.deeplearning4j.nn.layers.AbstractLayer; +import org.deeplearning4j.nn.workspace.LayerWorkspaceMgr; +import org.nd4j.TFGraphRunnerService; +import org.nd4j.linalg.api.buffer.DataType; +import org.nd4j.linalg.api.ndarray.INDArray; +import org.nd4j.linalg.factory.Nd4j; +import org.nd4j.linalg.primitives.Pair; +import org.tensorflow.framework.AttrValue; +import org.tensorflow.framework.GraphDef; +import org.tensorflow.framework.NodeDef; +import com.google.gson.Gson; +import org.nd4j.shade.protobuf.Message; +import org.nd4j.shade.protobuf.TextFormat; + +import java.util.*; +import java.util.List; + + +@Slf4j +@Data +public class TFOpLayerImpl extends AbstractLayer { + + + private Map nodeDef; + private Map constants; + private List inputNames; + TFGraphRunnerService graphRunnerService; + + public TFOpLayerImpl(Map nodeDef, Map constants, NeuralNetConfiguration conf, DataType dtype){ + super(conf, dtype); + this.nodeDef = nodeDef; + this.constants = constants; + setGraphRunner(); + } + + @Override + public Pair backpropGradient(INDArray epsilon, LayerWorkspaceMgr workspaceMgr){ + throw new RuntimeException("Backprop through TFOpLayerImpl is not supported yet." + + " TFOpLayerImpl is created when importing TensorFlow 2.0 Keras models " + + "(tf.keras) into DL4J, that contains TensorFlow operations not just Keras layers."); + } + + /** + * Converts a Map representation of Nodedef to a singleton TF Graph and instantiates a GraphRunner. + */ + private void setGraphRunner() { + try{ + String json = new Gson().toJson(nodeDef); + NodeDef.Builder builder = NodeDef.newBuilder(); + org.nd4j.shade.protobuf.util.JsonFormat.parser().merge(json, builder); + NodeDef nodeDef = builder.build(); + List allInputNames = new ArrayList<>(); // including constants + Map inputDataTypes = new HashMap<>(); + Map constArrays = new HashMap(); + this.inputNames = new ArrayList<>(); + List outputNames = Arrays.asList(nodeDef.getName()); + Map attrMap = nodeDef.getAttrMap(); + for (int i = 0; i < nodeDef.getInputCount(); i++){ + String inputName = nodeDef.getInput(i); + String[] split = inputName.split("/"); + String attrKey; + if (split.length == 1){ + attrKey = "T"; + } + else{ + attrKey = "T" + split[split.length - 1]; + } + allInputNames.add(nodeDef.getInput(i)); + inputDataTypes.put(nodeDef.getInput(i), attrMap.get(attrKey).getType().toString()); + if (constants.containsKey(String.valueOf(i))){ + constArrays.put(nodeDef.getInput(i), Nd4j.create((List)constants.get(String.valueOf(i)))); + } + else{ + this.inputNames.add(nodeDef.getInput(i)); + } + } + String graph = "node{\n" + nodeDef.toString() + "\n}\nversions {\n producer: 22\n}"; + for (int i = 0; i < allInputNames.size(); i++){ + String inpName = allInputNames.get(i); + String dtype = inputDataTypes.get(inpName); + graph = "node{\nname: \"" + inpName + "\"\nop: \"Placeholder\"\nattr{\nkey: \"dtype\"\n value {\n type: " + dtype + "}\n}\n}\n" + graph; + } + log.info(graph); + GraphDef.Builder graphDefBuilder = GraphDef.newBuilder(); + TextFormat.getParser().merge(graph, graphDefBuilder); + GraphDef graphDef = graphDefBuilder.build(); + org.nd4j.shade.protobuf.ByteString serialized = graphDef.toByteString(); + byte[] graphBytes = serialized.toByteArray(); + + ServiceLoader sl = ServiceLoader.load(TFGraphRunnerService.class); + Iterator iter = sl.iterator(); + if (!iter.hasNext()){ + throw new RuntimeException("The model contains a Tensorflow Op, which requires the nd4j-tensorflow dependency to execute."); + } + + this.graphRunnerService = iter.next().init(allInputNames, outputNames, graphBytes, constArrays, inputDataTypes); + } + catch (Exception e){ + throw new RuntimeException("Error parsing protobuf", e); + } + + } + + private INDArray runGraph(INDArray input){ + if (input.rank() == 3){ + // TODO make this a preprocessor + input = input.permute(0, 2, 1); + } + Map inputMap = new HashMap<>(); + inputMap.put(inputNames.get(0), input); + INDArray out = graphRunnerService.run(inputMap).values().toArray(new INDArray[0])[0]; + if (out.rank() == 3){ + out = out.permute(0, 2, 1); // TODO post-processing? + } + + return out; + } + + public long[] getOutputShape(long[] inputShape){ + long[] shape = ArrayUtils.clone(inputShape); + for(int i = 0; i < shape.length; i++){ + if (shape[i] < 0){ + shape[i] = 1; + } + } + INDArray dummyArr = Nd4j.zeros(shape); + return runGraph(dummyArr).shape(); + } + + @Override + public INDArray activate(boolean training, LayerWorkspaceMgr workspaceMgr){ + return runGraph(input); + } + + + @Override + public boolean isPretrainLayer(){ + return false; + } + + @Override + public void clearNoiseWeightParams(){ + + } + +} diff --git a/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/utils/KerasLayerUtils.java b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/utils/KerasLayerUtils.java index 1428b6322..3f69cb7d4 100644 --- a/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/utils/KerasLayerUtils.java +++ b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/utils/KerasLayerUtils.java @@ -21,10 +21,12 @@ import org.deeplearning4j.nn.conf.graph.ElementWiseVertex; import org.deeplearning4j.nn.conf.layers.Layer; import org.deeplearning4j.nn.conf.layers.samediff.SameDiffLambdaLayer; import org.deeplearning4j.nn.modelimport.keras.KerasLayer; +import org.deeplearning4j.nn.modelimport.keras.config.Keras2LayerConfiguration; import org.deeplearning4j.nn.modelimport.keras.config.KerasLayerConfiguration; import org.deeplearning4j.nn.modelimport.keras.exceptions.InvalidKerasConfigurationException; import org.deeplearning4j.nn.modelimport.keras.exceptions.UnsupportedKerasConfigurationException; import org.deeplearning4j.nn.modelimport.keras.layers.KerasInput; +import org.deeplearning4j.nn.modelimport.keras.layers.KerasTFOpLayer; import org.deeplearning4j.nn.modelimport.keras.layers.advanced.activations.*; import org.deeplearning4j.nn.modelimport.keras.layers.convolutional.*; import org.deeplearning4j.nn.modelimport.keras.layers.core.*; @@ -317,6 +319,11 @@ public class KerasLayerUtils { layer = new KerasELU(layerConfig, enforceTrainingConfig); } else if(layerClassName.equals(conf.getLAYER_CLASS_NAME_SOFTMAX())){ layer = new KerasSoftmax(layerConfig, enforceTrainingConfig); + } else if (conf instanceof Keras2LayerConfiguration){ + Keras2LayerConfiguration k2conf = (Keras2LayerConfiguration)conf; + if (layerClassName.equals(k2conf.getTENSORFLOW_OP_LAYER())){ + layer = new KerasTFOpLayer(layerConfig, enforceTrainingConfig); + } } if (layer == null){ Class customConfig = customLayers.get(layerClassName); @@ -402,6 +409,16 @@ public class KerasLayerUtils { public static String getLayerNameFromConfig(Map layerConfig, KerasLayerConfiguration conf) throws InvalidKerasConfigurationException { + if(conf instanceof Keras2LayerConfiguration){ + Keras2LayerConfiguration k2conf = (Keras2LayerConfiguration)conf; + if (getClassNameFromConfig(layerConfig, conf).equals(((Keras2LayerConfiguration) conf).getTENSORFLOW_OP_LAYER())){ + if (!layerConfig.containsKey(conf.getLAYER_FIELD_NAME())) + throw new InvalidKerasConfigurationException("Field " + conf.getLAYER_FIELD_NAME() + + " missing from layer config"); + return (String) layerConfig.get(conf.getLAYER_FIELD_NAME()); + } + } + Map innerConfig = KerasLayerUtils.getInnerLayerConfigFromConfig(layerConfig, conf); if (!innerConfig.containsKey(conf.getLAYER_FIELD_NAME())) throw new InvalidKerasConfigurationException("Field " + conf.getLAYER_FIELD_NAME() diff --git a/deeplearning4j/deeplearning4j-modelimport/src/test/java/org/deeplearning4j/nn/modelimport/keras/TFKerasTests.java b/deeplearning4j/deeplearning4j-modelimport/src/test/java/org/deeplearning4j/nn/modelimport/keras/TFKerasTests.java new file mode 100644 index 000000000..cb74b1ed1 --- /dev/null +++ b/deeplearning4j/deeplearning4j-modelimport/src/test/java/org/deeplearning4j/nn/modelimport/keras/TFKerasTests.java @@ -0,0 +1,50 @@ +/******************************************************************************* + * Copyright (c) 2020 Konduit K.K. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +package org.deeplearning4j.nn.modelimport.keras; + +import org.deeplearning4j.BaseDL4JTest; +import org.deeplearning4j.nn.graph.ComputationGraph; +import org.junit.Assert; +import org.junit.Test; +import org.nd4j.linalg.api.ndarray.INDArray; +import org.nd4j.linalg.factory.Nd4j; +import org.nd4j.resources.Resources; + +import java.io.File; +import java.util.Arrays; + +public class TFKerasTests extends BaseDL4JTest{ + + @Test + public void testModelWithTFOp1() throws Exception{ + File f = Resources.asFile("modelimport/keras/tfkeras/reshape.h5"); + ComputationGraph graph = KerasModelImport.importKerasModelAndWeights(f.getAbsolutePath()); + INDArray out = graph.outputSingle(Nd4j.zeros(12, 2, 3)); + Assert.assertArrayEquals(new long[]{12, 3}, out.shape()); + } + + @Test + public void testModelWithTFOp2() throws Exception{ + File f = Resources.asFile("modelimport/keras/tfkeras/permute.h5"); + ComputationGraph graph = KerasModelImport.importKerasModelAndWeights(f.getAbsolutePath()); + INDArray out = graph.outputSingle(Nd4j.zeros(12, 2, 3)); + // dl4j's feedforward doesn't support 3D output, so batch and time axes gets squashed + long[] expectedShape = new long[]{12 * 2, 5}; + Assert.assertArrayEquals(expectedShape, out.shape()); + } + +} diff --git a/deeplearning4j/deeplearning4j-nn/pom.xml b/deeplearning4j/deeplearning4j-nn/pom.xml index e92372fc8..77acb2dc7 100644 --- a/deeplearning4j/deeplearning4j-nn/pom.xml +++ b/deeplearning4j/deeplearning4j-nn/pom.xml @@ -77,7 +77,11 @@ nd4j-common ${nd4j.version} - + + com.google.code.gson + gson + ${gson.version} + org.nd4j diff --git a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/layers/AbstractLayer.java b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/layers/AbstractLayer.java index 750bca77d..ad8590b0b 100644 --- a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/layers/AbstractLayer.java +++ b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/layers/AbstractLayer.java @@ -62,6 +62,7 @@ public abstract class AbstractLayer inputNames, + List outputNames, + byte[] graphBytes, + Map constants, + Map inputDataTypes + ); + + Map run(Map inputs); +} diff --git a/nd4j/nd4j-tensorflow/src/main/java/org/nd4j/tensorflow/conversion/graphrunner/GraphRunner.java b/nd4j/nd4j-tensorflow/src/main/java/org/nd4j/tensorflow/conversion/graphrunner/GraphRunner.java index 9cb0a609b..49861e3fe 100644 --- a/nd4j/nd4j-tensorflow/src/main/java/org/nd4j/tensorflow/conversion/graphrunner/GraphRunner.java +++ b/nd4j/nd4j-tensorflow/src/main/java/org/nd4j/tensorflow/conversion/graphrunner/GraphRunner.java @@ -16,18 +16,16 @@ package org.nd4j.tensorflow.conversion.graphrunner; -import lombok.Builder; -import lombok.Singular; +import lombok.*; import org.apache.commons.io.FileUtils; import org.nd4j.base.Preconditions; +import org.nd4j.linalg.api.buffer.DataType; import org.nd4j.linalg.factory.Nd4j; import org.nd4j.linalg.io.ClassPathResource; import org.nd4j.linalg.primitives.Pair; import org.nd4j.shade.protobuf.ByteString; import org.nd4j.shade.protobuf.InvalidProtocolBufferException; import org.nd4j.shade.protobuf.util.JsonFormat; -import lombok.Getter; -import lombok.Setter; import lombok.extern.slf4j.Slf4j; import org.nd4j.tensorflow.conversion.TensorDataType; import org.apache.commons.io.IOUtils; @@ -56,6 +54,7 @@ import static org.bytedeco.tensorflow.global.tensorflow.*; * @author Adam Gibson */ @Slf4j +@NoArgsConstructor public class GraphRunner implements Closeable { private static boolean isTfWarmedUp = false; @@ -103,6 +102,9 @@ public class GraphRunner implements Closeable { * @param inputDataTypes the expected input data types * @param outputDataTypes the expected output data types */ + + + @Builder public GraphRunner(List inputNames, List outputNames, @@ -440,6 +442,7 @@ public class GraphRunner implements Closeable { * @return a map of the output names to the * ndarrays matching each output specified in the graph */ + public Map run(Map inputs) { if (!isTfWarmedUp && !isTfWarmingUp){ isTfWarmingUp = true; @@ -683,4 +686,7 @@ public class GraphRunner implements Closeable { return builder1.build(); } + + + } diff --git a/nd4j/nd4j-tensorflow/src/main/java/org/nd4j/tensorflow/conversion/graphrunner/GraphRunnerServiceProvider.java b/nd4j/nd4j-tensorflow/src/main/java/org/nd4j/tensorflow/conversion/graphrunner/GraphRunnerServiceProvider.java new file mode 100644 index 000000000..7459a40ea --- /dev/null +++ b/nd4j/nd4j-tensorflow/src/main/java/org/nd4j/tensorflow/conversion/graphrunner/GraphRunnerServiceProvider.java @@ -0,0 +1,52 @@ +package org.nd4j.tensorflow.conversion.graphrunner; + +import org.nd4j.TFGraphRunnerService; +import org.nd4j.linalg.api.buffer.DataType; +import org.nd4j.linalg.api.ndarray.INDArray; +import org.nd4j.tensorflow.conversion.TensorDataType; + +import java.util.HashMap; +import java.util.List; +import java.util.Map; + +public class GraphRunnerServiceProvider implements TFGraphRunnerService { + + private GraphRunner graphRunner; + Map inputs; + + @Override + public TFGraphRunnerService init( + List inputNames, + List outputNames, + byte[] graphBytes, + Map constants, + Map inputDataTypes){ + if (inputNames.size() != inputDataTypes.size()){ + throw new IllegalArgumentException("inputNames.size() != inputDataTypes.size()"); + } + Map convertedDataTypes = new HashMap<>(); + for (int i = 0; i < inputNames.size(); i++){ + convertedDataTypes.put(inputNames.get(i), TensorDataType.fromProtoValue(inputDataTypes.get(inputNames.get(i)))); + } + Map castConstants = new HashMap<>(); + for (Map.Entry e: constants.entrySet()) { + DataType requiredDtype = TensorDataType.toNd4jType(TensorDataType.fromProtoValue(inputDataTypes.get(e.getKey()))); + castConstants.put(e.getKey(), e.getValue().castTo(requiredDtype)); + } + this.inputs = castConstants; + graphRunner = GraphRunner.builder().inputNames(inputNames) + .outputNames(outputNames).graphBytes(graphBytes) + .inputDataTypes(convertedDataTypes).build(); + return this; + + } + + @Override + public Map run(Map inputs){ + if (graphRunner == null){ + throw new RuntimeException("GraphRunner not initialized."); + } + this.inputs.putAll(inputs); + return graphRunner.run(this.inputs); + } +} diff --git a/nd4j/nd4j-tensorflow/src/main/resources/META-INF/services/org.nd4j.TFGraphRunnerService b/nd4j/nd4j-tensorflow/src/main/resources/META-INF/services/org.nd4j.TFGraphRunnerService new file mode 100644 index 000000000..1b038ee6c --- /dev/null +++ b/nd4j/nd4j-tensorflow/src/main/resources/META-INF/services/org.nd4j.TFGraphRunnerService @@ -0,0 +1,17 @@ + ################################################################################ + # Copyright (c) 2020 Konduit K.K.. + # + # This program and the accompanying materials are made available under the + # terms of the Apache License, Version 2.0 which is available at + # https://www.apache.org/licenses/LICENSE-2.0. + # + # Unless required by applicable law or agreed to in writing, software + # distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + # WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + # License for the specific language governing permissions and limitations + # under the License. + # + # SPDX-License-Identifier: Apache-2.0 + ################################################################################ + +org.nd4j.tensorflow.conversion.graphrunner.GraphRunnerServiceProvider From e8cbf5255a43638ce827a25a0b7bd3486625253b Mon Sep 17 00:00:00 2001 From: Oleh Date: Wed, 25 Mar 2020 07:40:30 +0200 Subject: [PATCH 3/6] Backpropagation implementation of mergemax, mergeadd and mergeavg ops (#343) * libnd4j: first step of merge_max implementation Signed-off-by: Oleg * libnd4j fixed typos Signed-off-by: Oleg * libnd4j some corrections for mergeMaxBp Signed-off-by: Oleg * libnd4j some minor corrections Signed-off-by: Oleg * libnd4j test added for mergemax_bp Signed-off-by: Oleg * libnd4j fixed several problems tests added, check with gradCheck Signed-off-by: Oleg * libnd4j remove duplicated tests Signed-off-by: Oleg * libnd4j split implementation of transforms ops into separate file implementation Signed-off-by: Oleg * libnd4j code clean up, added mergeavg_bp and mergeadd_bp, need testing Signed-off-by: Oleg * libnd4j merge master, fixed typos and added tests Signed-off-by: Oleg * libnd4j some minor fixes Signed-off-by: Oleg * libnd4j added helper for mergeAddBp operation, this permits to skip nullify Signed-off-by: Oleg * libnd4j file renaming changes and cuda some corrections, need some additional corrections Signed-off-by: Oleg * libnd4j some additional corrections for merge ops Signed-off-by: Oleg * libnd4j more corrections per request for cuda more proper usage Signed-off-by: Oleg --- .../generic/transforms/merge_add.cpp | 42 +- .../generic/transforms/merge_avg.cpp | 40 +- .../generic/transforms/merge_max.cpp | 44 +- .../generic/transforms/merge_max_idx.cpp | 2 +- .../ops/declarable/headers/transforms.h | 3 + .../ops/declarable/helpers/cpu/clip.cpp | 274 ++++ .../ops/declarable/helpers/cpu/eye.cpp | 45 + .../helpers/cpu/gatherTransforms.cpp | 183 +++ .../helpers/cpu/invertPermutation.cpp | 51 + .../ops/declarable/helpers/cpu/merge.cpp | 277 ++++ .../ops/declarable/helpers/cpu/pad.cpp | 483 ++++++ .../declarable/helpers/cpu/randomShuffle.cpp | 126 ++ .../helpers/cpu/scatterUpdateAndSimple.cpp | 115 ++ .../ops/declarable/helpers/cpu/tile.cpp | 91 ++ .../ops/declarable/helpers/cpu/trace.cpp | 47 + .../ops/declarable/helpers/cpu/transforms.cpp | 1305 ----------------- .../ops/declarable/helpers/cpu/triu.cpp | 56 + .../ops/declarable/helpers/cuda/merge.cu | 415 +++++- .../ops/declarable/helpers/transforms.h | 11 +- .../layers_tests/DeclarableOpsTests13.cpp | 153 ++ 20 files changed, 2381 insertions(+), 1382 deletions(-) create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/clip.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/eye.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/gatherTransforms.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/invertPermutation.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/merge.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/pad.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/randomShuffle.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/scatterUpdateAndSimple.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/tile.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/trace.cpp delete mode 100644 libnd4j/include/ops/declarable/helpers/cpu/transforms.cpp create mode 100644 libnd4j/include/ops/declarable/helpers/cpu/triu.cpp diff --git a/libnd4j/include/ops/declarable/generic/transforms/merge_add.cpp b/libnd4j/include/ops/declarable/generic/transforms/merge_add.cpp index a68a4ce02..64858001a 100644 --- a/libnd4j/include/ops/declarable/generic/transforms/merge_add.cpp +++ b/libnd4j/include/ops/declarable/generic/transforms/merge_add.cpp @@ -33,7 +33,7 @@ OP_IMPL(mergeadd, -1, 1, false) { auto output = OUTPUT_VARIABLE(0); - std::vector inArrs(block.width()); + std::vector inArrs(block.width()); for(int i = 0; i < block.width(); ++i) inArrs[i] = INPUT_VARIABLE(i); @@ -42,7 +42,6 @@ OP_IMPL(mergeadd, -1, 1, false) { return Status::OK(); } - DECLARE_SYN(mergesum, mergeadd); DECLARE_SYN(add_n, mergeadd); DECLARE_SYN(addn, mergeadd); @@ -54,6 +53,45 @@ DECLARE_SYN(accumulate_n, mergeadd); ->setAllowedInputTypes(sd::DataType::ANY) ->setAllowedOutputTypes(sd::DataType::ANY); } + + + CUSTOM_OP_IMPL(mergeadd_bp, 2, 1, false, 0, 0) { + + auto inSize = block.width() - 1; + + REQUIRE_OK(this->validateInputDimensionsMatch(block)); + + std::vector outArrs(inSize); + + const auto gradient = INPUT_VARIABLE(inSize); + + for (int i = 0; i < inSize; ++i) { + outArrs[i] = OUTPUT_VARIABLE(i); + } + helpers::mergeAddBp(block.launchContext(), *gradient, outArrs); + + return Status::OK(); + } + + DECLARE_TYPES(mergeadd_bp) { + getOpDescriptor() + ->setAllowedInputTypes(sd::DataType::ANY) + ->setAllowedOutputTypes(sd::DataType::ANY); + } + DECLARE_SHAPE_FN(mergeadd_bp) { + + const int numOfInArrs = block.width() - 1; + + auto shapeList = SHAPELIST(); + + for (int e = 0; e < numOfInArrs; e++) { + auto inShape = inputShape->at(e); + shapeList->push_back(ConstantShapeHelper::getInstance()->createShapeInfo(ShapeDescriptor(ArrayOptions::dataType(inShape), shape::order(inShape), shape::shapeOf(inShape), shape::rank(inShape)))); + } + + return shapeList; + } + } } diff --git a/libnd4j/include/ops/declarable/generic/transforms/merge_avg.cpp b/libnd4j/include/ops/declarable/generic/transforms/merge_avg.cpp index 53e46c16e..83a448170 100644 --- a/libnd4j/include/ops/declarable/generic/transforms/merge_avg.cpp +++ b/libnd4j/include/ops/declarable/generic/transforms/merge_avg.cpp @@ -33,7 +33,7 @@ OP_IMPL(mergeavg, -1, 1, false) { auto output = OUTPUT_VARIABLE(0); - std::vector inArrs(block.width()); + std::vector inArrs(block.width()); for(int i = 0; i < block.width(); ++i) inArrs[i] = INPUT_VARIABLE(i); @@ -48,6 +48,44 @@ OP_IMPL(mergeavg, -1, 1, false) { ->setAllowedInputTypes({ALL_FLOATS}) ->setAllowedOutputTypes({ALL_FLOATS}); } + + + CUSTOM_OP_IMPL(mergeavg_bp, 2, 1, false, 0, 0) { + + auto inSize = block.width() - 1; + + REQUIRE_OK(this->validateInputDimensionsMatch(block)); + + std::vector outArrs(inSize); + + const auto gradient = INPUT_VARIABLE(inSize); + + for (int i = 0; i < inSize; ++i) { + outArrs[i] = OUTPUT_VARIABLE(i); + } + helpers::mergeAvgBp(block.launchContext(), *gradient, outArrs); + return Status::OK(); + } + + DECLARE_TYPES(mergeavg_bp) { + getOpDescriptor() + ->setAllowedInputTypes(sd::DataType::ANY) + ->setAllowedOutputTypes(sd::DataType::ANY); + } + DECLARE_SHAPE_FN(mergeavg_bp) { + + const int numOfInArrs = block.width() - 1; + + auto shapeList = SHAPELIST(); + + for (int e = 0; e < numOfInArrs; e++) { + auto inShape = inputShape->at(e); + shapeList->push_back(ConstantShapeHelper::getInstance()->createShapeInfo(ShapeDescriptor(ArrayOptions::dataType(inShape), shape::order(inShape), shape::shapeOf(inShape), shape::rank(inShape)))); + } + + return shapeList; + } + } } diff --git a/libnd4j/include/ops/declarable/generic/transforms/merge_max.cpp b/libnd4j/include/ops/declarable/generic/transforms/merge_max.cpp index 7a41f4c1b..49ab78f7c 100644 --- a/libnd4j/include/ops/declarable/generic/transforms/merge_max.cpp +++ b/libnd4j/include/ops/declarable/generic/transforms/merge_max.cpp @@ -33,7 +33,7 @@ OP_IMPL(mergemax, -1, 1, false) { auto output = OUTPUT_VARIABLE(0); - std::vector inArrs(block.width()); + std::vector inArrs(block.width()); for(int i = 0; i < block.width(); ++i) inArrs[i] = INPUT_VARIABLE(i); @@ -42,7 +42,6 @@ OP_IMPL(mergemax, -1, 1, false) { return Status::OK(); } - DECLARE_SYN(MergeMax, mergemax); DECLARE_TYPES(mergemax) { @@ -51,6 +50,47 @@ DECLARE_SYN(MergeMax, mergemax); ->setAllowedOutputTypes(sd::DataType::ANY); } + + CUSTOM_OP_IMPL(mergemax_bp, 2, 1, false, 0, 0) { + + auto inSize = block.width(); + + REQUIRE_OK(this->validateInputDimensionsMatch(block)); + + std::vector inArrs(inSize); + std::vector outArrs(inSize - 1); + + for (int i = 0; i < inSize; ++i) + inArrs[i] = INPUT_VARIABLE(i); + + for (int i = 0; i < (inSize - 1); ++i) { + outArrs[i] = OUTPUT_NULLIFIED(i); + } + + helpers::mergeMaxBp(block.launchContext(), inArrs, outArrs); + + return Status::OK(); + } + + DECLARE_TYPES(mergemax_bp) { + getOpDescriptor() + ->setAllowedInputTypes(sd::DataType::ANY) + ->setAllowedOutputTypes(sd::DataType::ANY); + } + DECLARE_SHAPE_FN(mergemax_bp) { + + const int numOfInArrs = block.width() - 1; + + auto shapeList = SHAPELIST(); + + for (int e = 0; e < numOfInArrs; e++) { + auto inShape = inputShape->at(e); + shapeList->push_back(ConstantShapeHelper::getInstance()->createShapeInfo(ShapeDescriptor(ArrayOptions::dataType(inShape), shape::order(inShape), shape::shapeOf(inShape), shape::rank(inShape)))); + } + + return shapeList; + } + } } diff --git a/libnd4j/include/ops/declarable/generic/transforms/merge_max_idx.cpp b/libnd4j/include/ops/declarable/generic/transforms/merge_max_idx.cpp index 7fe727452..1ffe42f4b 100644 --- a/libnd4j/include/ops/declarable/generic/transforms/merge_max_idx.cpp +++ b/libnd4j/include/ops/declarable/generic/transforms/merge_max_idx.cpp @@ -32,7 +32,7 @@ CUSTOM_OP_IMPL(mergemaxindex, -1, 1, false, 0, 0) { REQUIRE_OK(this->validateInputDimensionsMatch(block)); auto output = OUTPUT_VARIABLE(0); - std::vector inArrs(block.width()); + std::vector inArrs(block.width()); for(int i = 0; i < block.width(); ++i) inArrs[i] = INPUT_VARIABLE(i); diff --git a/libnd4j/include/ops/declarable/headers/transforms.h b/libnd4j/include/ops/declarable/headers/transforms.h index 0e14037df..29efc4a73 100644 --- a/libnd4j/include/ops/declarable/headers/transforms.h +++ b/libnd4j/include/ops/declarable/headers/transforms.h @@ -64,6 +64,7 @@ namespace sd { #if NOT_EXCLUDED(OP_mergemax) DECLARE_OP(mergemax, -1, 1, false); + DECLARE_CUSTOM_OP(mergemax_bp, 2, 1, false, 0, 0); #endif /* * Complete tensor with max indices merged from all input tensors list @@ -78,10 +79,12 @@ namespace sd { #if NOT_EXCLUDED(OP_mergeadd) DECLARE_OP(mergeadd, -1, 1, false); + DECLARE_CUSTOM_OP(mergeadd_bp, 2, 1, false, 0, 0); #endif #if NOT_EXCLUDED(OP_mergeavg) DECLARE_OP(mergeavg, -1, 1, false); + DECLARE_CUSTOM_OP(mergeavg_bp, 2, 1, false, 0, 0); #endif #if NOT_EXCLUDED(OP_scatter_update) diff --git a/libnd4j/include/ops/declarable/helpers/cpu/clip.cpp b/libnd4j/include/ops/declarable/helpers/cpu/clip.cpp new file mode 100644 index 000000000..d4240d780 --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/clip.cpp @@ -0,0 +1,274 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + + +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + +////////////////////////////////////////////////////////////////////////// +template +static void clipByNorm_(NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace) { + + const int rank = input.rankOf(); + const auto norm2 = input.reduceAlongDimension(reduce::Norm2, dimensions); + + const T normActual = norm2.e(0); + const T normClip = clipNorm.e(0); + + if (isInplace) { + + if(norm2.lengthOf() == 1) { + + if(normActual > normClip) + input *= (normClip / normActual); + } + else { + + auto listOfInSubArrs = input.allTensorsAlongDimension(dimensions); + + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) { + const T iNormActual = norm2.e(i); + if (iNormActual > normClip) + *listOfInSubArrs.at(i) *= normClip / iNormActual; + } + }; + samediff::Threads::parallel_tad(func, 0, listOfInSubArrs.size()); + } + } + else { + + if(norm2.lengthOf() == 1) { + + if(normActual > normClip) + output.assign(input * (normClip / normActual)); + else + output.assign(input); + } + else { + + auto listOfInSubArrs = input.allTensorsAlongDimension(dimensions); + auto listOfOutSubArrs = output.allTensorsAlongDimension(dimensions); + + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) { + auto inputSubArr = listOfInSubArrs.at(i); + auto outputSubArr = listOfOutSubArrs.at(i); + outputSubArr->assign(inputSubArr); + + const T iNormActual = norm2.e(i); + + if (iNormActual > clipNorm.e(0)) + *outputSubArr *= clipNorm / iNormActual; + } + }; + samediff::Threads::parallel_tad(func, 0, listOfInSubArrs.size()); + } + } +} + +////////////////////////////////////////////////////////////////////////// +void clipByNorm(sd::LaunchContext * context, NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace) { + BUILD_SINGLE_SELECTOR(output.dataType(), clipByNorm_, (input, output, dimensions, clipNorm, isInplace), FLOAT_TYPES); +} + + + template + static void clipByGlobalNorm_(std::vector const& inputs, double clipNorm, sd::memory::Workspace* workspace, std::vector& outputs, bool isInplace) { + T globalNorm = 0; //NDArrayFactory::create(0, inputs[0]->getContext()); //sqrt(sum([l2norm(t)**2 for t in t_list])) +// PRAGMA_OMP_PARALLEL_FOR_SIMD_REDUCTION(sumT : globalNorm) + for (size_t i = 0; i < inputs.size(); i++) { + auto input = inputs[i]; + auto l2norm = input->reduceNumber(reduce::Norm2); + globalNorm += l2norm.t(0) * l2norm.t(0); + } + + //globalNorm.applyTransform(transform::Sqrt, nullptr, nullptr);// = sd::math::nd4j_sqrt(globalNorm); + auto normS = sd::math::nd4j_sqrt(globalNorm); + outputs[inputs.size()]->p(0, normS); + + const T factor = clipNorm / normS; + +// PRAGMA_OMP_PARALLEL_FOR + for (size_t e = 0; e < inputs.size(); e++) { + // all-reduce + auto input = inputs[e]; + auto output = outputs[e]; + + if (normS <= clipNorm) { + output->assign(input); + } + else { + + auto lambda = LAMBDA_T(_x, factor) { return _x * factor; }; + input->applyLambda(lambda, *output); + } + } + } + void clipByGlobalNorm(sd::LaunchContext * context, std::vector const& inputs, double clipNorm, sd::memory::Workspace* workspace, std::vector& outputs, bool isInplace) { + BUILD_SINGLE_SELECTOR(outputs[0]->dataType(), clipByGlobalNorm_, (inputs, clipNorm, workspace, outputs, isInplace), FLOAT_TYPES); + } + + BUILD_SINGLE_TEMPLATE(template void clipByGlobalNorm_, (std::vector const& inputs, double clipNorm, sd::memory::Workspace* workspace, std::vector& outputs, bool isInplace), FLOAT_TYPES); + +////////////////////////////////////////////////////////////////////////// +template +static void clipByNormBP_(const NDArray& input, const NDArray& gradO, NDArray& gradI /*output*/, const std::vector& dimensions, const NDArray& clipNorm) { + + const int rank = input.rankOf(); + + auto norm2 = input.reduceAlongDimension(reduce::Norm2, dimensions); + + if(norm2.lengthOf() == 1) { + + const T N = norm2.e(0); + + auto cn = clipNorm.e(0); + + if(N > cn) { + + const T sumOfProd = (input * gradO).reduceNumber(reduce::Sum).e(0); // reduce to scalar + const T factor1 = static_cast(1.f) / N; + const T factor3 = factor1 / (N * N); // 1 / (N*N*N) + + auto lambda = LAMBDA_TT(elem1, elem2, cn, sumOfProd, factor1, factor3) { + return cn * (factor1 * elem2 - factor3 * elem1 * sumOfProd); + }; + + (const_cast(input)).applyPairwiseLambda(const_cast(gradO), lambda, gradI); + } + else + gradI.assign(gradO); + } + else { + + auto gradISubArrs = gradI.allTensorsAlongDimension({dimensions}); + auto gradOSubArrs = gradO.allTensorsAlongDimension({dimensions}); + auto inputSubArrs = input.allTensorsAlongDimension({dimensions}); + + auto cn = clipNorm.e(0); + + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) { + T N = norm2.e(i); + + auto gradOSubArr = gradOSubArrs.at(i); + auto gradISubArr = gradISubArrs.at(i); + + if (N > cn) { + auto inputSubArr = inputSubArrs.at(i); + const T sumOfProd = (*inputSubArr * *gradOSubArr).reduceNumber(reduce::Sum).e(0); // reduce to scalar + const T factor1 = static_cast(1.f) / N; + const T factor3 = factor1 / (N * N); // 1 / (N*N*N) + + auto lambda = LAMBDA_TT(elem1, elem2, cn, sumOfProd, factor1, factor3) { + return cn * (factor1 * elem2 - factor3 * elem1 * sumOfProd); + }; + + inputSubArr->applyPairwiseLambda(*gradOSubArr, lambda, *gradISubArr); + } else + gradISubArr->assign(gradOSubArr); + } + }; + samediff::Threads::parallel_tad(func, 0, gradISubArrs.size()); + } +} + + void clipByNormBP(sd::LaunchContext * context, const NDArray& input, const NDArray& gradO, NDArray& gradI /*output*/, const std::vector& dimensions, const NDArray& clipNorm) { + BUILD_SINGLE_SELECTOR(gradI.dataType(), clipByNormBP_, (input, gradO, gradI, dimensions, clipNorm), FLOAT_TYPES); + } + + BUILD_SINGLE_TEMPLATE(template void clipByNormBP_, (const NDArray& input, const NDArray& gradO, NDArray& gradI /*output*/, const std::vector& dimensions, const NDArray& clipNorm), FLOAT_TYPES); + + +////////////////////////////////////////////////////////////////////////// +template +static void clipByAveraged_(NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace) { + + auto cn = clipNorm.e(0); + if (dimensions.size() == 0) { + // all-reduce + T n2 = input.reduceNumber(reduce::Norm2).e(0) / input.lengthOf(); + if (n2 <= cn) { + if (!isInplace) + output.assign(input); + } + else { + const T factor = cn / n2; + auto lambda = LAMBDA_T(_x, factor) { return _x * factor; }; + input.applyLambda(lambda, output); + } + } + else { + // along dimension + auto norm2 = input.reduceAlongDimension(reduce::Norm2, dimensions, false); + if (!isInplace) + output.assign(input); + auto tads = output.allTensorsAlongDimension(dimensions); + // TODO: make this CUDA-compliant somehow + for (int e = 0; e < tads.size(); e++) { + T n2 = norm2.e(e) / tads.at(e)->lengthOf(); + const T factor = cn / n2; + if (n2 > cn) { + auto lambda = LAMBDA_T(_x, factor) {return _x * factor;}; + tads.at(e)->applyLambda(lambda, output); + } + } + } +} + + void clipByAveraged(sd::LaunchContext * context, NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace) { + BUILD_SINGLE_SELECTOR(input.dataType(), clipByAveraged_, (input, output, dimensions, clipNorm, isInplace), FLOAT_TYPES); + } + + BUILD_SINGLE_TEMPLATE(template void clipByAveraged_, (NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace), FLOAT_TYPES); + +/* + if (d1 > params[1]) + return params[1]; + else if (d1 < params[0]) + return params[0]; + else return d1; +*/ + + template + static void clipByValue_(NDArray& input, double leftBound, double rightBound, NDArray& output) { + auto routine = LAMBDA_T(_x, leftBound, rightBound) { + if (_x > rightBound) return rightBound; + if (_x < leftBound) return leftBound; + return _x; + }; + + input.applyLambda(routine, output); + } + + void clipByValue(sd::LaunchContext * context, NDArray& input, double leftBound, double rightBound, NDArray& output) { + BUILD_SINGLE_SELECTOR(input.dataType(), clipByValue_, (input, leftBound, rightBound, output), FLOAT_TYPES); + } + + BUILD_SINGLE_TEMPLATE(template void clipByValue_, (NDArray& input, double leftBound, double rightBound, NDArray& output);, FLOAT_TYPES); + +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/eye.cpp b/libnd4j/include/ops/declarable/helpers/cpu/eye.cpp new file mode 100644 index 000000000..30a83b871 --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/eye.cpp @@ -0,0 +1,45 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + + +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + +////////////////////////////////////////////////////////////////////////// +void eye(sd::LaunchContext * context, NDArray& output) { + + const int rank = output.rankOf(); + auto arrs = output.allTensorsAlongDimension({rank-2, rank-1}); + + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) + arrs.at(i)->setIdentity(); + }; + + samediff::Threads::parallel_tad(func, 0, arrs.size()); +} + +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/gatherTransforms.cpp b/libnd4j/include/ops/declarable/helpers/cpu/gatherTransforms.cpp new file mode 100644 index 000000000..f7cb1cf59 --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/gatherTransforms.cpp @@ -0,0 +1,183 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + + + +#include +#include +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + + +//////////////////////////////////////////////////////////////////////// +template +static void gatherND_(NDArray& input, NDArray& indices, NDArray& output) { + + const X* x = reinterpret_cast(input.getBuffer()); + const Y* y = reinterpret_cast(indices.getBuffer()); + X* z = reinterpret_cast(output.getBuffer()); + + const int xRank = input.rankOf(); + const int yRank = indices.rankOf(); + const int zRank = output.rankOf(); + const int maxRank = sd::math::nd4j_max(yRank, sd::math::nd4j_max(xRank, zRank)); + + const Nd4jLong zLen = output.lengthOf(); + + const uint yLastDim = indices.sizeAt(-1); + + const int diff = zRank - xRank; + const bool bEqual = yLastDim == xRank; + + auto func = PRAGMA_THREADS_FOR { + + int xCoords[MAX_RANK], zCoords[MAX_RANK], temp; + + for (auto i = start; i < stop; i++) { + + shape::index2coordsCPU(start, i, output.getShapeInfo(), zCoords); + + const auto zOffset = shape::getOffset(output.getShapeInfo(), zCoords); + + temp = zCoords[yRank - 1]; + zCoords[yRank - 1] = 0; + const auto yOffset = shape::getOffset(indices.getShapeInfo(), zCoords); + zCoords[yRank - 1] = temp; + + if(bEqual) + memcpy(xCoords, zCoords, zRank * sizeof(int)); + else if(diff >= 0) + memcpy(xCoords, zCoords + diff, xRank * sizeof(int)); + else + memcpy(xCoords - diff, zCoords, zRank * sizeof(int)); + + for (uint j = 0; j < yLastDim; ++j) + xCoords[j] = y[yOffset + j * indices.stridesOf()[yRank - 1]]; // last stride + + const auto xOffset = shape::getOffset(input.getShapeInfo(), xCoords); + + z[zOffset] = x[xOffset]; + } + }; + + samediff::Threads::parallel_tad(func, 0, zLen); +} + +//////////////////////////////////////////////////////////////////////// +void gatherND(sd::LaunchContext * context, NDArray& input, NDArray& indices, NDArray& output) { + BUILD_DOUBLE_SELECTOR(input.dataType(), indices.dataType(), gatherND_, (input, indices, output), LIBND4J_TYPES, INDEXING_TYPES); +} + + +//////////////////////////////////////////////////////////////////////// +template +static void gather_(NDArray* input, const NDArray* indices, NDArray* output, const std::vector& intArgs) { + + int axis = intArgs.size() > 0 ? intArgs[0] : 0; + const int inputRank = input->rankOf(); + if(axis < 0) + axis += inputRank; + + const int numOfIntArgs = intArgs.size(); + + if (indices != nullptr) { + + for(Nd4jLong i = 0; i < indices->lengthOf(); ++i) + if(indices->e(i) >= input->sizeAt(axis)) + throw std::runtime_error("helpers::gather function: indices array contains wrong elements, each element must be smaller than corresponding dimension of input array !"); + + // first case: indices consist of only one scalar + if(indices->isScalar()) { + if(input->rankOf() <= 1){ + //For scalar indices, rank 0 or 1 input: can't do tensor along dimension 0 as this is whole array... instead, we want to get a scalar + auto idx = indices->e(0); + auto scalarNDArray = input->e(idx); + output->assign(scalarNDArray); + } else { + auto dimensions = ShapeUtils::evalDimsToExclude(input->rankOf(), {axis}); + auto tadPack = sd::ConstantTadHelper::getInstance()->tadForDimensions(input->getShapeInfo(), dimensions); + + auto tadArr = NDArray(reinterpret_cast(reinterpret_cast(input->getBuffer()) + tadPack.primaryOffsets()[indices->e(0)]), tadPack.primaryShapeInfo(), output->getContext()); + output->assign(&tadArr); + } + } + else if (input->rankOf() == 1 && indices->isVector()) { + // special case + auto func = PRAGMA_THREADS_FOR { + for (auto e = start; e < stop; e++) + output->p(e, input->e(indices->e(e))); + }; + + samediff::Threads::parallel_for(func, 0, indices->lengthOf()); + } + else { + + std::vector dimsOut(indices->rankOf()); + std::iota(dimsOut.begin(), dimsOut.end(), axis); // fill with axis, axis+1, ... indices->rankOf()-1 + const Nd4jLong numOfSubArrs = ShapeUtils::getNumOfSubArrs(output->getShapeInfo(), dimsOut); + + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) { + NDArray subArrOut = (*output)(i, dimsOut); + NDArray subArrIn = (*input)(indices->e(i), {axis}); + subArrOut.assign(subArrIn); + } + }; + + samediff::Threads::parallel_tad(func, 0, numOfSubArrs); + } + } + else { + + for(int i = 1; i < numOfIntArgs; ++i) + if(intArgs[i] >= input->sizeAt(axis)) + throw std::runtime_error("helpers::gather function: some of input indexes is larger than corresponding shape of input array !"); + + // we only allow scalar/vector case here + if (numOfIntArgs == 2) { // scalar case + output->assign((*input)(intArgs[1], {axis})); + } + else { // vector case + const Nd4jLong numOfSubArrs = ShapeUtils::getNumOfSubArrs(output->getShapeInfo(), {axis}); + + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) { + NDArray subArrOut = (*output)(i, {axis}); + NDArray subArrIn = (*input)(intArgs[i + 1], {axis}); + subArrOut.assign(subArrIn); + } + }; + + samediff::Threads::parallel_tad(func, 0, numOfSubArrs); + } + } +} + + void gather(NDArray* input, const NDArray* indices, NDArray* output, const std::vector& intArgs) { + BUILD_SINGLE_SELECTOR(input->dataType(), gather_, (input, indices, output, intArgs), LIBND4J_TYPES); + } + +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/invertPermutation.cpp b/libnd4j/include/ops/declarable/helpers/cpu/invertPermutation.cpp new file mode 100644 index 000000000..5325ac282 --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/invertPermutation.cpp @@ -0,0 +1,51 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + + +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + +//////////////////////////////////////////////////////////////////////// +void invertPermutation(sd::LaunchContext * context, const NDArray& input, NDArray& output) { + + std::set uniqueElems; + const int length = input.lengthOf(); + + for(int i = 0; i < length; ++i) { + + int elem = input.e(i); + + if(!uniqueElems.insert(elem).second) // this operation forbids us to use #pragma omp + throw std::runtime_error("helpers::invertPermutation function: input array contains duplicates !"); + + if(elem < 0 || elem > length - 1) + throw std::runtime_error("helpers::invertPermutation function: element of input array is out of range (0, length-1) !"); + + output.p(elem, i); + } +} + +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/merge.cpp b/libnd4j/include/ops/declarable/helpers/cpu/merge.cpp new file mode 100644 index 000000000..74007635f --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/merge.cpp @@ -0,0 +1,277 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * Copyright (c) 2019-2020 Konduit K.K. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// @author Oleh Semeniv (oleg.semeniv@gmail.com) +// + +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + + +////////////////////////////////////////////////////////////////////////// +template +static void mergeMaxIndex_(const std::vector& inArrs, NDArray& output) { + + const Nd4jLong numArgs = inArrs.size(); + auto x = inArrs[0]; + + auto func = PRAGMA_THREADS_FOR { + for (auto e = start; e < stop; e++) { + T max = -DataTypeUtils::max(); + Nd4jLong idx = 0; + + for (Nd4jLong i = 0; i < numArgs; i++) { + T v = inArrs[i]->e(e); + if (v > max) { + max = v; + idx = i; + } + } + output.p(e, idx); + } + }; + + samediff::Threads::parallel_for(func, 0, x->lengthOf()); +} + +void mergeMaxIndex(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { + BUILD_SINGLE_SELECTOR(inArrs[0]->dataType(), mergeMaxIndex_, (inArrs, output), LIBND4J_TYPES); +} + + +////////////////////////////////////////////////////////////////////////// +template +static void mergeMax_(const std::vector& inArrs, NDArray& output) { + + const Nd4jLong numArgs = inArrs.size(); + auto x = inArrs[0]; + + auto func = PRAGMA_THREADS_FOR { + for (auto e = start; e < stop; e++) { + T max = -DataTypeUtils::max(); + for (Nd4jLong i = 0; i < numArgs; i++) { + T v = inArrs[i]->e(e); + if (v > max) + max = v; + } + output.p(e, max); + } + }; + + samediff::Threads::parallel_for(func, 0, x->lengthOf()); +} + +void mergeMax(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { + BUILD_SINGLE_SELECTOR(output.dataType(), mergeMax_, (inArrs, output), LIBND4J_TYPES); +} + + +////////////////////////////////////////////////////////////////////////// +template +static void mergeMaxBp_(const std::vector& inArrs, std::vector& outArrs) { + + // outArrs.size() == inArrs.size() - 1 + const Nd4jLong numArgs = outArrs.size(); + // last array is gradient + const auto gradient = inArrs[numArgs]->bufferAsT(); + auto length = inArrs[numArgs]->lengthOf(); + + bool bSameOrderAndEws1 = (1 == inArrs[numArgs]->ews()); + + if (bSameOrderAndEws1) { + auto gradOrdering = inArrs[numArgs]->ordering(); + + for (int i = 0; i < numArgs; ++i) { + bSameOrderAndEws1 &= (gradOrdering == inArrs[i]->ordering()); + bSameOrderAndEws1 &= (1 == inArrs[i]->ews()); + bSameOrderAndEws1 &= (gradOrdering == outArrs[i]->ordering()); + bSameOrderAndEws1 &= (1 == outArrs[i]->ews()); + } + } + + + if(bSameOrderAndEws1){ + auto func = PRAGMA_THREADS_FOR{ + for (auto e = start; e < stop; e++) { + T max = -DataTypeUtils::max(); + Nd4jLong nMaxIndex = 0; + for (Nd4jLong i = 0; i < numArgs; i++) { + const T* v = inArrs[i]->bufferAsT(); + if (v[e] > max) { + max = v[e]; + nMaxIndex = i; + } + } + T* z = outArrs[nMaxIndex]->bufferAsT(); + z[e] = gradient[e]; + } + }; + + samediff::Threads::parallel_for(func, 0, length); + return; + } + + auto gradShape = inArrs[numArgs]->getShapeInfo(); + std::vector vbSameShaepeAndStrides(numArgs); + for (int i = 0; i < numArgs; ++i) { + vbSameShaepeAndStrides[i] = shape::haveSameShapeAndStrides(gradShape, inArrs[i]->getShapeInfo()); + } + + auto func = PRAGMA_THREADS_FOR{ + + int coords[MAX_RANK]; + for (auto e = start; e < stop; e++) { + + shape::index2coordsCPU(start, e, gradShape, coords); + + const auto gradOffset = shape::getOffset(gradShape, coords); + + T max = -DataTypeUtils::max(); + Nd4jLong nMaxIndex = 0; + + for (Nd4jLong i = 0; i < numArgs; i++) { + + const auto xOffset = vbSameShaepeAndStrides[i] ? gradOffset : shape::getOffset(inArrs[i]->getShapeInfo(), coords); + const T* v = inArrs[i]->bufferAsT(); + if (v[xOffset] > max) { + max = v[xOffset]; + nMaxIndex = i; + } + } + + const auto zOffset = vbSameShaepeAndStrides[nMaxIndex] ? gradOffset : shape::getOffset(outArrs[nMaxIndex]->getShapeInfo(), coords); + + T* z = outArrs[nMaxIndex]->bufferAsT(); + z[zOffset] = gradient[gradOffset]; + } + }; + + samediff::Threads::parallel_for(func, 0, length); + return; +} + +void mergeMaxBp(sd::LaunchContext* context, const std::vector& inArrs, std::vector& outArrs) { + BUILD_SINGLE_SELECTOR(outArrs[0]->dataType(), mergeMaxBp_, (inArrs, outArrs), LIBND4J_TYPES); +} + +////////////////////////////////////////////////////////////////////////// +template +static void mergeAvg_(const std::vector& inArrs, NDArray& output) { + const Nd4jLong numArgs = inArrs.size(); + const T factor = 1.f / numArgs; + auto x = inArrs[0]; + + auto func = PRAGMA_THREADS_FOR { + for (auto e = start; e < stop; e++) { + T sum = 0.; + for (Nd4jLong i = 0; i < numArgs; i++) { + T v = inArrs[i]->e(e); + sum += v; + } + output.p(e, sum * factor); + } + }; + + samediff::Threads::parallel_for(func, 0, x->lengthOf()); +} + +void mergeAvg(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { + BUILD_SINGLE_SELECTOR(output.dataType(), mergeAvg_, (inArrs, output), LIBND4J_TYPES); +} + +////////////////////////////////////////////////////////////////////////// +template +static void mergeAvgBp_(const NDArray& gradient, std::vector& outArrs) { + + const Nd4jLong numArgs = outArrs.size(); + + auto func = PRAGMA_THREADS_FOR{ + for (auto e = start; e < stop; e++) { + + T v = gradient.e(e) / numArgs; + + for (Nd4jLong i = 0; i < numArgs; i++) { + outArrs[i]->p(e, v); + } + } + }; + + samediff::Threads::parallel_for(func, 0, gradient.lengthOf()); +} + +void mergeAvgBp(sd::LaunchContext* context, const NDArray& gradient, std::vector& outArrs) { + BUILD_SINGLE_SELECTOR(gradient.dataType(), mergeAvgBp_, (gradient, outArrs), LIBND4J_TYPES); +} + + +////////////////////////////////////////////////////////////////////////// +template +static void mergeAdd_(const std::vector& inArrs, NDArray& output) { + + const Nd4jLong numArgs = inArrs.size(); + auto x = inArrs[0]; + + auto func = PRAGMA_THREADS_FOR { + for (auto e = start; e < stop; e++) { + T sum = (T) 0.f; + for (Nd4jLong i = 0; i < numArgs; i++) + sum += inArrs[i]->e(e); + + output.p(e, sum); + } + }; + + samediff::Threads::parallel_for(func, 0, x->lengthOf()); +} + void mergeAdd(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { + BUILD_SINGLE_SELECTOR(output.dataType(), mergeAdd_, (inArrs, output), LIBND4J_TYPES); + } + +////////////////////////////////////////////////////////////////////////// +template +static void mergeAddBp_(const NDArray& gradient, std::vector& outArrs) { + + const Nd4jLong numArgs = outArrs.size(); + + auto func = PRAGMA_THREADS_FOR{ + for (auto e = start; e < stop; e++) { + + T v = gradient.e(e); + + for (Nd4jLong i = 0; i < numArgs; i++) { + outArrs[i]->p(e, v); + } + } + }; + + samediff::Threads::parallel_for(func, 0, gradient.lengthOf()); +} + +void mergeAddBp(sd::LaunchContext* context, const NDArray& gradient, std::vector& outArrs) { + BUILD_SINGLE_SELECTOR(gradient.dataType(), mergeAddBp_, (gradient, outArrs), LIBND4J_TYPES); +} + + +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/pad.cpp b/libnd4j/include/ops/declarable/helpers/cpu/pad.cpp new file mode 100644 index 000000000..b303d95ae --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/pad.cpp @@ -0,0 +1,483 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + + +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + + +////////////////////////////////////////////////////////////////////////// +template +void pad_(const int mode, const NDArray& input, const NDArray& paddings, NDArray& output, const NDArray& padValue) { + + const T* x = input.bufferAsT(); + T* z = output.bufferAsT(); + + const Nd4jLong* xShape = input.shapeOf(); + const Nd4jLong* zShape = output.shapeOf(); + + const int rank = input.rankOf(); // both input and output have the same rank + const int rankMinusOne = rank - 1; + + const auto zLen = output.lengthOf(); + + if(mode == 0) { // CONSTANT case + + const T padVal = padValue.e(0); + + auto func = PRAGMA_THREADS_FOR { + + int zCoords[MAX_RANK], xCoords[MAX_RANK]; + + for (auto i = start; i < stop; i++) { + + shape::index2coordsCPU(start, i, output.getShapeInfo(), zCoords); + const auto zOffset = shape::getOffset(output.getShapeInfo(), zCoords); + + memcpy(xCoords, zCoords, rank * sizeof(int)); + + bool within = true; + + for (int j = rankMinusOne; j >= 0; --j) { + + if (xShape[j] == zShape[j]) + continue; + + const auto left = paddings.e(j, 0); + + if (zCoords[j] < left || zCoords[j] >= left + xShape[j]) { + within = false; + break; + } + else + xCoords[j] = zCoords[j] - left; + } + + if (within) + z[zOffset] = x[shape::getOffset(input.getShapeInfo(), xCoords)]; + else + z[zOffset] = padVal; + } + }; + + samediff::Threads::parallel_tad(func, 0, zLen); + } + else { // REFLECT and SYMMETRIC cases + + const Nd4jLong shift1 = mode == 1 ? 0 : 1; // REFLECT : SYMMETRIC + const Nd4jLong shift2 = mode == 1 ? 2 : 1; // REFLECT : SYMMETRIC + + auto func = PRAGMA_THREADS_FOR { + + int zCoords[MAX_RANK], xCoords[MAX_RANK]; + + for (auto i = start; i < stop; i++) { + + shape::index2coordsCPU(start, i, output.getShapeInfo(), zCoords); + const auto zOffset = shape::getOffset(output.getShapeInfo(), zCoords); + + memcpy(xCoords, zCoords, rank * sizeof(int)); + + for (int j = rankMinusOne; j >= 0; --j) { + + if (xShape[j] == zShape[j]) + continue; + + xCoords[j] = zCoords[j] - paddings.e(j, 0); // are ready to fill middle (within input dimension range) + + if (xCoords[j] < 0) + xCoords[j] = -xCoords[j] - shift1; // means fill from left + else if (xCoords[j] >= xShape[j]) + xCoords[j] = 2 * xShape[j] - xCoords[j] - shift2; // means fill from right + } + + const auto xOffset = shape::getOffset(input.getShapeInfo(), xCoords); + z[zOffset] = x[xOffset]; + } + }; + + samediff::Threads::parallel_tad(func, 0, zLen); + } +} + +// ////////////////////////////////////////////////////////////////////////// +// template +// void pad2_(const int mode, const NDArray& input, const NDArray& paddings, NDArray& output, NDArray const& padValue) { + +// const int rank = output.rankOf(); +// std::vector dimsToExclude(rank); +// std::iota(dimsToExclude.begin(), dimsToExclude.end(), 0); // fill with 0, 1, ... rank-1 + +// Nd4jLong numLeft = paddings.e(rank-1,0); +// Nd4jLong numRight = paddings.e(rank-1,1); +// Nd4jLong inDimSize = input.sizeAt(rank-1); +// Nd4jLong outDimSize = output.sizeAt(rank-1); + +// std::vector> outIdx = { std::vector(2*rank), {numLeft, numLeft + inDimSize}, {0, numLeft}, {numLeft + inDimSize, outDimSize} }; + +// for(int i = 0; i < rank-1; ++i) { +// outIdx[0][2*i] = paddings.e(i, 0); +// outIdx[0][2*i + 1] = outIdx[0][2*i] + input.sizeAt(i); +// } +// outIdx[0][2*rank-1] = outIdx[0][2*rank-2] = 0; + +// // ***** populate innermost sub-arrays firstly ***** // +// dimsToExclude.pop_back(); + +// Nd4jLong startL = mode == 1 ? 1 : 0; // REFLECT or SYMMETRIC +// Nd4jLong startR = mode == 1 ? inDimSize-2 : inDimSize-1; // REFLECT or SYMMETRIC + +// Nd4jLong numOfSubArrs = ShapeUtils::getNumOfSubArrs(input.getShapeInfo(), dimsToExclude); + +// NDArray outSubArr0 = output(outIdx[0], true); + +// PRAGMA_OMP_PARALLEL_FOR +// for(Nd4jLong j = 0; j < numOfSubArrs; ++j) { + +// NDArray outSubArr1 = outSubArr0(j, dimsToExclude); +// NDArray inSubArr = input(j, dimsToExclude); +// NDArray outSubArrMid = outSubArr1(outIdx[1]); + +// outSubArrMid.assign(inSubArr); // assign middle + +// if(mode == 0) { // CONSTANT +// if(numLeft != 0) { +// NDArray temp = outSubArr1(outIdx[2]); +// temp.assign(padValue); // assign left +// } +// if(numRight != 0) { +// NDArray temp = outSubArr1(outIdx[3]); +// temp.assign(padValue); // assign right +// } +// } +// else { // REFLECT or SYMMETRIC + +// for(Nd4jLong k = numLeft-1, e = startL; k >= 0; --k, ++e) // fill left side +// outSubArr1.t(k) = inSubArr.t(e); + +// for(Nd4jLong k = numLeft + inDimSize, e = startR; k < outDimSize; ++k, --e) // fill right side +// outSubArr1.t(k) = inSubArr.t(e); +// } +// } + +// // ***** fill rest of outer sub-arrays ***** // +// std::vector outIdxInner(2, 0); +// std::vector outIdxOuter(2, 0); + +// for(int i = rankBorder - 1; i >= 0; --i) { + +// dimsToExclude.pop_back(); + +// outIdxInner.push_back(0), outIdxInner.push_back(0); +// outIdxOuter.push_back(0), outIdxOuter.push_back(0); + +// Nd4jLong numLeft = paddings.e(i, 0); +// Nd4jLong numRight = paddings.e(i, 1); + +// if(numLeft == 0 && numRight == 0) +// continue; + +// Nd4jLong inDimSize = input.sizeAt(i); +// Nd4jLong outDimSize = output.sizeAt(i); + +// if(mode == 0) { +// outIdxOuter[0] = 0; outIdxOuter[1] = numLeft; +// outIdxInner[0] = numLeft + inDimSize; outIdxInner[1] = outDimSize; +// } + +// startL = mode == 1 ? numLeft + 1 : numLeft; // REFLECT or SYMMETRIC +// startR = mode == 1 ? numLeft + inDimSize - 2 : numLeft + inDimSize-1; // REFLECT or SYMMETRIC + +// numOfSubArrs = ShapeUtils::getNumOfSubArrs(output.getShapeInfo(), dimsToExclude); + +// PRAGMA_OMP_PARALLEL_FOR_ARGS(firstprivate(outIdxOuter, outIdxInner)) +// for(Nd4jLong j = 0; j < numOfSubArrs; ++j) { + +// NDArray outSubArr = output(j, dimsToExclude); + +// if(mode == 0) { // CONSTANT + +// if(numLeft != 0) { +// NDArray tempO = outSubArr(outIdxOuter); +// tempO.assign(padValue); // assign left +// } + +// if(numRight != 0) { +// NDArray tempI = outSubArr(outIdxInner); +// tempI.assign(padValue); // assign right +// } +// } +// else { // REFLECT or SYMMETRIC + +// for(Nd4jLong k = numLeft-1, e = startL; k >= 0; --k, ++e) { // fill left side +// outIdxOuter[0] = k; +// outIdxOuter[1] = k+1; +// outIdxInner[0] = e; +// outIdxInner[1] = e+1; +// NDArray outSubArrInner = outSubArr(outIdxInner); +// NDArray outSubArrOuter = outSubArr(outIdxOuter); +// outSubArrOuter.assign(outSubArrInner); +// } + +// for(Nd4jLong k = numLeft + inDimSize, e = startR; k < outDimSize; ++k, --e) { // fill right side +// outIdxOuter[0] = k; +// outIdxOuter[1] = k+1; +// outIdxInner[0] = e; +// outIdxInner[1] = e+1; +// NDArray outSubArrInner = outSubArr(outIdxInner); +// NDArray outSubArrOuter = outSubArr(outIdxOuter); +// outSubArrOuter.assign(outSubArrInner); +// } +// } +// } +// } +// } + +void pad(sd::LaunchContext * context, const int mode, const NDArray& input, const NDArray& paddings, NDArray& output, NDArray const& padValue) { + BUILD_SINGLE_SELECTOR(input.dataType(), pad_, (mode, input, paddings, output, padValue), LIBND4J_TYPES); +} + +////////////////////////////////////////////////////////////////////////// +template +static void mirrorPad_(const NDArray& input, const NDArray& paddings, NDArray& output, const int mode) { + + // mode: 0 - REFLECT, else - SYMMETRIC + const int reflBorder = (bool)mode ? 1 : 0; + const int rank = input.rankOf(); + const Nd4jLong outLen = output.lengthOf(); + + if(rank <= 1) { + + const Nd4jLong inLen = input.lengthOf(); + const auto leftSide = paddings.e(0); + const auto leftSideCorrected = leftSide - reflBorder; + const Nd4jLong len = 2*(inLen-1) + leftSide + reflBorder; + + for(int i = 0; i < outLen; ++i) { + + if (i < leftSide) // left side + output.p(i, input.e(leftSideCorrected - i)); + + else if(i >= leftSide && i < leftSide + inLen) // middle + output.p(i, input.e(i - leftSide)); + + else // right side + output.p(i, input.e(len - i)); + } + } + else { + + auto func = PRAGMA_THREADS_FOR { + + int inIdx[MAX_RANK], outIdx[MAX_RANK]; + + for (auto i = start; i < stop; i++) { + + shape::index2coordsCPU(start, i, output.getShapeInfo(), outIdx); + + for (int j = 0; j < rank; ++j) { + const Nd4jLong inLen = input.sizeAt(j); + const auto leftSide = paddings.e(j, 0); + const auto leftSideCorrected = leftSide - reflBorder; + const Nd4jLong len = 2 * (inLen - 1) + leftSide + reflBorder; + + if (outIdx[j] < leftSide) // left side + inIdx[j] = leftSideCorrected - outIdx[j]; + + else if (outIdx[j] >= leftSide && outIdx[j] < leftSide + inLen) // middle + inIdx[j] = outIdx[j] - leftSide; + + else // right side + inIdx[j] = len - outIdx[j]; + } + + auto outOffset = shape::getOffset(output.getShapeInfo(), outIdx); + auto inOffset = shape::getOffset(input.getShapeInfo(), inIdx); + reinterpret_cast(output.buffer())[outOffset] = reinterpret_cast(input.getBuffer())[inOffset]; + } + }; + + samediff::Threads::parallel_for(func, 0, outLen); + } +} + + void mirrorPad(sd::LaunchContext * context, const NDArray& input, const NDArray& paddings, NDArray& output, const int mode) { + BUILD_SINGLE_SELECTOR(input.dataType(), mirrorPad_, (input, paddings, output, mode), LIBND4J_TYPES); + } + + BUILD_SINGLE_TEMPLATE(template void mirrorPad_, (const NDArray& input, const NDArray& paddings, NDArray& output, const int mode), LIBND4J_TYPES); + + +//////////////////////////////////////////////////////////////////////// +/*// initial values of inIdx, outIdx, dim must be equal to zero +template +static void recursiveLoopForPad_(const int mode, NDArray& input, const NDArray& paddings, NDArray& output, std::vector dimensions, int dim, int inIdx, int outIdx, NDArray& padValue ) { + + int leftOffset; + // dimensions are array of input dimensions, it is sorted in increasing order + // every time at the beginning we erase first element from it (not good idea to use vector for this purpose, but luckily it is small enough) + // then we use this array for tads building, every time while recursion the number of built tads becomes bigger + dimensions.erase(dimensions.begin()); + // build tad basing on output array, also create auxiliary arrays pointing on required output array ranges + shape::TAD tadOut(output.getShapeInfo(), dimensions.data(), dimensions.size()); + tadOut.createTadOnlyShapeInfo(); + tadOut.createOffsets(); + auto subArrOut = NDArray(output.getBuffer(), tadOut.tadOnlyShapeInfo, output.getContext()); + auto subArr = NDArray(output.getBuffer(), tadOut.tadOnlyShapeInfo, output.getContext()); + // build tad basing on input array, also create auxiliary array pointing on required input array range + shape::TAD tadIn(input.getShapeInfo(), dimensions.data(), dimensions.size()); + tadIn.createTadOnlyShapeInfo(); + tadIn.createOffsets(); + auto subArrIn = NDArray(input.getBuffer(), tadIn.tadOnlyShapeInfo, output.getContext()); + // these indices take into account recursion and always point to actual tads numbers + if (input.rankOf() > 1 && output.rankOf() > 1) {// only for non-vector cases + outIdx = outIdx * output.sizeAt(dim + 1); + inIdx = inIdx * input.sizeAt(dim + 1); + } + // current input tad number, we add to it unity in a loop + int k = -1; + // loop through current dimension + for(int i = 0; i < output.sizeAt(dim); ++i) { + // corresponds to outer range (relevant indices are absent in input) + leftOffset = paddings.e(dim, 0); + if(i < leftOffset || i >= (input.sizeAt(dim) + leftOffset)) + continue; + + // increase input tads number + ++k; + // recursion condition allows for the fact that tad can't reduce to scalar + if(dim < input.rankOf() - 2) + recursiveLoopForPad(mode, input, paddings, output, dimensions, dim + 1, inIdx + k, outIdx + i, padValue); + else if (paddings.sizeAt(0) > dim + 1){ + leftOffset = paddings.e(dim + 1, 0); + // shift buffers pointers to actual element position + if (output.rankOf() > 1) { + subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + i]); + subArrIn.setBuffer(reinterpret_cast(input.getBuffer()) + tadIn.tadOffsets[inIdx + i - paddings.e(dim, 0)]); + } + else { + subArrOut.p(i, subArrIn.e(i - leftOffset)); + } + // most inner loop, corresponds to last dim = rank-1 + switch (mode) { + case 0: // CONSTANT mode + for(int j = 0; j < subArrOut.lengthOf(); ++j) + if(j < leftOffset || j >= (subArrIn.lengthOf() + leftOffset) ) // firstly fill with zeros outer ranges + subArrOut.p(j, (T)0.f); + else + subArrOut.p(j, subArrIn.e(j - leftOffset)); // fill middle with elements of input array + break; + + case 1: // REFLECT mode + for(int j = 1; j <= leftOffset; ++j) // fill firstly left side + subArrOut.p(leftOffset - j, subArrIn.e(j)); + for(int j = 0; j < subArrIn.lengthOf(); ++j) // fill middle + subArrOut.p(leftOffset + j, subArrIn.e(j)); + for(int j = (subArrOut.lengthOf() - leftOffset); j < subArrOut.lengthOf(); ++j) // fill right side + subArrOut.p(j, subArrIn.e(subArrOut.lengthOf() - j - 1)); + break; + + case 2: // SYMMETRIC mode + for(int j = 1; j <= leftOffset; ++j) // fill firstly left side + subArrOut.p(leftOffset - j, subArrIn.e(j-1)); + for(int j = 0; j < subArrIn.lengthOf(); ++j) // fill middle + subArrOut.p(leftOffset + j, subArrIn.e(j)); + for(int j = (subArrOut.lengthOf() - leftOffset); j < subArrOut.lengthOf(); ++j) // fill right side + subArrOut.p(j, subArrIn.e(subArrOut.lengthOf() - j)); + break; + } + } + else { + + if (mode == 0 && input.rankOf() < 2) + subArrOut.p(i, subArrIn.e(i - leftOffset)); // fill middle with elements of input array + } + } + // populate sub-array formed previously + leftOffset = paddings.e(dim,0); + switch (mode) { + case 0: // CONSTANT mode + for(int j = 1; j <= leftOffset; ++j) { + // fill left side with padValue + if (output.rankOf() > 1) { + subArrOut.setBuffer( + reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset - j]); + subArrOut.assign(padValue); + } + else { + subArrOut.p(j - 1, padValue); + } + } +// output.printIndexedBuffer("Output at"); + for(int j = (output.sizeAt(dim) - leftOffset); j < output.sizeAt(dim); ++j) { // fill left side with zeros + if (output.rankOf() > 1) { + subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + j]); + subArrOut.assign(padValue); + } + else { + subArrOut.p(j, padValue); + } + } + break; + + case 1: // REFLECT mode + for(int j = 1; j <= leftOffset; ++j) { // fill left side + subArr.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset + j]); + subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset - j]); + subArrOut.assign(&subArr); + } + for(int j = (output.sizeAt(dim) - leftOffset); j < output.sizeAt(dim); ++j) { // fill right side + subArr.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + output.sizeAt(dim) + leftOffset - 1 - j]); + subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + j]); + subArrOut.assign(&subArr); + } + break; + + case 2: // SYMMETRIC mode + for(int j = 1; j <= leftOffset; ++j) { // fill left side + subArr.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset + j - 1]); + subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset - j]); + subArrOut.assign(&subArr); + } + for(int j = (output.sizeAt(dim) - leftOffset); j < output.sizeAt(dim); ++j) { // fill right side + subArr.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + output.sizeAt(dim) + leftOffset - j]); + subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + j]); + subArrOut.assign(&subArr); + } + break; + } +} + */ +/* + void recursiveLoopForPad(const int mode, NDArray& input, const NDArray& paddings, NDArray& output, std::vector dimensions, int dim, int inIdx, int outIdx, NDArray& padValue ) { + BUILD_SINGLE_SELECTOR(input.dataType(), recursiveLoopForPad_, (mode, input, paddings, output, dimensions, dim, inIdx, outIdx, padValue), LIBND4J_TYPES); + } + + BUILD_SINGLE_TEMPLATE(template void recursiveLoopForPad_, (const int mode, NDArray& input, const NDArray& paddings, NDArray& output, std::vector dimensions, int dim, int inIdx, int outIdx, NDArray& padValue), LIBND4J_TYPES); + +*/ + +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/randomShuffle.cpp b/libnd4j/include/ops/declarable/helpers/cpu/randomShuffle.cpp new file mode 100644 index 000000000..7323c3937 --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/randomShuffle.cpp @@ -0,0 +1,126 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + + + +#include +#include +#include +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + +////////////////////////////////////////////////////////////////////////// +template +void randomShuffle_(NDArray& input, NDArray& output, sd::graph::RandomGenerator& rng, const bool isInplace) { + + // check edge cases first + int temp; + const int firstDim = input.sizeAt(0); + if(input.lengthOf() == 1 || firstDim == 1) { + + if(!isInplace) + output.assign(input); + } + else if (input.isVector() || shape::isLikeVector(input.getShapeInfo(), temp)) { + + // apply Fisher-Yates shuffle + if(isInplace) { + //PRAGMA_OMP_PARALLEL_FOR_IF((firstDim-1) > Environment::getInstance()->tadThreshold()) + for(int i = firstDim-1; i > 0; --i) { + int r = rng.relativeInt(i) % i; + if(i == r) + continue; + T t0 = input.t(i); + T t1 = input.t(r); + //math::nd4j_swap(input(i), input(r)); + input.t(i) = t1; + input.t(r) = t0; + } + } + else { + std::vector indices(firstDim); + std::iota(indices.begin(), indices.end(), 0); + output.p(Nd4jLong(0), input.e(0)); + + // FIXME: parallelism!! + for(int i = firstDim-1; i > 0; --i) { + int r = rng.relativeInt(i) % i; + output.t(i) = input.t(indices[r]); + if(i == r) + continue; + + output.t(r) = input.t(indices[i]); + math::nd4j_swap(indices[i], indices[r]); + } + rng.rewindH(firstDim-1); + } + } + else { + + // evaluate sub-arrays list of input array through all dimensions excluding first one + std::vector dimensions = ShapeUtils::evalDimsToExclude(input.rankOf(), {0}); + auto subArrsListIn = input.allTensorsAlongDimension(dimensions); + + // apply Fisher-Yates shuffle + if(isInplace) { + //PRAGMA_OMP_PARALLEL_FOR_IF((firstDim-1) > Environment::getInstance()->elementwiseThreshold()) + for(int i = firstDim - 1; i > 0; --i) { + int r = rng.relativeInt(i) % i; + + if(i == r) + continue; + subArrsListIn.at(i)->swapUnsafe(*subArrsListIn.at(r)); + } + } + else { + // evaluate sub-arrays list of output array through all dimensions excluding first one + auto subArrsListOut = output.allTensorsAlongDimension(dimensions); + std::vector indices(firstDim); + std::iota(indices.begin(), indices.end(), 0); + bool isZeroShuffled = false; + //PRAGMA_OMP_PARALLEL_FOR_IF((firstDim-1) > Environment::getInstance()->tadThreshold()) + for(int i = firstDim - 1; i > 0; --i) { + int r = rng.relativeInt(i) % i; + subArrsListOut.at(i)->assign(subArrsListIn.at(indices[r])); + if(r == 0) + isZeroShuffled = true; + if(i == r) + continue; + subArrsListOut.at(r)->assign(subArrsListIn.at(indices[i])); + math::nd4j_swap(indices[i], indices[r]); + } + if(!isZeroShuffled) + subArrsListOut.at(0)->assign(subArrsListIn.at(0)); + } + rng.rewindH(firstDim-1); + } + +} + + void randomShuffle(sd::LaunchContext * context, NDArray& input, NDArray& output, sd::graph::RandomGenerator& rng, const bool isInplace) { + BUILD_SINGLE_SELECTOR(input.dataType(), randomShuffle_, (input, output, rng, isInplace), LIBND4J_TYPES); + } +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/scatterUpdateAndSimple.cpp b/libnd4j/include/ops/declarable/helpers/cpu/scatterUpdateAndSimple.cpp new file mode 100644 index 000000000..fe41c5d43 --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/scatterUpdateAndSimple.cpp @@ -0,0 +1,115 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + +#include +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + +////////////////////////////////////////////////////////////////////////// +void scatterUpdate(sd::LaunchContext * context, NDArray& input, NDArray& updates, const std::vector* intArgs) { + + int opCode = (*intArgs)[0]; + int dimSize = (*intArgs)[1]; + Nd4jLong e; + Nd4jLong limg = 2 + dimSize; + std::vector tadDimensions(dimSize); + for (e = 2; e < limg; e++) + tadDimensions[e-2] = (*intArgs)[e]; + + std::vector dimsToExclude = ShapeUtils::evalDimsToExclude(input.rankOf(), tadDimensions); + + // increasing counter to skip numIndices + e++; + std::vector indices; + for (; e < static_cast(intArgs->size()); e++) + indices.push_back((*intArgs)[e]); + + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) { + auto inSubArr = input(indices[i], dimsToExclude, true); + auto updSubArr = updates(i, dimsToExclude, true); + + if (inSubArr.lengthOf() != updSubArr.lengthOf()) + continue; + + switch (opCode) { + case 0: + inSubArr.applyPairwiseTransform(pairwise::Add, updSubArr, inSubArr); + break; + case 1: + inSubArr.applyPairwiseTransform(pairwise::Subtract, updSubArr, inSubArr); + break; + case 2: + inSubArr.applyPairwiseTransform(pairwise::Multiply, updSubArr, inSubArr); + break; + case 3: + inSubArr.applyPairwiseTransform(pairwise::Divide, updSubArr, inSubArr); + break; + case 4: + inSubArr.applyPairwiseTransform(pairwise::ReverseSubtract, updSubArr, inSubArr); + break; + case 5: + inSubArr.applyPairwiseTransform(pairwise::ReverseDivide, updSubArr, inSubArr); + break; + case 6: + inSubArr.applyPairwiseTransform(pairwise::CopyPws, updSubArr, inSubArr); + break; + default: + continue; + } + } + }; + + samediff::Threads::parallel_tad(func, 0, indices.size()); +} + + +////////////////////////////////////////////////////////////////////////// +void scatterSimple(sd::LaunchContext * context, const int opId, NDArray& input, const NDArray& updates, const NDArray& indices, const std::vector& dimensions) { + + // updates and indices have same length + const Nd4jLong len = indices.lengthOf(); + + switch (opId) { + + case 6: { // copy + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) { + auto inSubArr = input(i, dimensions); + inSubArr.p(indices.t(i), updates.e(i)); + } + }; + + samediff::Threads::parallel_for(func, 0, len); + } + break; + + default: + throw std::invalid_argument("helpers::scatterSimple: operation is not implemented for given id !"); + } +} + +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/tile.cpp b/libnd4j/include/ops/declarable/helpers/cpu/tile.cpp new file mode 100644 index 000000000..8f2a10bc9 --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/tile.cpp @@ -0,0 +1,91 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + + +#include +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + +////////////////////////////////////////////////////////////////////////// +template +static void tileBP_(const NDArray& gradO /*input*/, NDArray& gradI /*output*/, const std::vector reps) { + + T* gradIBuff = reinterpret_cast(gradI.getBuffer()); + const T* gradOBuff = reinterpret_cast(gradO.getBuffer()); + const Nd4jLong gradILen = gradI.lengthOf(); + const Nd4jLong gradOLen = gradO.lengthOf(); // gradOLen >= gradILen + const Nd4jLong gradIEWS = sd::math::nd4j_abs(gradI.ews()); + const Nd4jLong gradOEWS = gradO.ews(); + + // initial zeroing of gradI content + if(gradIEWS == 1) + memset(gradIBuff, 0, gradILen * sizeof(T)); + else { + //PRAGMA_OMP_PARALLEL_FOR_SIMD + for (Nd4jLong i = 0; i < gradILen * gradIEWS; i += gradIEWS) + gradIBuff[i] = static_cast(0.f); + } + + + if(gradO.ordering() == 'c' && gradOEWS == 1) { + + //PRAGMA_OMP_PARALLEL_FOR_SIMD + for(Nd4jLong i=0; i(idx) + gradOBuff[i]); + } + } + else if(gradO.ordering() == 'c' && gradOEWS > 1) { + + //PRAGMA_OMP_PARALLEL_FOR_SIMD + for(Nd4jLong i=0; i(idx) + gradOBuff[i * gradOEWS]); + } + } + else { + + //PRAGMA_OMP_PARALLEL_FOR_SIMD + for(Nd4jLong i=0; i(fidx) + gradOBuff[shape::getIndexOffset(i, gradO.getShapeInfo())]); + } + } +} + +void tileBP(sd::LaunchContext * context, const NDArray& gradO /*input*/, NDArray& gradI /*output*/, const std::vector reps) { + BUILD_SINGLE_SELECTOR(gradI.dataType(), tileBP_, (gradO, gradI, reps), FLOAT_TYPES); +} + + +BUILD_SINGLE_TEMPLATE(template void tileBP_, (const NDArray& gradO /*input*/, NDArray& gradI /*output*/, const std::vector reps), FLOAT_TYPES); + + + + + +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/trace.cpp b/libnd4j/include/ops/declarable/helpers/cpu/trace.cpp new file mode 100644 index 000000000..d544fa24e --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/trace.cpp @@ -0,0 +1,47 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + + +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + +////////////////////////////////////////////////////////////////////////// +template +static void trace_(const NDArray& input, NDArray& output) { + const int inRank = input.rankOf(); + auto setOfSubArrs = input.allTensorsAlongDimension({inRank-2, inRank-1}); + + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) + output.p(i, setOfSubArrs.at(i)->getTrace()); + }; + samediff::Threads::parallel_for(func, 0, setOfSubArrs.size()); +} + + void trace(sd::LaunchContext * context, const NDArray& input, NDArray& output) { + BUILD_SINGLE_SELECTOR(input.dataType(), trace_, (input, output), LIBND4J_TYPES); + } +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/transforms.cpp b/libnd4j/include/ops/declarable/helpers/cpu/transforms.cpp deleted file mode 100644 index 7169cca4a..000000000 --- a/libnd4j/include/ops/declarable/helpers/cpu/transforms.cpp +++ /dev/null @@ -1,1305 +0,0 @@ -/******************************************************************************* - * Copyright (c) 2015-2018 Skymind, Inc. - * - * This program and the accompanying materials are made available under the - * terms of the Apache License, Version 2.0 which is available at - * https://www.apache.org/licenses/LICENSE-2.0. - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT - * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the - * License for the specific language governing permissions and limitations - * under the License. - * - * SPDX-License-Identifier: Apache-2.0 - ******************************************************************************/ - -// -// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 -// - - -#include -#include -#include -#include -#include -#include -#include -#include -#include - -namespace sd { -namespace ops { -namespace helpers { - - -////////////////////////////////////////////////////////////////////////// -template -static void triuBP_(sd::LaunchContext * context, const NDArray& input, const NDArray& gradO, NDArray& gradI, const int diagonal) { - - auto dOdI = NDArray(&gradO); // dO/dI - const_cast(input).fillAsTriangular(0, diagonal, dOdI.sizeAt(-1), dOdI, 'b'); - int dLen = dOdI.lengthOf(); - - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) { - if (dOdI.t(i) != static_cast(0.f)) - dOdI.t(i) = static_cast(1.f); - } - }; - samediff::Threads::parallel_for(func, 0, dLen); - - // FIXME: !!! - gradI.assign(dOdI * gradO); // chain rule: dLoss/dI = dO/dI * dLoss/dO -} - - void triuBP(sd::LaunchContext * context, const NDArray& input, const NDArray& gradO, NDArray& gradI, const int diagonal) { - BUILD_SINGLE_SELECTOR(gradO.dataType(), triuBP_, (context, input, gradO, gradI, diagonal), LIBND4J_TYPES); - } - -////////////////////////////////////////////////////////////////////////// -template -static void trace_(const NDArray& input, NDArray& output) { - const int inRank = input.rankOf(); - auto setOfSubArrs = input.allTensorsAlongDimension({inRank-2, inRank-1}); - - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) - output.p(i, setOfSubArrs.at(i)->getTrace()); - }; - samediff::Threads::parallel_for(func, 0, setOfSubArrs.size()); -} - - void trace(sd::LaunchContext * context, const NDArray& input, NDArray& output) { - BUILD_SINGLE_SELECTOR(input.dataType(), trace_, (input, output), LIBND4J_TYPES); - } - -////////////////////////////////////////////////////////////////////////// -template -void randomShuffle_(NDArray& input, NDArray& output, sd::graph::RandomGenerator& rng, const bool isInplace) { - - // check edge cases first - int temp; - const int firstDim = input.sizeAt(0); - if(input.lengthOf() == 1 || firstDim == 1) { - - if(!isInplace) - output.assign(input); - } - else if (input.isVector() || shape::isLikeVector(input.getShapeInfo(), temp)) { - - // apply Fisher-Yates shuffle - if(isInplace) { - //PRAGMA_OMP_PARALLEL_FOR_IF((firstDim-1) > Environment::getInstance()->tadThreshold()) - for(int i = firstDim-1; i > 0; --i) { - int r = rng.relativeInt(i) % i; - if(i == r) - continue; - T t0 = input.t(i); - T t1 = input.t(r); - //math::nd4j_swap(input(i), input(r)); - input.t(i) = t1; - input.t(r) = t0; - } - } - else { - std::vector indices(firstDim); - std::iota(indices.begin(), indices.end(), 0); - output.p(Nd4jLong(0), input.e(0)); - - // FIXME: parallelism!! - for(int i = firstDim-1; i > 0; --i) { - int r = rng.relativeInt(i) % i; - output.t(i) = input.t(indices[r]); - if(i == r) - continue; - - output.t(r) = input.t(indices[i]); - math::nd4j_swap(indices[i], indices[r]); - } - rng.rewindH(firstDim-1); - } - } - else { - - // evaluate sub-arrays list of input array through all dimensions excluding first one - std::vector dimensions = ShapeUtils::evalDimsToExclude(input.rankOf(), {0}); - auto subArrsListIn = input.allTensorsAlongDimension(dimensions); - - // apply Fisher-Yates shuffle - if(isInplace) { - //PRAGMA_OMP_PARALLEL_FOR_IF((firstDim-1) > Environment::getInstance()->elementwiseThreshold()) - for(int i = firstDim - 1; i > 0; --i) { - int r = rng.relativeInt(i) % i; - - if(i == r) - continue; - subArrsListIn.at(i)->swapUnsafe(*subArrsListIn.at(r)); - } - } - else { - // evaluate sub-arrays list of output array through all dimensions excluding first one - auto subArrsListOut = output.allTensorsAlongDimension(dimensions); - std::vector indices(firstDim); - std::iota(indices.begin(), indices.end(), 0); - bool isZeroShuffled = false; - //PRAGMA_OMP_PARALLEL_FOR_IF((firstDim-1) > Environment::getInstance()->tadThreshold()) - for(int i = firstDim - 1; i > 0; --i) { - int r = rng.relativeInt(i) % i; - subArrsListOut.at(i)->assign(subArrsListIn.at(indices[r])); - if(r == 0) - isZeroShuffled = true; - if(i == r) - continue; - subArrsListOut.at(r)->assign(subArrsListIn.at(indices[i])); - math::nd4j_swap(indices[i], indices[r]); - } - if(!isZeroShuffled) - subArrsListOut.at(0)->assign(subArrsListIn.at(0)); - } - rng.rewindH(firstDim-1); - } - -} - - void randomShuffle(sd::LaunchContext * context, NDArray& input, NDArray& output, sd::graph::RandomGenerator& rng, const bool isInplace) { - BUILD_SINGLE_SELECTOR(input.dataType(), randomShuffle_, (input, output, rng, isInplace), LIBND4J_TYPES); - } - - -////////////////////////////////////////////////////////////////////////// -template -void pad_(const int mode, const NDArray& input, const NDArray& paddings, NDArray& output, const NDArray& padValue) { - - const T* x = input.bufferAsT(); - T* z = output.bufferAsT(); - - const Nd4jLong* xShape = input.shapeOf(); - const Nd4jLong* zShape = output.shapeOf(); - - const int rank = input.rankOf(); // both input and output have the same rank - const int rankMinusOne = rank - 1; - - const auto zLen = output.lengthOf(); - - if(mode == 0) { // CONSTANT case - - const T padVal = padValue.e(0); - - auto func = PRAGMA_THREADS_FOR { - - int zCoords[MAX_RANK], xCoords[MAX_RANK]; - - for (auto i = start; i < stop; i++) { - - shape::index2coordsCPU(start, i, output.getShapeInfo(), zCoords); - const auto zOffset = shape::getOffset(output.getShapeInfo(), zCoords); - - memcpy(xCoords, zCoords, rank * sizeof(int)); - - bool within = true; - - for (int j = rankMinusOne; j >= 0; --j) { - - if (xShape[j] == zShape[j]) - continue; - - const auto left = paddings.e(j, 0); - - if (zCoords[j] < left || zCoords[j] >= left + xShape[j]) { - within = false; - break; - } - else - xCoords[j] = zCoords[j] - left; - } - - if (within) - z[zOffset] = x[shape::getOffset(input.getShapeInfo(), xCoords)]; - else - z[zOffset] = padVal; - } - }; - - samediff::Threads::parallel_tad(func, 0, zLen); - } - else { // REFLECT and SYMMETRIC cases - - const Nd4jLong shift1 = mode == 1 ? 0 : 1; // REFLECT : SYMMETRIC - const Nd4jLong shift2 = mode == 1 ? 2 : 1; // REFLECT : SYMMETRIC - - auto func = PRAGMA_THREADS_FOR { - - int zCoords[MAX_RANK], xCoords[MAX_RANK]; - - for (auto i = start; i < stop; i++) { - - shape::index2coordsCPU(start, i, output.getShapeInfo(), zCoords); - const auto zOffset = shape::getOffset(output.getShapeInfo(), zCoords); - - memcpy(xCoords, zCoords, rank * sizeof(int)); - - for (int j = rankMinusOne; j >= 0; --j) { - - if (xShape[j] == zShape[j]) - continue; - - xCoords[j] = zCoords[j] - paddings.e(j, 0); // are ready to fill middle (within input dimension range) - - if (xCoords[j] < 0) - xCoords[j] = -xCoords[j] - shift1; // means fill from left - else if (xCoords[j] >= xShape[j]) - xCoords[j] = 2 * xShape[j] - xCoords[j] - shift2; // means fill from right - } - - const auto xOffset = shape::getOffset(input.getShapeInfo(), xCoords); - z[zOffset] = x[xOffset]; - } - }; - - samediff::Threads::parallel_tad(func, 0, zLen); - } -} - -// ////////////////////////////////////////////////////////////////////////// -// template -// void pad2_(const int mode, const NDArray& input, const NDArray& paddings, NDArray& output, NDArray const& padValue) { - -// const int rank = output.rankOf(); -// std::vector dimsToExclude(rank); -// std::iota(dimsToExclude.begin(), dimsToExclude.end(), 0); // fill with 0, 1, ... rank-1 - -// Nd4jLong numLeft = paddings.e(rank-1,0); -// Nd4jLong numRight = paddings.e(rank-1,1); -// Nd4jLong inDimSize = input.sizeAt(rank-1); -// Nd4jLong outDimSize = output.sizeAt(rank-1); - -// std::vector> outIdx = { std::vector(2*rank), {numLeft, numLeft + inDimSize}, {0, numLeft}, {numLeft + inDimSize, outDimSize} }; - -// for(int i = 0; i < rank-1; ++i) { -// outIdx[0][2*i] = paddings.e(i, 0); -// outIdx[0][2*i + 1] = outIdx[0][2*i] + input.sizeAt(i); -// } -// outIdx[0][2*rank-1] = outIdx[0][2*rank-2] = 0; - -// // ***** populate innermost sub-arrays firstly ***** // -// dimsToExclude.pop_back(); - -// Nd4jLong startL = mode == 1 ? 1 : 0; // REFLECT or SYMMETRIC -// Nd4jLong startR = mode == 1 ? inDimSize-2 : inDimSize-1; // REFLECT or SYMMETRIC - -// Nd4jLong numOfSubArrs = ShapeUtils::getNumOfSubArrs(input.getShapeInfo(), dimsToExclude); - -// NDArray outSubArr0 = output(outIdx[0], true); - -// PRAGMA_OMP_PARALLEL_FOR -// for(Nd4jLong j = 0; j < numOfSubArrs; ++j) { - -// NDArray outSubArr1 = outSubArr0(j, dimsToExclude); -// NDArray inSubArr = input(j, dimsToExclude); -// NDArray outSubArrMid = outSubArr1(outIdx[1]); - -// outSubArrMid.assign(inSubArr); // assign middle - -// if(mode == 0) { // CONSTANT -// if(numLeft != 0) { -// NDArray temp = outSubArr1(outIdx[2]); -// temp.assign(padValue); // assign left -// } -// if(numRight != 0) { -// NDArray temp = outSubArr1(outIdx[3]); -// temp.assign(padValue); // assign right -// } -// } -// else { // REFLECT or SYMMETRIC - -// for(Nd4jLong k = numLeft-1, e = startL; k >= 0; --k, ++e) // fill left side -// outSubArr1.t(k) = inSubArr.t(e); - -// for(Nd4jLong k = numLeft + inDimSize, e = startR; k < outDimSize; ++k, --e) // fill right side -// outSubArr1.t(k) = inSubArr.t(e); -// } -// } - -// // ***** fill rest of outer sub-arrays ***** // -// std::vector outIdxInner(2, 0); -// std::vector outIdxOuter(2, 0); - -// for(int i = rankBorder - 1; i >= 0; --i) { - -// dimsToExclude.pop_back(); - -// outIdxInner.push_back(0), outIdxInner.push_back(0); -// outIdxOuter.push_back(0), outIdxOuter.push_back(0); - -// Nd4jLong numLeft = paddings.e(i, 0); -// Nd4jLong numRight = paddings.e(i, 1); - -// if(numLeft == 0 && numRight == 0) -// continue; - -// Nd4jLong inDimSize = input.sizeAt(i); -// Nd4jLong outDimSize = output.sizeAt(i); - -// if(mode == 0) { -// outIdxOuter[0] = 0; outIdxOuter[1] = numLeft; -// outIdxInner[0] = numLeft + inDimSize; outIdxInner[1] = outDimSize; -// } - -// startL = mode == 1 ? numLeft + 1 : numLeft; // REFLECT or SYMMETRIC -// startR = mode == 1 ? numLeft + inDimSize - 2 : numLeft + inDimSize-1; // REFLECT or SYMMETRIC - -// numOfSubArrs = ShapeUtils::getNumOfSubArrs(output.getShapeInfo(), dimsToExclude); - -// PRAGMA_OMP_PARALLEL_FOR_ARGS(firstprivate(outIdxOuter, outIdxInner)) -// for(Nd4jLong j = 0; j < numOfSubArrs; ++j) { - -// NDArray outSubArr = output(j, dimsToExclude); - -// if(mode == 0) { // CONSTANT - -// if(numLeft != 0) { -// NDArray tempO = outSubArr(outIdxOuter); -// tempO.assign(padValue); // assign left -// } - -// if(numRight != 0) { -// NDArray tempI = outSubArr(outIdxInner); -// tempI.assign(padValue); // assign right -// } -// } -// else { // REFLECT or SYMMETRIC - -// for(Nd4jLong k = numLeft-1, e = startL; k >= 0; --k, ++e) { // fill left side -// outIdxOuter[0] = k; -// outIdxOuter[1] = k+1; -// outIdxInner[0] = e; -// outIdxInner[1] = e+1; -// NDArray outSubArrInner = outSubArr(outIdxInner); -// NDArray outSubArrOuter = outSubArr(outIdxOuter); -// outSubArrOuter.assign(outSubArrInner); -// } - -// for(Nd4jLong k = numLeft + inDimSize, e = startR; k < outDimSize; ++k, --e) { // fill right side -// outIdxOuter[0] = k; -// outIdxOuter[1] = k+1; -// outIdxInner[0] = e; -// outIdxInner[1] = e+1; -// NDArray outSubArrInner = outSubArr(outIdxInner); -// NDArray outSubArrOuter = outSubArr(outIdxOuter); -// outSubArrOuter.assign(outSubArrInner); -// } -// } -// } -// } -// } - -void pad(sd::LaunchContext * context, const int mode, const NDArray& input, const NDArray& paddings, NDArray& output, NDArray const& padValue) { - BUILD_SINGLE_SELECTOR(input.dataType(), pad_, (mode, input, paddings, output, padValue), LIBND4J_TYPES); -} - -//////////////////////////////////////////////////////////////////////// -/*// initial values of inIdx, outIdx, dim must be equal to zero -template -static void recursiveLoopForPad_(const int mode, NDArray& input, const NDArray& paddings, NDArray& output, std::vector dimensions, int dim, int inIdx, int outIdx, NDArray& padValue ) { - - int leftOffset; - // dimensions are array of input dimensions, it is sorted in increasing order - // every time at the beginning we erase first element from it (not good idea to use vector for this purpose, but luckily it is small enough) - // then we use this array for tads building, every time while recursion the number of built tads becomes bigger - dimensions.erase(dimensions.begin()); - // build tad basing on output array, also create auxiliary arrays pointing on required output array ranges - shape::TAD tadOut(output.getShapeInfo(), dimensions.data(), dimensions.size()); - tadOut.createTadOnlyShapeInfo(); - tadOut.createOffsets(); - auto subArrOut = NDArray(output.getBuffer(), tadOut.tadOnlyShapeInfo, output.getContext()); - auto subArr = NDArray(output.getBuffer(), tadOut.tadOnlyShapeInfo, output.getContext()); - // build tad basing on input array, also create auxiliary array pointing on required input array range - shape::TAD tadIn(input.getShapeInfo(), dimensions.data(), dimensions.size()); - tadIn.createTadOnlyShapeInfo(); - tadIn.createOffsets(); - auto subArrIn = NDArray(input.getBuffer(), tadIn.tadOnlyShapeInfo, output.getContext()); - // these indices take into account recursion and always point to actual tads numbers - if (input.rankOf() > 1 && output.rankOf() > 1) {// only for non-vector cases - outIdx = outIdx * output.sizeAt(dim + 1); - inIdx = inIdx * input.sizeAt(dim + 1); - } - // current input tad number, we add to it unity in a loop - int k = -1; - // loop through current dimension - for(int i = 0; i < output.sizeAt(dim); ++i) { - // corresponds to outer range (relevant indices are absent in input) - leftOffset = paddings.e(dim, 0); - if(i < leftOffset || i >= (input.sizeAt(dim) + leftOffset)) - continue; - - // increase input tads number - ++k; - // recursion condition allows for the fact that tad can't reduce to scalar - if(dim < input.rankOf() - 2) - recursiveLoopForPad(mode, input, paddings, output, dimensions, dim + 1, inIdx + k, outIdx + i, padValue); - else if (paddings.sizeAt(0) > dim + 1){ - leftOffset = paddings.e(dim + 1, 0); - // shift buffers pointers to actual element position - if (output.rankOf() > 1) { - subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + i]); - subArrIn.setBuffer(reinterpret_cast(input.getBuffer()) + tadIn.tadOffsets[inIdx + i - paddings.e(dim, 0)]); - } - else { - subArrOut.p(i, subArrIn.e(i - leftOffset)); - } - // most inner loop, corresponds to last dim = rank-1 - switch (mode) { - case 0: // CONSTANT mode - for(int j = 0; j < subArrOut.lengthOf(); ++j) - if(j < leftOffset || j >= (subArrIn.lengthOf() + leftOffset) ) // firstly fill with zeros outer ranges - subArrOut.p(j, (T)0.f); - else - subArrOut.p(j, subArrIn.e(j - leftOffset)); // fill middle with elements of input array - break; - - case 1: // REFLECT mode - for(int j = 1; j <= leftOffset; ++j) // fill firstly left side - subArrOut.p(leftOffset - j, subArrIn.e(j)); - for(int j = 0; j < subArrIn.lengthOf(); ++j) // fill middle - subArrOut.p(leftOffset + j, subArrIn.e(j)); - for(int j = (subArrOut.lengthOf() - leftOffset); j < subArrOut.lengthOf(); ++j) // fill right side - subArrOut.p(j, subArrIn.e(subArrOut.lengthOf() - j - 1)); - break; - - case 2: // SYMMETRIC mode - for(int j = 1; j <= leftOffset; ++j) // fill firstly left side - subArrOut.p(leftOffset - j, subArrIn.e(j-1)); - for(int j = 0; j < subArrIn.lengthOf(); ++j) // fill middle - subArrOut.p(leftOffset + j, subArrIn.e(j)); - for(int j = (subArrOut.lengthOf() - leftOffset); j < subArrOut.lengthOf(); ++j) // fill right side - subArrOut.p(j, subArrIn.e(subArrOut.lengthOf() - j)); - break; - } - } - else { - - if (mode == 0 && input.rankOf() < 2) - subArrOut.p(i, subArrIn.e(i - leftOffset)); // fill middle with elements of input array - } - } - // populate sub-array formed previously - leftOffset = paddings.e(dim,0); - switch (mode) { - case 0: // CONSTANT mode - for(int j = 1; j <= leftOffset; ++j) { - // fill left side with padValue - if (output.rankOf() > 1) { - subArrOut.setBuffer( - reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset - j]); - subArrOut.assign(padValue); - } - else { - subArrOut.p(j - 1, padValue); - } - } -// output.printIndexedBuffer("Output at"); - for(int j = (output.sizeAt(dim) - leftOffset); j < output.sizeAt(dim); ++j) { // fill left side with zeros - if (output.rankOf() > 1) { - subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + j]); - subArrOut.assign(padValue); - } - else { - subArrOut.p(j, padValue); - } - } - break; - - case 1: // REFLECT mode - for(int j = 1; j <= leftOffset; ++j) { // fill left side - subArr.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset + j]); - subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset - j]); - subArrOut.assign(&subArr); - } - for(int j = (output.sizeAt(dim) - leftOffset); j < output.sizeAt(dim); ++j) { // fill right side - subArr.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + output.sizeAt(dim) + leftOffset - 1 - j]); - subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + j]); - subArrOut.assign(&subArr); - } - break; - - case 2: // SYMMETRIC mode - for(int j = 1; j <= leftOffset; ++j) { // fill left side - subArr.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset + j - 1]); - subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + leftOffset - j]); - subArrOut.assign(&subArr); - } - for(int j = (output.sizeAt(dim) - leftOffset); j < output.sizeAt(dim); ++j) { // fill right side - subArr.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + output.sizeAt(dim) + leftOffset - j]); - subArrOut.setBuffer(reinterpret_cast(output.getBuffer()) + tadOut.tadOffsets[outIdx + j]); - subArrOut.assign(&subArr); - } - break; - } -} - */ -/* - void recursiveLoopForPad(const int mode, NDArray& input, const NDArray& paddings, NDArray& output, std::vector dimensions, int dim, int inIdx, int outIdx, NDArray& padValue ) { - BUILD_SINGLE_SELECTOR(input.dataType(), recursiveLoopForPad_, (mode, input, paddings, output, dimensions, dim, inIdx, outIdx, padValue), LIBND4J_TYPES); - } - - BUILD_SINGLE_TEMPLATE(template void recursiveLoopForPad_, (const int mode, NDArray& input, const NDArray& paddings, NDArray& output, std::vector dimensions, int dim, int inIdx, int outIdx, NDArray& padValue), LIBND4J_TYPES); - -*/ - -//////////////////////////////////////////////////////////////////////// -void invertPermutation(sd::LaunchContext * context, const NDArray& input, NDArray& output) { - - std::set uniqueElems; - const int length = input.lengthOf(); - - for(int i = 0; i < length; ++i) { - - int elem = input.e(i); - - if(!uniqueElems.insert(elem).second) // this operation forbids us to use #pragma omp - throw std::runtime_error("helpers::invertPermutation function: input array contains duplicates !"); - - if(elem < 0 || elem > length - 1) - throw std::runtime_error("helpers::invertPermutation function: element of input array is out of range (0, length-1) !"); - - output.p(elem, i); - } -} - -//////////////////////////////////////////////////////////////////////// -template -static void gatherND_(NDArray& input, NDArray& indices, NDArray& output) { - - const X* x = reinterpret_cast(input.getBuffer()); - const Y* y = reinterpret_cast(indices.getBuffer()); - X* z = reinterpret_cast(output.getBuffer()); - - const int xRank = input.rankOf(); - const int yRank = indices.rankOf(); - const int zRank = output.rankOf(); - const int maxRank = sd::math::nd4j_max(yRank, sd::math::nd4j_max(xRank, zRank)); - - const Nd4jLong zLen = output.lengthOf(); - - const uint yLastDim = indices.sizeAt(-1); - - const int diff = zRank - xRank; - const bool bEqual = yLastDim == xRank; - - auto func = PRAGMA_THREADS_FOR { - - int xCoords[MAX_RANK], zCoords[MAX_RANK], temp; - - for (auto i = start; i < stop; i++) { - - shape::index2coordsCPU(start, i, output.getShapeInfo(), zCoords); - - const auto zOffset = shape::getOffset(output.getShapeInfo(), zCoords); - - temp = zCoords[yRank - 1]; - zCoords[yRank - 1] = 0; - const auto yOffset = shape::getOffset(indices.getShapeInfo(), zCoords); - zCoords[yRank - 1] = temp; - - if(bEqual) - memcpy(xCoords, zCoords, zRank * sizeof(int)); - else if(diff >= 0) - memcpy(xCoords, zCoords + diff, xRank * sizeof(int)); - else - memcpy(xCoords - diff, zCoords, zRank * sizeof(int)); - - for (uint j = 0; j < yLastDim; ++j) - xCoords[j] = y[yOffset + j * indices.stridesOf()[yRank - 1]]; // last stride - - const auto xOffset = shape::getOffset(input.getShapeInfo(), xCoords); - - z[zOffset] = x[xOffset]; - } - }; - - samediff::Threads::parallel_tad(func, 0, zLen); -} - -//////////////////////////////////////////////////////////////////////// -void gatherND(sd::LaunchContext * context, NDArray& input, NDArray& indices, NDArray& output) { - BUILD_DOUBLE_SELECTOR(input.dataType(), indices.dataType(), gatherND_, (input, indices, output), LIBND4J_TYPES, INDEXING_TYPES); -} - - -//////////////////////////////////////////////////////////////////////// -template -static void gather_(NDArray* input, const NDArray* indices, NDArray* output, const std::vector& intArgs) { - - int axis = intArgs.size() > 0 ? intArgs[0] : 0; - const int inputRank = input->rankOf(); - if(axis < 0) - axis += inputRank; - - const int numOfIntArgs = intArgs.size(); - - if (indices != nullptr) { - - for(Nd4jLong i = 0; i < indices->lengthOf(); ++i) - if(indices->e(i) >= input->sizeAt(axis)) - throw std::runtime_error("helpers::gather function: indices array contains wrong elements, each element must be smaller than corresponding dimension of input array !"); - - // first case: indices consist of only one scalar - if(indices->isScalar()) { - if(input->rankOf() <= 1){ - //For scalar indices, rank 0 or 1 input: can't do tensor along dimension 0 as this is whole array... instead, we want to get a scalar - auto idx = indices->e(0); - auto scalarNDArray = input->e(idx); - output->assign(scalarNDArray); - } else { - auto dimensions = ShapeUtils::evalDimsToExclude(input->rankOf(), {axis}); - auto tadPack = sd::ConstantTadHelper::getInstance()->tadForDimensions(input->getShapeInfo(), dimensions); - - auto tadArr = NDArray(reinterpret_cast(reinterpret_cast(input->getBuffer()) + tadPack.primaryOffsets()[indices->e(0)]), tadPack.primaryShapeInfo(), output->getContext()); - output->assign(&tadArr); - } - } - else if (input->rankOf() == 1 && indices->isVector()) { - // special case - auto func = PRAGMA_THREADS_FOR { - for (auto e = start; e < stop; e++) - output->p(e, input->e(indices->e(e))); - }; - - samediff::Threads::parallel_for(func, 0, indices->lengthOf()); - } - else { - - std::vector dimsOut(indices->rankOf()); - std::iota(dimsOut.begin(), dimsOut.end(), axis); // fill with axis, axis+1, ... indices->rankOf()-1 - const Nd4jLong numOfSubArrs = ShapeUtils::getNumOfSubArrs(output->getShapeInfo(), dimsOut); - - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) { - NDArray subArrOut = (*output)(i, dimsOut); - NDArray subArrIn = (*input)(indices->e(i), {axis}); - subArrOut.assign(subArrIn); - } - }; - - samediff::Threads::parallel_tad(func, 0, numOfSubArrs); - } - } - else { - - for(int i = 1; i < numOfIntArgs; ++i) - if(intArgs[i] >= input->sizeAt(axis)) - throw std::runtime_error("helpers::gather function: some of input indexes is larger than corresponding shape of input array !"); - - // we only allow scalar/vector case here - if (numOfIntArgs == 2) { // scalar case - output->assign((*input)(intArgs[1], {axis})); - } - else { // vector case - const Nd4jLong numOfSubArrs = ShapeUtils::getNumOfSubArrs(output->getShapeInfo(), {axis}); - - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) { - NDArray subArrOut = (*output)(i, {axis}); - NDArray subArrIn = (*input)(intArgs[i + 1], {axis}); - subArrOut.assign(subArrIn); - } - }; - - samediff::Threads::parallel_tad(func, 0, numOfSubArrs); - } - } -} - - void gather(NDArray* input, const NDArray* indices, NDArray* output, const std::vector& intArgs) { - BUILD_SINGLE_SELECTOR(input->dataType(), gather_, (input, indices, output, intArgs), LIBND4J_TYPES); - } - -////////////////////////////////////////////////////////////////////////// -void eye(sd::LaunchContext * context, NDArray& output) { - - const int rank = output.rankOf(); - auto arrs = output.allTensorsAlongDimension({rank-2, rank-1}); - - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) - arrs.at(i)->setIdentity(); - }; - - samediff::Threads::parallel_tad(func, 0, arrs.size()); -} - -////////////////////////////////////////////////////////////////////////// -void scatterUpdate(sd::LaunchContext * context, NDArray& input, NDArray& updates, const std::vector* intArgs) { - - int opCode = (*intArgs)[0]; - int dimSize = (*intArgs)[1]; - Nd4jLong e; - Nd4jLong limg = 2 + dimSize; - std::vector tadDimensions(dimSize); - for (e = 2; e < limg; e++) - tadDimensions[e-2] = (*intArgs)[e]; - - std::vector dimsToExclude = ShapeUtils::evalDimsToExclude(input.rankOf(), tadDimensions); - - // increasing counter to skip numIndices - e++; - std::vector indices; - for (; e < static_cast(intArgs->size()); e++) - indices.push_back((*intArgs)[e]); - - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) { - auto inSubArr = input(indices[i], dimsToExclude, true); - auto updSubArr = updates(i, dimsToExclude, true); - - if (inSubArr.lengthOf() != updSubArr.lengthOf()) - continue; - - switch (opCode) { - case 0: - inSubArr.applyPairwiseTransform(pairwise::Add, updSubArr, inSubArr); - break; - case 1: - inSubArr.applyPairwiseTransform(pairwise::Subtract, updSubArr, inSubArr); - break; - case 2: - inSubArr.applyPairwiseTransform(pairwise::Multiply, updSubArr, inSubArr); - break; - case 3: - inSubArr.applyPairwiseTransform(pairwise::Divide, updSubArr, inSubArr); - break; - case 4: - inSubArr.applyPairwiseTransform(pairwise::ReverseSubtract, updSubArr, inSubArr); - break; - case 5: - inSubArr.applyPairwiseTransform(pairwise::ReverseDivide, updSubArr, inSubArr); - break; - case 6: - inSubArr.applyPairwiseTransform(pairwise::CopyPws, updSubArr, inSubArr); - break; - default: - continue; - } - } - }; - - samediff::Threads::parallel_tad(func, 0, indices.size()); -} - - -////////////////////////////////////////////////////////////////////////// -void scatterSimple(sd::LaunchContext * context, const int opId, NDArray& input, const NDArray& updates, const NDArray& indices, const std::vector& dimensions) { - - // updates and indices have same length - const Nd4jLong len = indices.lengthOf(); - - switch (opId) { - - case 6: { // copy - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) { - auto inSubArr = input(i, dimensions); - inSubArr.p(indices.t(i), updates.e(i)); - } - }; - - samediff::Threads::parallel_for(func, 0, len); - } - break; - - default: - throw std::invalid_argument("helpers::scatterSimple: operation is not implemented for given id !"); - } -} - -////////////////////////////////////////////////////////////////////////// -template -static void mergeMaxIndex_(const std::vector& inArrs, NDArray& output) { - - const Nd4jLong numArgs = inArrs.size(); - auto x = inArrs[0]; - - auto func = PRAGMA_THREADS_FOR { - for (auto e = start; e < stop; e++) { - T max = -DataTypeUtils::max(); - Nd4jLong idx = 0; - - for (Nd4jLong i = 0; i < numArgs; i++) { - T v = inArrs[i]->e(e); - if (v > max) { - max = v; - idx = i; - } - } - output.p(e, idx); - } - }; - - samediff::Threads::parallel_for(func, 0, x->lengthOf()); -} - -void mergeMaxIndex(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - BUILD_SINGLE_SELECTOR(inArrs[0]->dataType(), mergeMaxIndex_, (inArrs, output), LIBND4J_TYPES); -} - - -////////////////////////////////////////////////////////////////////////// -template -static void mergeMax_(const std::vector& inArrs, NDArray& output) { - const Nd4jLong numArgs = inArrs.size(); - auto x = inArrs[0]; - - auto func = PRAGMA_THREADS_FOR { - for (auto e = start; e < stop; e++) { - T max = -DataTypeUtils::max(); - for (Nd4jLong i = 0; i < numArgs; i++) { - T v = inArrs[i]->e(e); - if (v > max) - max = v; - } - output.p(e, max); - } - }; - - samediff::Threads::parallel_for(func, 0, x->lengthOf()); -} - -void mergeMax(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - BUILD_SINGLE_SELECTOR(output.dataType(), mergeMax_, (inArrs, output), LIBND4J_TYPES); -} - -////////////////////////////////////////////////////////////////////////// -template -static void mergeAvg_(const std::vector& inArrs, NDArray& output) { - const Nd4jLong numArgs = inArrs.size(); - const T factor = 1.f / numArgs; - auto x = inArrs[0]; - - auto func = PRAGMA_THREADS_FOR { - for (auto e = start; e < stop; e++) { - T sum = 0.; - for (Nd4jLong i = 0; i < numArgs; i++) { - T v = inArrs[i]->e(e); - sum += v; - } - output.p(e, sum * factor); - } - }; - - samediff::Threads::parallel_for(func, 0, x->lengthOf()); -} - -void mergeAvg(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - BUILD_SINGLE_SELECTOR(output.dataType(), mergeAvg_, (inArrs, output), LIBND4J_TYPES); -} - - -////////////////////////////////////////////////////////////////////////// -template -static void mergeAdd_(const std::vector& inArrs, NDArray& output) { - - const Nd4jLong numArgs = inArrs.size(); - auto x = inArrs[0]; - - auto func = PRAGMA_THREADS_FOR { - for (auto e = start; e < stop; e++) { - T sum = (T) 0.f; - for (Nd4jLong i = 0; i < numArgs; i++) - sum += inArrs[i]->e(e); - - output.p(e, sum); - } - }; - - samediff::Threads::parallel_for(func, 0, x->lengthOf()); -} - void mergeAdd(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - BUILD_SINGLE_SELECTOR(output.dataType(), mergeAdd_, (inArrs, output), LIBND4J_TYPES); - } - -////////////////////////////////////////////////////////////////////////// -template -static void clipByNorm_(NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace) { - - const int rank = input.rankOf(); - const auto norm2 = input.reduceAlongDimension(reduce::Norm2, dimensions); - - const T normActual = norm2.e(0); - const T normClip = clipNorm.e(0); - - if (isInplace) { - - if(norm2.lengthOf() == 1) { - - if(normActual > normClip) - input *= (normClip / normActual); - } - else { - - auto listOfInSubArrs = input.allTensorsAlongDimension(dimensions); - - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) { - const T iNormActual = norm2.e(i); - if (iNormActual > normClip) - *listOfInSubArrs.at(i) *= normClip / iNormActual; - } - }; - samediff::Threads::parallel_tad(func, 0, listOfInSubArrs.size()); - } - } - else { - - if(norm2.lengthOf() == 1) { - - if(normActual > normClip) - output.assign(input * (normClip / normActual)); - else - output.assign(input); - } - else { - - auto listOfInSubArrs = input.allTensorsAlongDimension(dimensions); - auto listOfOutSubArrs = output.allTensorsAlongDimension(dimensions); - - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) { - auto inputSubArr = listOfInSubArrs.at(i); - auto outputSubArr = listOfOutSubArrs.at(i); - outputSubArr->assign(inputSubArr); - - const T iNormActual = norm2.e(i); - - if (iNormActual > clipNorm.e(0)) - *outputSubArr *= clipNorm / iNormActual; - } - }; - samediff::Threads::parallel_tad(func, 0, listOfInSubArrs.size()); - } - } -} - -////////////////////////////////////////////////////////////////////////// -void clipByNorm(sd::LaunchContext * context, NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace) { - BUILD_SINGLE_SELECTOR(output.dataType(), clipByNorm_, (input, output, dimensions, clipNorm, isInplace), FLOAT_TYPES); -} - - - - - - - - - - - template - static void clipByGlobalNorm_(std::vector const& inputs, double clipNorm, sd::memory::Workspace* workspace, std::vector& outputs, bool isInplace) { - T globalNorm = 0; //NDArrayFactory::create(0, inputs[0]->getContext()); //sqrt(sum([l2norm(t)**2 for t in t_list])) -// PRAGMA_OMP_PARALLEL_FOR_SIMD_REDUCTION(sumT : globalNorm) - for (size_t i = 0; i < inputs.size(); i++) { - auto input = inputs[i]; - auto l2norm = input->reduceNumber(reduce::Norm2); - globalNorm += l2norm.t(0) * l2norm.t(0); - } - - //globalNorm.applyTransform(transform::Sqrt, nullptr, nullptr);// = sd::math::nd4j_sqrt(globalNorm); - auto normS = sd::math::nd4j_sqrt(globalNorm); - outputs[inputs.size()]->p(0, normS); - - const T factor = clipNorm / normS; - -// PRAGMA_OMP_PARALLEL_FOR - for (size_t e = 0; e < inputs.size(); e++) { - // all-reduce - auto input = inputs[e]; - auto output = outputs[e]; - - if (normS <= clipNorm) { - output->assign(input); - } - else { - - auto lambda = LAMBDA_T(_x, factor) { return _x * factor; }; - input->applyLambda(lambda, *output); - } - } - } - void clipByGlobalNorm(sd::LaunchContext * context, std::vector const& inputs, double clipNorm, sd::memory::Workspace* workspace, std::vector& outputs, bool isInplace) { - BUILD_SINGLE_SELECTOR(outputs[0]->dataType(), clipByGlobalNorm_, (inputs, clipNorm, workspace, outputs, isInplace), FLOAT_TYPES); - } - - BUILD_SINGLE_TEMPLATE(template void clipByGlobalNorm_, (std::vector const& inputs, double clipNorm, sd::memory::Workspace* workspace, std::vector& outputs, bool isInplace), FLOAT_TYPES); - -////////////////////////////////////////////////////////////////////////// -template -static void clipByNormBP_(const NDArray& input, const NDArray& gradO, NDArray& gradI /*output*/, const std::vector& dimensions, const NDArray& clipNorm) { - - const int rank = input.rankOf(); - - auto norm2 = input.reduceAlongDimension(reduce::Norm2, dimensions); - - if(norm2.lengthOf() == 1) { - - const T N = norm2.e(0); - - auto cn = clipNorm.e(0); - - if(N > cn) { - - const T sumOfProd = (input * gradO).reduceNumber(reduce::Sum).e(0); // reduce to scalar - const T factor1 = static_cast(1.f) / N; - const T factor3 = factor1 / (N * N); // 1 / (N*N*N) - - auto lambda = LAMBDA_TT(elem1, elem2, cn, sumOfProd, factor1, factor3) { - return cn * (factor1 * elem2 - factor3 * elem1 * sumOfProd); - }; - - (const_cast(input)).applyPairwiseLambda(const_cast(gradO), lambda, gradI); - } - else - gradI.assign(gradO); - } - else { - - auto gradISubArrs = gradI.allTensorsAlongDimension({dimensions}); - auto gradOSubArrs = gradO.allTensorsAlongDimension({dimensions}); - auto inputSubArrs = input.allTensorsAlongDimension({dimensions}); - - auto cn = clipNorm.e(0); - - auto func = PRAGMA_THREADS_FOR { - for (auto i = start; i < stop; i++) { - T N = norm2.e(i); - - auto gradOSubArr = gradOSubArrs.at(i); - auto gradISubArr = gradISubArrs.at(i); - - if (N > cn) { - auto inputSubArr = inputSubArrs.at(i); - const T sumOfProd = (*inputSubArr * *gradOSubArr).reduceNumber(reduce::Sum).e(0); // reduce to scalar - const T factor1 = static_cast(1.f) / N; - const T factor3 = factor1 / (N * N); // 1 / (N*N*N) - - auto lambda = LAMBDA_TT(elem1, elem2, cn, sumOfProd, factor1, factor3) { - return cn * (factor1 * elem2 - factor3 * elem1 * sumOfProd); - }; - - inputSubArr->applyPairwiseLambda(*gradOSubArr, lambda, *gradISubArr); - } else - gradISubArr->assign(gradOSubArr); - } - }; - samediff::Threads::parallel_tad(func, 0, gradISubArrs.size()); - } -} - - void clipByNormBP(sd::LaunchContext * context, const NDArray& input, const NDArray& gradO, NDArray& gradI /*output*/, const std::vector& dimensions, const NDArray& clipNorm) { - BUILD_SINGLE_SELECTOR(gradI.dataType(), clipByNormBP_, (input, gradO, gradI, dimensions, clipNorm), FLOAT_TYPES); - } - - BUILD_SINGLE_TEMPLATE(template void clipByNormBP_, (const NDArray& input, const NDArray& gradO, NDArray& gradI /*output*/, const std::vector& dimensions, const NDArray& clipNorm), FLOAT_TYPES); - - -////////////////////////////////////////////////////////////////////////// -template -static void clipByAveraged_(NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace) { - - auto cn = clipNorm.e(0); - if (dimensions.size() == 0) { - // all-reduce - T n2 = input.reduceNumber(reduce::Norm2).e(0) / input.lengthOf(); - if (n2 <= cn) { - if (!isInplace) - output.assign(input); - } - else { - const T factor = cn / n2; - auto lambda = LAMBDA_T(_x, factor) { return _x * factor; }; - input.applyLambda(lambda, output); - } - } - else { - // along dimension - auto norm2 = input.reduceAlongDimension(reduce::Norm2, dimensions, false); - if (!isInplace) - output.assign(input); - auto tads = output.allTensorsAlongDimension(dimensions); - // TODO: make this CUDA-compliant somehow - for (int e = 0; e < tads.size(); e++) { - T n2 = norm2.e(e) / tads.at(e)->lengthOf(); - const T factor = cn / n2; - if (n2 > cn) { - auto lambda = LAMBDA_T(_x, factor) {return _x * factor;}; - tads.at(e)->applyLambda(lambda, output); - } - } - } -} - - void clipByAveraged(sd::LaunchContext * context, NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace) { - BUILD_SINGLE_SELECTOR(input.dataType(), clipByAveraged_, (input, output, dimensions, clipNorm, isInplace), FLOAT_TYPES); - } - - BUILD_SINGLE_TEMPLATE(template void clipByAveraged_, (NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace), FLOAT_TYPES); - -/* - if (d1 > params[1]) - return params[1]; - else if (d1 < params[0]) - return params[0]; - else return d1; -*/ - - template - static void clipByValue_(NDArray& input, double leftBound, double rightBound, NDArray& output) { - auto routine = LAMBDA_T(_x, leftBound, rightBound) { - if (_x > rightBound) return rightBound; - if (_x < leftBound) return leftBound; - return _x; - }; - - input.applyLambda(routine, output); - } - - void clipByValue(sd::LaunchContext * context, NDArray& input, double leftBound, double rightBound, NDArray& output) { - BUILD_SINGLE_SELECTOR(input.dataType(), clipByValue_, (input, leftBound, rightBound, output), FLOAT_TYPES); - } - - BUILD_SINGLE_TEMPLATE(template void clipByValue_, (NDArray& input, double leftBound, double rightBound, NDArray& output);, FLOAT_TYPES); - -////////////////////////////////////////////////////////////////////////// -template -static void mirrorPad_(const NDArray& input, const NDArray& paddings, NDArray& output, const int mode) { - - // mode: 0 - REFLECT, else - SYMMETRIC - const int reflBorder = (bool)mode ? 1 : 0; - const int rank = input.rankOf(); - const Nd4jLong outLen = output.lengthOf(); - - if(rank <= 1) { - - const Nd4jLong inLen = input.lengthOf(); - const auto leftSide = paddings.e(0); - const auto leftSideCorrected = leftSide - reflBorder; - const Nd4jLong len = 2*(inLen-1) + leftSide + reflBorder; - - for(int i = 0; i < outLen; ++i) { - - if (i < leftSide) // left side - output.p(i, input.e(leftSideCorrected - i)); - - else if(i >= leftSide && i < leftSide + inLen) // middle - output.p(i, input.e(i - leftSide)); - - else // right side - output.p(i, input.e(len - i)); - } - } - else { - - auto func = PRAGMA_THREADS_FOR { - - int inIdx[MAX_RANK], outIdx[MAX_RANK]; - - for (auto i = start; i < stop; i++) { - - shape::index2coordsCPU(start, i, output.getShapeInfo(), outIdx); - - for (int j = 0; j < rank; ++j) { - const Nd4jLong inLen = input.sizeAt(j); - const auto leftSide = paddings.e(j, 0); - const auto leftSideCorrected = leftSide - reflBorder; - const Nd4jLong len = 2 * (inLen - 1) + leftSide + reflBorder; - - if (outIdx[j] < leftSide) // left side - inIdx[j] = leftSideCorrected - outIdx[j]; - - else if (outIdx[j] >= leftSide && outIdx[j] < leftSide + inLen) // middle - inIdx[j] = outIdx[j] - leftSide; - - else // right side - inIdx[j] = len - outIdx[j]; - } - - auto outOffset = shape::getOffset(output.getShapeInfo(), outIdx); - auto inOffset = shape::getOffset(input.getShapeInfo(), inIdx); - reinterpret_cast(output.buffer())[outOffset] = reinterpret_cast(input.getBuffer())[inOffset]; - } - }; - - samediff::Threads::parallel_for(func, 0, outLen); - } -} - - void mirrorPad(sd::LaunchContext * context, const NDArray& input, const NDArray& paddings, NDArray& output, const int mode) { - BUILD_SINGLE_SELECTOR(input.dataType(), mirrorPad_, (input, paddings, output, mode), LIBND4J_TYPES); - } - - BUILD_SINGLE_TEMPLATE(template void mirrorPad_, (const NDArray& input, const NDArray& paddings, NDArray& output, const int mode), LIBND4J_TYPES); - - -////////////////////////////////////////////////////////////////////////// -template -static void tileBP_(const NDArray& gradO /*input*/, NDArray& gradI /*output*/, const std::vector reps) { - - T* gradIBuff = reinterpret_cast(gradI.getBuffer()); - const T* gradOBuff = reinterpret_cast(gradO.getBuffer()); - const Nd4jLong gradILen = gradI.lengthOf(); - const Nd4jLong gradOLen = gradO.lengthOf(); // gradOLen >= gradILen - const Nd4jLong gradIEWS = sd::math::nd4j_abs(gradI.ews()); - const Nd4jLong gradOEWS = gradO.ews(); - - // initial zeroing of gradI content - if(gradIEWS == 1) - memset(gradIBuff, 0, gradILen * sizeof(T)); - else { - //PRAGMA_OMP_PARALLEL_FOR_SIMD - for (Nd4jLong i = 0; i < gradILen * gradIEWS; i += gradIEWS) - gradIBuff[i] = static_cast(0.f); - } - - - if(gradO.ordering() == 'c' && gradOEWS == 1) { - - //PRAGMA_OMP_PARALLEL_FOR_SIMD - for(Nd4jLong i=0; i(idx) + gradOBuff[i]); - } - } - else if(gradO.ordering() == 'c' && gradOEWS > 1) { - - //PRAGMA_OMP_PARALLEL_FOR_SIMD - for(Nd4jLong i=0; i(idx) + gradOBuff[i * gradOEWS]); - } - } - else { - - //PRAGMA_OMP_PARALLEL_FOR_SIMD - for(Nd4jLong i=0; i(fidx) + gradOBuff[shape::getIndexOffset(i, gradO.getShapeInfo())]); - } - } -} - -void tileBP(sd::LaunchContext * context, const NDArray& gradO /*input*/, NDArray& gradI /*output*/, const std::vector reps) { - BUILD_SINGLE_SELECTOR(gradI.dataType(), tileBP_, (gradO, gradI, reps), FLOAT_TYPES); -} - - -BUILD_SINGLE_TEMPLATE(template void tileBP_, (const NDArray& gradO /*input*/, NDArray& gradI /*output*/, const std::vector reps), FLOAT_TYPES); - - - - - -} -} -} diff --git a/libnd4j/include/ops/declarable/helpers/cpu/triu.cpp b/libnd4j/include/ops/declarable/helpers/cpu/triu.cpp new file mode 100644 index 000000000..4194e976c --- /dev/null +++ b/libnd4j/include/ops/declarable/helpers/cpu/triu.cpp @@ -0,0 +1,56 @@ +/******************************************************************************* + * Copyright (c) 2015-2018 Skymind, Inc. + * + * This program and the accompanying materials are made available under the + * terms of the Apache License, Version 2.0 which is available at + * https://www.apache.org/licenses/LICENSE-2.0. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, WITHOUT + * WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the + * License for the specific language governing permissions and limitations + * under the License. + * + * SPDX-License-Identifier: Apache-2.0 + ******************************************************************************/ + +// +// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 +// + + +#include +#include + +namespace sd { +namespace ops { +namespace helpers { + + +////////////////////////////////////////////////////////////////////////// +template +static void triuBP_(sd::LaunchContext * context, const NDArray& input, const NDArray& gradO, NDArray& gradI, const int diagonal) { + + auto dOdI = NDArray(&gradO); // dO/dI + const_cast(input).fillAsTriangular(0, diagonal, dOdI.sizeAt(-1), dOdI, 'b'); + int dLen = dOdI.lengthOf(); + + auto func = PRAGMA_THREADS_FOR { + for (auto i = start; i < stop; i++) { + if (dOdI.t(i) != static_cast(0.f)) + dOdI.t(i) = static_cast(1.f); + } + }; + samediff::Threads::parallel_for(func, 0, dLen); + + // FIXME: !!! + gradI.assign(dOdI * gradO); // chain rule: dLoss/dI = dO/dI * dLoss/dO +} + + void triuBP(sd::LaunchContext * context, const NDArray& input, const NDArray& gradO, NDArray& gradI, const int diagonal) { + BUILD_SINGLE_SELECTOR(gradO.dataType(), triuBP_, (context, input, gradO, gradI, diagonal), LIBND4J_TYPES); + } + +} +} +} diff --git a/libnd4j/include/ops/declarable/helpers/cuda/merge.cu b/libnd4j/include/ops/declarable/helpers/cuda/merge.cu index b448fbd35..a7dd9b199 100644 --- a/libnd4j/include/ops/declarable/helpers/cuda/merge.cu +++ b/libnd4j/include/ops/declarable/helpers/cuda/merge.cu @@ -14,9 +14,9 @@ * SPDX-License-Identifier: Apache-2.0 ******************************************************************************/ -// -// @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 -// + // + // @author Yurii Shyrma (iuriish@yahoo.com), created on 20.04.2018 + // #include @@ -34,7 +34,7 @@ namespace sd { namespace helpers { ////////////////////////////////////////////////////////////////////////// template - static __global__ void global_mergeMaxIndex_(void **inArrs, void **inShapes, const int numArrays, void *voutput, Nd4jLong *outputShape, Nd4jLong length) { + static __global__ void mergeMaxIndexCudaLauncher(void** inArrs, void** inShapes, const int numArrays, void* voutput, Nd4jLong* outputShape, Nd4jLong length) { auto output = reinterpret_cast(voutput); const auto tid = blockIdx.x * blockDim.x + threadIdx.x; @@ -46,54 +46,56 @@ namespace sd { for (int i = 0; i < numArrays; i++) { auto x = reinterpret_cast(inArrs[i]); - auto xShape = reinterpret_cast(inShapes[i]); + auto xShape = reinterpret_cast(inShapes[i]); auto val = x[shape::getIndexOffset(e, xShape)];; if (mVal < val) { mIdx = static_cast(i); mVal = val; } } - __syncthreads(); - + output[shape::getIndexOffset(e, outputShape)] = mIdx; } } template - static void mergeMaxIndex_(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - std::vector inBuffers(inArrs.size()); - std::vector inShapes(inArrs.size()); + static void mergeMaxIndex_(sd::LaunchContext* context, const std::vector& inArrs, NDArray& output) { + + int nArrSize = static_cast(inArrs.size()); + std::vector inBuffers(nArrSize), inShapes(nArrSize); - for (int e = 0; e < inArrs.size(); e++) { + for (int e = 0; e < nArrSize; e++) { inBuffers[e] = inArrs[e]->getSpecialBuffer(); inShapes[e] = inArrs[e]->getSpecialShapeInfo(); } PointersManager manager(context, "mergeMaxIndex"); - auto pInBuffers = reinterpret_cast(manager.replicatePointer(inBuffers.data(), inBuffers.size() * sizeof(void *))); - auto pInShapes = reinterpret_cast(manager.replicatePointer(inShapes.data(), inShapes.size() * sizeof(void *))); + auto pInBuffers = reinterpret_cast(manager.replicatePointer(inBuffers.data(), inBuffers.size() * sizeof(void*))); + auto pInShapes = reinterpret_cast(manager.replicatePointer(inShapes.data(), inShapes.size() * sizeof(void*))); auto length = output.lengthOf(); - global_mergeMaxIndex_<<<512, 512, 512, *context->getCudaStream()>>>(pInBuffers, pInShapes, (int) inArrs.size(), output.getSpecialBuffer(), output.getSpecialShapeInfo(), length); + const int threadsPerBlock = MAX_NUM_THREADS / 2; + const int blocksPerGrid = (length + threadsPerBlock - 1) / threadsPerBlock; + + mergeMaxIndexCudaLauncher << getCudaStream() >> > (pInBuffers, pInShapes, nArrSize, output.getSpecialBuffer(), output.getSpecialShapeInfo(), length); manager.synchronize(); } - void mergeMaxIndex(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - NDArray::prepareSpecialUse({&output}, {}); - for (auto v:inArrs) - v->syncToDevice(); - + void mergeMaxIndex(sd::LaunchContext* context, const std::vector& inArrs, NDArray& output) { + + NDArray::prepareSpecialUse({ &output }, inArrs); + BUILD_DOUBLE_SELECTOR(inArrs[0]->dataType(), output.dataType(), mergeMaxIndex_, (context, inArrs, output), LIBND4J_TYPES, INDEXING_TYPES); - NDArray::registerSpecialUse({&output}, {}); + NDArray::registerSpecialUse({ &output }, inArrs); } ////////////////////////////////////////////////////////////////////////// template - static __global__ void global_mergeMax_(void **inArrs, void **inShapes, const int numArrays, void *voutput, Nd4jLong *outputShape, Nd4jLong length) { + static __global__ void mergeMaxCudaLauncher(void** inArrs, void** inShapes, const int numArrays, void* voutput, Nd4jLong* outputShape, Nd4jLong length) { auto output = reinterpret_cast(voutput); const auto tid = blockIdx.x * blockDim.x + threadIdx.x; @@ -103,51 +105,163 @@ namespace sd { T mVal = -DataTypeUtils::max(); for (int i = 0; i < numArrays; i++) { - auto x = reinterpret_cast(inArrs[i]); - auto xShape = reinterpret_cast(inShapes[i]); + auto x = reinterpret_cast(inArrs[i]); + auto xShape = reinterpret_cast(inShapes[i]); auto val = x[shape::getIndexOffset(e, xShape)];; if (mVal < val) mVal = val; } - __syncthreads(); output[shape::getIndexOffset(e, outputShape)] = mVal; } } template - static void mergeMax_(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - std::vector inBuffers(inArrs.size()); - std::vector inShapes(inArrs.size()); + static void mergeMax_(sd::LaunchContext* context, const std::vector& inArrs, NDArray& output) { + + int nArrsSize = static_cast(inArrs.size()); - for (int e = 0; e < inArrs.size(); e++) { + std::vector inBuffers(nArrsSize), inShapes(nArrsSize); + + for (int e = 0; e < nArrsSize; e++) { inBuffers[e] = inArrs[e]->getSpecialBuffer(); inShapes[e] = inArrs[e]->getSpecialShapeInfo(); } PointersManager manager(context, "mergeMax"); - auto pInBuffers = reinterpret_cast(manager.replicatePointer(inBuffers.data(), inBuffers.size() * sizeof(void *))); - auto pInShapes = reinterpret_cast(manager.replicatePointer(inShapes.data(), inShapes.size() * sizeof(void *))); + auto pInBuffers = reinterpret_cast(manager.replicatePointer(inBuffers.data(), inBuffers.size() * sizeof(void*))); + auto pInShapes = reinterpret_cast(manager.replicatePointer(inShapes.data(), inShapes.size() * sizeof(void*))); auto length = output.lengthOf(); - global_mergeMax_<<<512, 512, 512, *context->getCudaStream()>>>(pInBuffers, pInShapes, (int) inArrs.size(), output.getSpecialBuffer(), output.getSpecialShapeInfo(), length); + const int threadsPerBlock = MAX_NUM_THREADS / 2; + const int blocksPerGrid = (length + threadsPerBlock - 1) / threadsPerBlock; + + mergeMaxCudaLauncher << getCudaStream() >> > (pInBuffers, pInShapes, nArrsSize, output.getSpecialBuffer(), output.getSpecialShapeInfo(), length); manager.synchronize(); } - void mergeMax(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - NDArray::prepareSpecialUse({&output}, {}); - for (auto v:inArrs) - v->syncToDevice(); - + void mergeMax(sd::LaunchContext* context, const std::vector& inArrs, NDArray& output) { + + NDArray::prepareSpecialUse({ &output }, inArrs); + BUILD_SINGLE_SELECTOR(output.dataType(), mergeMax_, (context, inArrs, output), LIBND4J_TYPES); - NDArray::registerSpecialUse({&output}, {}); + + NDArray::registerSpecialUse({ &output }, inArrs); } ////////////////////////////////////////////////////////////////////////// template - static __global__ void global_mergeAvg_(void **inArrs, void **inShapes, const int numArrays, void *voutput, Nd4jLong *outputShape, Nd4jLong length) { + static __global__ void mergeMaxBpCudaLauncher(void** inArrs, void** inShapes, void* vgradient, Nd4jLong* gradientShape, const int numArrays, + void** outArrs, void** outShapes, Nd4jLong length, bool bSameOrderAndEws1) { + + auto grad = reinterpret_cast(vgradient); + + const auto tid = blockIdx.x * blockDim.x + threadIdx.x; + const auto step = gridDim.x * blockDim.x; + + int coords[MAX_RANK]; + + for (Nd4jLong e = tid; e < length; e += step) { + + T mVal = -DataTypeUtils::max(); + int nMaxIndex = 0; + auto xOffset = e, zOffset = e, gradOffset = e; + + if (!bSameOrderAndEws1) { + shape::index2coords(e, gradientShape, coords); + gradOffset = shape::getOffset(gradientShape, coords); + } + + for (int i = 0; i < numArrays; i++) { + auto x = reinterpret_cast(inArrs[i]); + + if (!bSameOrderAndEws1) { + auto xShape = reinterpret_cast(inShapes[i]); + xOffset = shape::getOffset(xShape, coords); + } + + auto val = x[xOffset]; + if (mVal < val) { + mVal = val; + nMaxIndex = i; + } + } + + // outputs have to be pre-nullify + if (!bSameOrderAndEws1) { + auto outShape = reinterpret_cast(outShapes[nMaxIndex]); + zOffset = shape::getOffset(outShape, coords); + } + + auto output = reinterpret_cast(outArrs[nMaxIndex]); + + output[zOffset] = grad[gradOffset]; + } + } + + template + static void mergeMaxBp_(sd::LaunchContext* context, const std::vector& inArrs, std::vector& outArrs, int nArrSize, bool bSameOrderAndEws1) { + + std::vector inBuffers(nArrSize), inShapes(nArrSize), outBuffers(nArrSize), outShapes(nArrSize); + + for (int e = 0; e < nArrSize; e++) { + inBuffers[e] = inArrs[e]->getSpecialBuffer(); + inShapes[e] = inArrs[e]->getSpecialShapeInfo(); + outBuffers[e] = outArrs[e]->getSpecialBuffer(); + outShapes[e] = outArrs[e]->getSpecialShapeInfo(); + } + + PointersManager manager(context, "mergeMaxBp"); + + auto pInBuffers = reinterpret_cast(manager.replicatePointer(inBuffers.data(), inBuffers.size() * sizeof(void*))); + auto pInShapes = reinterpret_cast(manager.replicatePointer(inShapes.data(), inShapes.size() * sizeof(void*))); + + auto pOutBuffers = reinterpret_cast(manager.replicatePointer(outBuffers.data(), outBuffers.size() * sizeof(void*))); + auto pOutShapes = reinterpret_cast(manager.replicatePointer(outShapes.data(), outShapes.size() * sizeof(void*))); + + auto length = inArrs[nArrSize]->lengthOf(); + + const int threadsPerBlock = MAX_NUM_THREADS / 2; + const int blocksPerGrid = (length + threadsPerBlock - 1) / threadsPerBlock; + + mergeMaxBpCudaLauncher << getCudaStream() >> > (pInBuffers, pInShapes, inArrs[nArrSize]->getSpecialBuffer(), + inArrs[nArrSize]->getSpecialShapeInfo(), nArrSize, pOutBuffers, pOutShapes, + length, bSameOrderAndEws1); + + manager.synchronize(); + } + + void mergeMaxBp(sd::LaunchContext* context, const std::vector& inArrs, std::vector& outArrs) { + + // not use gradient + int nArrSize = static_cast(inArrs.size() - 1); + + const std::vector& out = reinterpret_cast&>(outArrs); + + NDArray::prepareSpecialUse(out, inArrs); + + bool bSameOrderAndEws1 = (1 == inArrs[nArrSize]->ews()); + auto ordering = inArrs[nArrSize]->ordering(); + + for (int i = 0; i < nArrSize; ++i) { + bSameOrderAndEws1 &= (ordering == inArrs[i]->ordering()); + bSameOrderAndEws1 &= (1 == inArrs[i]->ews()); + + bSameOrderAndEws1 &= (ordering == outArrs[i]->ordering()); + bSameOrderAndEws1 &= (1 == outArrs[i]->ews()); + } + + BUILD_SINGLE_SELECTOR(inArrs[nArrSize]->dataType(), mergeMaxBp_, (context, inArrs, outArrs, nArrSize, bSameOrderAndEws1), LIBND4J_TYPES); + + NDArray::registerSpecialUse( out, inArrs ); + } + + + ////////////////////////////////////////////////////////////////////////// + template + static __global__ void mergeAvgCudaLauncher(void** inArrs, void** inShapes, const int numArrays, void* voutput, Nd4jLong* outputShape, Nd4jLong length) { auto output = reinterpret_cast(voutput); const auto tid = blockIdx.x * blockDim.x + threadIdx.x; @@ -158,7 +272,7 @@ namespace sd { for (int i = 0; i < numArrays; i++) { auto x = reinterpret_cast(inArrs[i]); - auto xShape = reinterpret_cast(inShapes[i]); + auto xShape = reinterpret_cast(inShapes[i]); sum += x[shape::getIndexOffset(e, xShape)]; } @@ -168,9 +282,9 @@ namespace sd { } template - static void mergeAvg_(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - std::vector inBuffers(inArrs.size()); - std::vector inShapes(inArrs.size()); + static void mergeAvg_(sd::LaunchContext* context, const std::vector& inArrs, NDArray& output) { + + std::vector inBuffers(inArrs.size()), inShapes(inArrs.size()); for (int e = 0; e < inArrs.size(); e++) { inBuffers[e] = inArrs[e]->getSpecialBuffer(); @@ -179,28 +293,111 @@ namespace sd { PointersManager manager(context, "mergeAvg"); - auto pInBuffers = reinterpret_cast(manager.replicatePointer(inBuffers.data(), inBuffers.size() * sizeof(void *))); - auto pInShapes = reinterpret_cast(manager.replicatePointer(inShapes.data(), inShapes.size() * sizeof(void *))); + auto pInBuffers = reinterpret_cast(manager.replicatePointer(inBuffers.data(), inBuffers.size() * sizeof(void*))); + auto pInShapes = reinterpret_cast(manager.replicatePointer(inShapes.data(), inShapes.size() * sizeof(void*))); auto length = output.lengthOf(); - global_mergeAvg_<<<512, 512, 512, *context->getCudaStream()>>>(pInBuffers, pInShapes, (int) inArrs.size(), output.getSpecialBuffer(), output.getSpecialShapeInfo(), length); + const int threadsPerBlock = MAX_NUM_THREADS / 2; + const int blocksPerGrid = (length + threadsPerBlock - 1) / threadsPerBlock; + + mergeAvgCudaLauncher << getCudaStream() >> > (pInBuffers, pInShapes, (int)inArrs.size(), output.getSpecialBuffer(), output.getSpecialShapeInfo(), length); manager.synchronize(); } - void mergeAvg(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - NDArray::prepareSpecialUse({&output}, {}); - for (auto v:inArrs) - v->syncToDevice(); + void mergeAvg(sd::LaunchContext* context, const std::vector& inArrs, NDArray& output) { + + NDArray::prepareSpecialUse({ &output }, inArrs); BUILD_SINGLE_SELECTOR(output.dataType(), mergeAvg_, (context, inArrs, output), FLOAT_TYPES); - NDArray::registerSpecialUse({&output}, {}); + NDArray::registerSpecialUse({ &output }, inArrs); + } + ////////////////////////////////////////////////////////////////////////// + template + static __global__ void mergeAvgBpCudaLauncher(void* vgradient, Nd4jLong* gradientShape, void** outArrs, void** outShapes, + const int numArrays, Nd4jLong length, bool bSameOrderAndEws1) { + + auto grad = reinterpret_cast(vgradient); + + const auto tid = blockIdx.x * blockDim.x + threadIdx.x; + const auto step = gridDim.x * blockDim.x; + + int coords[MAX_RANK]; + + for (Nd4jLong e = tid; e < length; e += step) { + + auto zOffset = e, gradOffset = e; + if (!bSameOrderAndEws1) { + shape::index2coords(e, gradientShape, coords); + gradOffset = shape::getOffset(gradientShape, coords); + } + + for (int i = 0; i < numArrays; i++) { + + if (!bSameOrderAndEws1) { + auto outShape = reinterpret_cast(outShapes[i]); + zOffset = shape::getOffset(outShape, coords); + } + + auto output = reinterpret_cast(outArrs[i]); + + output[zOffset] = grad[gradOffset] / numArrays; + } + } + } + + template + static void mergeAvgBp_(sd::LaunchContext* context, const NDArray& gradient, std::vector& outArrs, bool bSameOrderAndEws1) { + + int nArrSize = static_cast(outArrs.size()); + + std::vector outBuffers(nArrSize), outShapes(nArrSize); + + for (int e = 0; e < nArrSize; e++) { + outBuffers[e] = outArrs[e]->getSpecialBuffer(); + outShapes[e] = outArrs[e]->getSpecialShapeInfo(); + } + + PointersManager manager(context, "mergeAvgBp"); + + auto pOutBuffers = reinterpret_cast(manager.replicatePointer(outBuffers.data(), outBuffers.size() * sizeof(void*))); + auto pOutShapes = reinterpret_cast(manager.replicatePointer(outShapes.data(), outShapes.size() * sizeof(void*))); + + auto length = gradient.lengthOf(); + + const int threadsPerBlock = MAX_NUM_THREADS / 2; + const int blocksPerGrid = (length + threadsPerBlock - 1) / threadsPerBlock; + + mergeAvgBpCudaLauncher << getCudaStream() >> > (gradient.getSpecialBuffer(), gradient.getSpecialShapeInfo(), + pOutBuffers, pOutShapes, nArrSize, length, bSameOrderAndEws1); + + manager.synchronize(); + } + + void mergeAvgBp(sd::LaunchContext* context, const NDArray& gradient, std::vector& outArrs) { + + const std::vector& out = reinterpret_cast&>(outArrs); + + NDArray::prepareSpecialUse( out, { &gradient }); + + bool bSameOrderAndEws1 = (1 == gradient.ews()); + auto ordering = gradient.ordering(); + + for (const auto& v : outArrs) { + bSameOrderAndEws1 &= (ordering == v->ordering()); + bSameOrderAndEws1 &= (1 == v->ews()); + } + + BUILD_SINGLE_SELECTOR(gradient.dataType(), mergeAvgBp_, (context, gradient, outArrs, bSameOrderAndEws1), LIBND4J_TYPES); + + NDArray::prepareSpecialUse(out, { &gradient }); } ////////////////////////////////////////////////////////////////////////// template - static __global__ void global_mergeAdd_(void **inArrs, void **inShapes, const int numArrays, void *voutput, Nd4jLong *outputShape, Nd4jLong length) { + static __global__ void mergeAddCudaLauncher(void** inArrs, void** inShapes, const int numArrays, void* voutput, Nd4jLong* outputShape, Nd4jLong length) { + auto output = reinterpret_cast(voutput); const auto tid = blockIdx.x * blockDim.x + threadIdx.x; @@ -211,7 +408,7 @@ namespace sd { for (int i = 0; i < numArrays; i++) { auto x = reinterpret_cast(inArrs[i]); - auto xShape = reinterpret_cast(inShapes[i]); + auto xShape = reinterpret_cast(inShapes[i]); sum += x[shape::getIndexOffset(e, xShape)]; } @@ -221,36 +418,120 @@ namespace sd { } template - static void mergeAdd_(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - std::vector inBuffers(inArrs.size()); - std::vector inShapes(inArrs.size()); + static void mergeAdd_(sd::LaunchContext* context, const std::vector& inArrs, NDArray& output) { + + int nArrSize = static_cast(inArrs.size()); + std::vector inBuffers(nArrSize), inShapes(nArrSize); - for (int e = 0; e < inArrs.size(); e++) { + for (int e = 0; e < nArrSize; e++) { inBuffers[e] = inArrs[e]->getSpecialBuffer(); inShapes[e] = inArrs[e]->getSpecialShapeInfo(); } PointersManager manager(context, "mergeAdd"); - auto pInBuffers = reinterpret_cast(manager.replicatePointer(inBuffers.data(), inBuffers.size() * sizeof(void *))); - auto pInShapes = reinterpret_cast(manager.replicatePointer(inShapes.data(), inShapes.size() * sizeof(void *))); + auto pInBuffers = reinterpret_cast(manager.replicatePointer(inBuffers.data(), inBuffers.size() * sizeof(void*))); + auto pInShapes = reinterpret_cast(manager.replicatePointer(inShapes.data(), inShapes.size() * sizeof(void*))); auto length = output.lengthOf(); - global_mergeAdd_<<<512, 512, 512, *context->getCudaStream()>>>(pInBuffers, pInShapes, (int) inArrs.size(), output.getSpecialBuffer(), output.getSpecialShapeInfo(), length); + const int threadsPerBlock = MAX_NUM_THREADS / 2; + const int blocksPerGrid = (length + threadsPerBlock - 1) / threadsPerBlock; + + mergeAddCudaLauncher << getCudaStream() >> > (pInBuffers, pInShapes, nArrSize, output.getSpecialBuffer(), output.getSpecialShapeInfo(), length); manager.synchronize(); } - BUILD_SINGLE_TEMPLATE(template void mergeAdd_, (sd::LaunchContext * context, const std::vector& inArrs, NDArray& output), NUMERIC_TYPES); - - void mergeAdd(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output) { - NDArray::prepareSpecialUse({&output}, {}); - for (auto v:inArrs) - v->syncToDevice(); + BUILD_SINGLE_TEMPLATE(template void mergeAdd_, (sd::LaunchContext* context, const std::vector& inArrs, NDArray& output), NUMERIC_TYPES); + void mergeAdd(sd::LaunchContext* context, const std::vector& inArrs, NDArray& output) { + + NDArray::prepareSpecialUse({ &output }, inArrs); + BUILD_SINGLE_SELECTOR(output.dataType(), mergeAdd_, (context, inArrs, output), NUMERIC_TYPES); - NDArray::registerSpecialUse({&output}, {}); + NDArray::registerSpecialUse({ &output }, inArrs); } + + ////////////////////////////////////////////////////////////////////////// + template + static __global__ void mergeAddBpCudaLauncher(void* vgradient, Nd4jLong* gradientShape, void** outArrs, void** outShapes, + const int numArrays, Nd4jLong length, bool bSameOrderAndEws1) { + + auto grad = reinterpret_cast(vgradient); + + const auto tid = blockIdx.x * blockDim.x + threadIdx.x; + const auto step = gridDim.x * blockDim.x; + + int coords[MAX_RANK]; + + for (Nd4jLong e = tid; e < length; e += step) { + + auto zOffset = e, gradOffset = e; + if (!bSameOrderAndEws1) { + shape::index2coords(e, gradientShape, coords); + gradOffset = shape::getOffset(gradientShape, coords); + } + + for (int i = 0; i < numArrays; i++) { + + if (!bSameOrderAndEws1) { + auto outShape = reinterpret_cast(outShapes[i]); + zOffset = shape::getOffset(outShape, coords); + } + + auto output = reinterpret_cast(outArrs[i]); + + output[zOffset] = grad[gradOffset]; + } + } + } + + template + static void mergeAddBp_(sd::LaunchContext* context, const NDArray& gradient, std::vector& outArrs, bool bSameOrderAndEws1) { + + int nArrSize = static_cast(outArrs.size()); + + std::vector outBuffers(nArrSize), outShapes(nArrSize); + + for (int e = 0; e < nArrSize; e++) { + outBuffers[e] = outArrs[e]->getSpecialBuffer(); + outShapes[e] = outArrs[e]->getSpecialShapeInfo(); + } + + PointersManager manager(context, "mergeAddBp"); + + auto pOutBuffers = reinterpret_cast(manager.replicatePointer(outBuffers.data(), outBuffers.size() * sizeof(void*))); + auto pOutShapes = reinterpret_cast(manager.replicatePointer(outShapes.data(), outShapes.size() * sizeof(void*))); + + auto length = gradient.lengthOf(); + + const int threadsPerBlock = MAX_NUM_THREADS / 2; + const int blocksPerGrid = (length + threadsPerBlock - 1) / threadsPerBlock; + + mergeAddBpCudaLauncher << getCudaStream() >> > (gradient.getSpecialBuffer(), gradient.getSpecialShapeInfo(), + pOutBuffers, pOutShapes, nArrSize, length, bSameOrderAndEws1); + + manager.synchronize(); + } + + void mergeAddBp(sd::LaunchContext* context, const NDArray& gradient, std::vector& outArrs) { + + const std::vector& out = reinterpret_cast& >(outArrs); + NDArray::prepareSpecialUse( out, { &gradient }); + + bool bSameOrderAndEws1 = (1 == gradient.ews()); + auto ordering = gradient.ordering(); + + for (const auto& v : outArrs) { + bSameOrderAndEws1 &= (ordering == v->ordering()); + bSameOrderAndEws1 &= (1 == v->ews()); + } + + BUILD_SINGLE_SELECTOR(gradient.dataType(), mergeAddBp_, (context, gradient, outArrs, bSameOrderAndEws1), LIBND4J_TYPES); + + NDArray::prepareSpecialUse( out, { &gradient }); + } + } } -} \ No newline at end of file +} diff --git a/libnd4j/include/ops/declarable/helpers/transforms.h b/libnd4j/include/ops/declarable/helpers/transforms.h index 8678c16fd..6ebecd8f7 100644 --- a/libnd4j/include/ops/declarable/helpers/transforms.h +++ b/libnd4j/include/ops/declarable/helpers/transforms.h @@ -52,13 +52,16 @@ namespace helpers { void scatterSimple(sd::LaunchContext * context, const int opId, NDArray& input, const NDArray& updates, const NDArray& indices, const std::vector& dimensions); - void mergeMaxIndex(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output); + void mergeMaxIndex(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output); - void mergeMax(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output); + void mergeMax(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output); + void mergeMaxBp(sd::LaunchContext* context, const std::vector& inArrs, std::vector& outArrs); - void mergeAvg(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output); + void mergeAvg(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output); + void mergeAvgBp(sd::LaunchContext* context, const NDArray& gradient, std::vector& outArrs); - void mergeAdd(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output); + void mergeAdd(sd::LaunchContext * context, const std::vector& inArrs, NDArray& output); + void mergeAddBp(sd::LaunchContext* context, const NDArray& gradient, std::vector& outArrs); void clipByNorm(sd::LaunchContext * context, NDArray& input, NDArray& output, const std::vector& dimensions, const NDArray& clipNorm, const bool isInplace); void clipByGlobalNorm(sd::LaunchContext * context, std::vector const& inputs, double clipNorm, sd::memory::Workspace* workspace, std::vector& outputs, bool isInplace); diff --git a/libnd4j/tests_cpu/layers_tests/DeclarableOpsTests13.cpp b/libnd4j/tests_cpu/layers_tests/DeclarableOpsTests13.cpp index 20665e5d3..4b5a24bb9 100644 --- a/libnd4j/tests_cpu/layers_tests/DeclarableOpsTests13.cpp +++ b/libnd4j/tests_cpu/layers_tests/DeclarableOpsTests13.cpp @@ -955,7 +955,160 @@ TEST_F(DeclarableOpsTests13, mergemax_2) { ASSERT_EQ(20, status); } +///////////////////////////////////////////////////////////////////////// +TEST_F(DeclarableOpsTests13, mergemax_bp_1) { + NDArray x1('c', { 5, 5 }, sd::DataType::FLOAT32); + NDArray x2('c', { 5, 5 }, sd::DataType::FLOAT32); + NDArray x3('c', { 5, 5 }, sd::DataType::FLOAT32); + NDArray grad('c', { 5, 5 }, sd::DataType::FLOAT32); + + x1.assign(3); + x2.assign(1); + x3.assign(2); + grad.linspace(.1, .1); + + + sd::ops::mergemax_bp op; + auto result = op.evaluate({ &x1, &x2, &x3, &grad }, {}, {}); + ASSERT_EQ(Status::OK(), result.status()); + ASSERT_EQ(3, result.size()); + + auto z = result.at(0); + + ASSERT_TRUE(grad.isSameShape(z)); + ASSERT_TRUE(grad.equalsTo(z)); + +} +///////////////////////////////////////////////////////////////////////// +TEST_F(DeclarableOpsTests13, mergemax_bp_2) { + + NDArray x1('c', { 2, 5 }, { 1,2,3,4,5,4,3,2,1,0 }, sd::DataType::FLOAT32); + NDArray x2('c', { 2, 5 }, { 0,1,2,3,4,5,6,7,8,9 }, sd::DataType::FLOAT32); + NDArray x3('c', { 2, 5 }, { 0,1,1,2,3,4,7,5,8,10 }, sd::DataType::FLOAT32); + NDArray grad('c', { 2, 5 }, sd::DataType::FLOAT32); + + grad.linspace(.1, .1); + + NDArray exp1('c', { 2, 5 }, { 0.1, 0.2, 0.3, 0.4, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0 }, sd::DataType::FLOAT32); + NDArray exp2('c', { 2, 5 }, { 0.0, 0.0, 0.0, 0.0, 0.0, 0.6, 0.0, 0.8, 0.9, 0.0 }, sd::DataType::FLOAT32); + NDArray exp3('c', { 2, 5 }, { 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.7, 0.0, 0.0, 1.0 }, sd::DataType::FLOAT32); + + sd::ops::mergemax_bp op; + auto result = op.evaluate({ &x1, &x2, &x3, &grad }, {}, {}); + ASSERT_EQ(Status::OK(), result.status()); + ASSERT_EQ(3, result.size()); + + auto z1 = result.at(0); + auto z2 = result.at(1); + auto z3 = result.at(2); + + ASSERT_TRUE(exp1.isSameShape(z1)); + ASSERT_TRUE(exp1.equalsTo(z1)); + ASSERT_TRUE(exp2.isSameShape(z2)); + ASSERT_TRUE(exp2.equalsTo(z2)); + ASSERT_TRUE(exp3.isSameShape(z3)); + ASSERT_TRUE(exp3.equalsTo(z3)); + +} +///////////////////////////////////////////////////////////////////////// +TEST_F(DeclarableOpsTests13, mergemax_bp_3) { + + NDArray x1C('c', { 2, 5 }, { 1,2,3,4,5,4,3,2,1,0 }, sd::DataType::FLOAT32); + NDArray x2C('c', { 2, 5 }, { 0,1,2,3,4,5,6,7,8,9 }, sd::DataType::FLOAT32); + NDArray x3C('c', { 2, 5 }, { 0,1,1,2,3,4,7,5,8,10 }, sd::DataType::FLOAT32); + NDArray grad('c', { 2, 5 }, sd::DataType::FLOAT32); + + grad.linspace(.1, .1); + + NDArray x1('f', { 2, 5 }, sd::DataType::FLOAT32); + NDArray x2('f', { 2, 5 }, sd::DataType::FLOAT32); + NDArray x3('f', { 2, 5 }, sd::DataType::FLOAT32); + + NDArray exp1C('c', { 2, 5 }, { 0.1, 0.2, 0.3, 0.4, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0 }, sd::DataType::FLOAT32); + NDArray exp2C('c', { 2, 5 }, { 0.0, 0.0, 0.0, 0.0, 0.0, 0.6, 0.0, 0.8, 0.9, 0.0 }, sd::DataType::FLOAT32); + NDArray exp3C('c', { 2, 5 }, { 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.7, 0.0, 0.0, 1.0 }, sd::DataType::FLOAT32); + + NDArray exp1('f', { 2, 5 }, sd::DataType::FLOAT32); + NDArray exp2('f', { 2, 5 }, sd::DataType::FLOAT32); + NDArray exp3('f', { 2, 5 }, sd::DataType::FLOAT32); + + x1.assign(x1C); + x2.assign(x2C); + x3.assign(x3C); + + exp1.assign(exp1C); + exp2.assign(exp2C); + exp3.assign(exp3C); + + sd::ops::mergemax_bp op; + auto result = op.evaluate({ &x1, &x2, &x3, &grad }, {}, {}); + ASSERT_EQ(Status::OK(), result.status()); + ASSERT_EQ(3, result.size()); + + auto z1 = result.at(0); + auto z2 = result.at(1); + auto z3 = result.at(2); + + ASSERT_TRUE(exp1.isSameShape(z1)); + ASSERT_TRUE(exp1.equalsTo(z1)); + ASSERT_TRUE(exp2.isSameShape(z2)); + ASSERT_TRUE(exp2.equalsTo(z2)); + ASSERT_TRUE(exp3.isSameShape(z3)); + ASSERT_TRUE(exp3.equalsTo(z3)); + +} +///////////////////////////////////////////////////////////////////////// +TEST_F(DeclarableOpsTests13, mergeadd_bp_1) { + + NDArray x1('c', { 5, 5 }, sd::DataType::FLOAT32); + NDArray x2('c', { 5, 5 }, sd::DataType::FLOAT32); + NDArray x3('c', { 5, 5 }, sd::DataType::FLOAT32); + NDArray grad('c', { 5, 5 }, sd::DataType::FLOAT32); + + x1.assign(3); + x2.assign(1); + x3.assign(2); + grad.linspace(.1, .1); + + sd::ops::mergeadd_bp op; + auto result = op.evaluate({ &x1, &x2, &x3, &grad }, {}, {}); + ASSERT_EQ(Status::OK(), result.status()); + ASSERT_EQ(3, result.size()); + + for (int i = 0; i < 3; i++) { + auto z = result.at(0); + ASSERT_TRUE(grad.isSameShape(z)); + ASSERT_TRUE(grad.equalsTo(z)); + } +} +///////////////////////////////////////////////////////////////////////// +TEST_F(DeclarableOpsTests13, mergeavg_bp_1) { + + NDArray x1('c', { 5, 5 }, sd::DataType::FLOAT32); + NDArray x2('c', { 5, 5 }, sd::DataType::FLOAT32); + NDArray x3('c', { 5, 5 }, sd::DataType::FLOAT32); + NDArray grad('c', { 5, 5 }, sd::DataType::FLOAT32); + + x1.assign(3); + x2.assign(1); + x3.assign(2); + grad.linspace(.1, .1); + + sd::ops::mergeavg_bp op; + auto result = op.evaluate({ &x1, &x2, &x3, &grad }, {}, {}); + ASSERT_EQ(Status::OK(), result.status()); + ASSERT_EQ(3, result.size()); + + grad.applyScalar(sd::scalar::Divide, 3, grad); + + for (int i = 0; i < 3; i++) { + auto z = result.at(i); + ASSERT_TRUE(grad.isSameShape(z)); + ASSERT_TRUE(grad.equalsTo(z)); + } + +} /////////////////////////////////////////////////////////////////// TEST_F(DeclarableOpsTests13, lstmLayer_1) { From 9970aadc5a0f3ed69938211b153c9bae17d3ff04 Mon Sep 17 00:00:00 2001 From: Alex Black Date: Wed, 25 Mar 2020 19:42:08 +1100 Subject: [PATCH 4/6] Upgrade python version to 3.7.7 from 3.7.6 (#346) * Upgrade python version to 3.7.7 Signed-off-by: Alex Black * Numpy 1.18.1 -> 1.18.2 Signed-off-by: Alex Black --- pom.xml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/pom.xml b/pom.xml index 898e666a1..5a8d49d88 100644 --- a/pom.xml +++ b/pom.xml @@ -292,9 +292,9 @@ 1.5.3-SNAPSHOT 1.5.3-SNAPSHOT - 3.7.6 + 3.7.7 ${python.version}-${javacpp-presets.version} - 1.18.1 + 1.18.2 ${numpy.version}-${javacpp-presets.version} 0.3.9 From 63c9223bc2d95e3311e7f551aac829abf3693875 Mon Sep 17 00:00:00 2001 From: Alex Black Date: Fri, 27 Mar 2020 00:33:13 +1100 Subject: [PATCH 5/6] Assorted fixes (#318) * #8777 MultiLayerNetwork.evaluate(MultiDataSetIterator) overload Signed-off-by: Alex Black * #8768 SameDiff.equals Signed-off-by: Alex Black * #8750 shade freemarker library and switch to it in DL4J UI Signed-off-by: Alex Black * #8704 DL4J UI redirect Signed-off-by: Alex Black * #8776 RecordReaderDataSetIterator builder collectMetaData fix Signed-off-by: Alex Black * #8718 Fix DL4J doEvaluation metadata Signed-off-by: Alex Black * #8715 ArchiveUtils - Add option to not log every extracted file Signed-off-by: Alex Black * No exception for evaluations that don't support metadata Signed-off-by: Alex Black * Fixes Signed-off-by: Alex Black * #8765 CompGraph+MDS fix for SharedTrainingMaster Signed-off-by: Alex Black * small fix Signed-off-by: Alex Black * Timeout Signed-off-by: Alex Black * Ignore Signed-off-by: Alex Black * Revert freemarker shading Signed-off-by: Alex Black * Ignore Signed-off-by: Alex Black --- .../RecordReaderDataSetiteratorTest.java | 13 + .../org/deeplearning4j/eval/EvalTest.java | 55 ++-- .../eval/RegressionEvalTest.java | 2 +- .../CacheableExtractableDataSetFetcher.java | 2 +- .../datavec/RecordReaderDataSetIterator.java | 1 + .../util/KuromojiBinFilesFetcher.java | 2 +- .../nn/graph/ComputationGraph.java | 3 +- .../nn/graph/util/ComputationGraphUtil.java | 8 +- .../nn/multilayer/MultiLayerNetwork.java | 41 ++- .../threshold/AdaptiveThresholdAlgorithm.java | 5 +- .../spark/dl4j-spark-parameterserver/pom.xml | 16 +- .../VirtualMultiDataSetIterator.java | 22 +- .../pw/SharedTrainingWrapper.java | 12 +- .../spark/parameterserver/BaseSparkTest.java | 3 +- .../train/GradientSharingTrainingTest.java | 250 ++++++++++-------- .../src/test/resources/log4j.properties | 8 +- .../src/test/resources/logback.xml | 2 +- .../deeplearning4j-ui-components/pom.xml | 4 - .../org/deeplearning4j/ui/TestStandAlone.java | 1 + .../deeplearning4j-vertx/pom.xml | 2 +- .../ui/module/train/TrainModule.java | 1 + .../nd4j/autodiff/samediff/SDVariable.java | 45 ++-- .../org/nd4j/autodiff/samediff/SameDiff.java | 15 +- .../autodiff/samediff/internal/Variable.java | 6 +- .../classification/EvaluationBinary.java | 3 - .../classification/EvaluationCalibration.java | 2 +- .../regression/RegressionEvaluation.java | 2 +- .../nd4j/autodiff/samediff/SameDiffTests.java | 48 ++++ .../nd4j/linalg/nativ/OpsMappingTests.java | 2 +- .../java/org/nd4j/resources/Downloader.java | 2 +- .../main/java/org/nd4j/util/ArchiveUtils.java | 12 +- 31 files changed, 381 insertions(+), 209 deletions(-) diff --git a/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/datasets/datavec/RecordReaderDataSetiteratorTest.java b/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/datasets/datavec/RecordReaderDataSetiteratorTest.java index c20b5855f..2b7121af4 100644 --- a/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/datasets/datavec/RecordReaderDataSetiteratorTest.java +++ b/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/datasets/datavec/RecordReaderDataSetiteratorTest.java @@ -1381,4 +1381,17 @@ public class RecordReaderDataSetiteratorTest extends BaseDL4JTest { assertNotNull(ds.getFeatures()); assertNull(ds.getLabels()); } + + + @Test + public void testCollectMetaData(){ + RecordReaderDataSetIterator trainIter = new RecordReaderDataSetIterator.Builder(new CollectionRecordReader(Collections.>emptyList()), 1) + .collectMetaData(true) + .build(); + assertTrue(trainIter.isCollectMetaData()); + trainIter.setCollectMetaData(false); + assertFalse(trainIter.isCollectMetaData()); + trainIter.setCollectMetaData(true); + assertTrue(trainIter.isCollectMetaData()); + } } diff --git a/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/eval/EvalTest.java b/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/eval/EvalTest.java index 812ea2b08..bd65af6a3 100644 --- a/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/eval/EvalTest.java +++ b/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/eval/EvalTest.java @@ -33,7 +33,6 @@ import org.deeplearning4j.datasets.iterator.IteratorMultiDataSetIterator; import org.deeplearning4j.datasets.iterator.impl.IrisDataSetIterator; import org.deeplearning4j.datasets.iterator.impl.ListDataSetIterator; import org.deeplearning4j.datasets.iterator.impl.SingletonMultiDataSetIterator; -import org.deeplearning4j.eval.meta.Prediction; import org.deeplearning4j.nn.api.OptimizationAlgorithm; import org.deeplearning4j.nn.conf.*; import org.deeplearning4j.nn.conf.layers.*; @@ -52,19 +51,13 @@ import org.nd4j.linalg.dataset.api.MultiDataSet; import org.nd4j.linalg.dataset.api.iterator.DataSetIterator; import org.nd4j.linalg.dataset.api.preprocessor.NormalizerStandardize; import org.nd4j.linalg.factory.Nd4j; -import org.nd4j.linalg.indexing.INDArrayIndex; -import org.nd4j.linalg.indexing.NDArrayIndex; -import org.nd4j.linalg.io.ClassPathResource; import org.nd4j.linalg.learning.config.Sgd; import org.nd4j.linalg.lossfunctions.LossFunctions; -import org.nd4j.linalg.util.FeatureUtil; import org.nd4j.resources.Resources; import java.util.*; import static org.junit.Assert.*; -import static org.nd4j.linalg.indexing.NDArrayIndex.all; -import static org.nd4j.linalg.indexing.NDArrayIndex.interval; /** * Created by agibsonccc on 12/22/14. @@ -165,7 +158,7 @@ public class EvalTest extends BaseDL4JTest { assertEquals(evalExpected.getConfusionMatrix(), evalActual.getConfusionMatrix()); } - @Test(timeout = 300000) + @Test public void testEvaluationWithMetaData() throws Exception { RecordReader csv = new CSVRecordReader(); @@ -256,6 +249,30 @@ public class EvalTest extends BaseDL4JTest { assertEquals(actualCounts[i], actualClassI.size()); assertEquals(predictedCounts[i], predictedClassI.size()); } + + + //Finally: test doEvaluation methods + rrdsi.reset(); + org.nd4j.evaluation.classification.Evaluation e2 = new org.nd4j.evaluation.classification.Evaluation(); + net.doEvaluation(rrdsi, e2); + for (int i = 0; i < 3; i++) { + List actualClassI = e2.getPredictionsByActualClass(i); + List predictedClassI = e2.getPredictionByPredictedClass(i); + assertEquals(actualCounts[i], actualClassI.size()); + assertEquals(predictedCounts[i], predictedClassI.size()); + } + + ComputationGraph cg = net.toComputationGraph(); + rrdsi.reset(); + e2 = new org.nd4j.evaluation.classification.Evaluation(); + cg.doEvaluation(rrdsi, e2); + for (int i = 0; i < 3; i++) { + List actualClassI = e2.getPredictionsByActualClass(i); + List predictedClassI = e2.getPredictionByPredictedClass(i); + assertEquals(actualCounts[i], actualClassI.size()); + assertEquals(predictedCounts[i], predictedClassI.size()); + } + } private static void apply(org.nd4j.evaluation.classification.Evaluation e, int nTimes, INDArray predicted, INDArray actual) { @@ -504,11 +521,11 @@ public class EvalTest extends BaseDL4JTest { list.add(new org.nd4j.linalg.dataset.MultiDataSet(new INDArray[]{ds.getFeatures()}, new INDArray[]{ds.getLabels(), ds.getLabels()})); } - Evaluation e = new Evaluation(); - RegressionEvaluation e2 = new RegressionEvaluation(); - Map evals = new HashMap<>(); - evals.put(0, new IEvaluation[]{(IEvaluation) e}); - evals.put(1, new IEvaluation[]{(IEvaluation) e2}); + org.nd4j.evaluation.classification.Evaluation e = new org.nd4j.evaluation.classification.Evaluation(); + org.nd4j.evaluation.regression.RegressionEvaluation e2 = new org.nd4j.evaluation.regression.RegressionEvaluation(); + Map evals = new HashMap<>(); + evals.put(0, new org.nd4j.evaluation.IEvaluation[]{e}); + evals.put(1, new org.nd4j.evaluation.IEvaluation[]{e2}); cg.evaluate(new IteratorMultiDataSetIterator(list.iterator(), 30), evals); @@ -567,14 +584,14 @@ public class EvalTest extends BaseDL4JTest { } try { - net.evaluateROC(iter); + net.evaluateROC(iter, 0); fail("Expected exception"); } catch (IllegalStateException e){ assertTrue(e.getMessage().contains("Classifier") && e.getMessage().contains("ROC")); } try { - net.evaluateROCMultiClass(iter); + net.evaluateROCMultiClass(iter, 0); fail("Expected exception"); } catch (IllegalStateException e){ assertTrue(e.getMessage().contains("Classifier") && e.getMessage().contains("ROCMultiClass")); @@ -589,14 +606,14 @@ public class EvalTest extends BaseDL4JTest { } try { - cg.evaluateROC(iter); + cg.evaluateROC(iter, 0); fail("Expected exception"); } catch (IllegalStateException e){ assertTrue(e.getMessage().contains("Classifier") && e.getMessage().contains("ROC")); } try { - cg.evaluateROCMultiClass(iter); + cg.evaluateROCMultiClass(iter, 0); fail("Expected exception"); } catch (IllegalStateException e){ assertTrue(e.getMessage().contains("Classifier") && e.getMessage().contains("ROCMultiClass")); @@ -606,10 +623,10 @@ public class EvalTest extends BaseDL4JTest { //Disable validation, and check same thing: net.getLayerWiseConfigurations().setValidateOutputLayerConfig(false); net.evaluate(iter); - net.evaluateROCMultiClass(iter); + net.evaluateROCMultiClass(iter, 0); cg.getConfiguration().setValidateOutputLayerConfig(false); cg.evaluate(iter); - cg.evaluateROCMultiClass(iter); + cg.evaluateROCMultiClass(iter, 0); } } diff --git a/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/eval/RegressionEvalTest.java b/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/eval/RegressionEvalTest.java index 7df75f6da..ba469546d 100644 --- a/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/eval/RegressionEvalTest.java +++ b/deeplearning4j/deeplearning4j-core/src/test/java/org/deeplearning4j/eval/RegressionEvalTest.java @@ -61,7 +61,7 @@ public class RegressionEvalTest extends BaseDL4JTest { DataSet ds = new DataSet(f, l); DataSetIterator iter = new ExistingDataSetIterator(Collections.singletonList(ds)); - RegressionEvaluation re = net.evaluateRegression(iter); + org.nd4j.evaluation.regression.RegressionEvaluation re = net.evaluateRegression(iter); for (int i = 0; i < 5; i++) { assertEquals(1.0, re.meanSquaredError(i), 1e-6); diff --git a/deeplearning4j/deeplearning4j-data/deeplearning4j-datasets/src/main/java/org/deeplearning4j/datasets/fetchers/CacheableExtractableDataSetFetcher.java b/deeplearning4j/deeplearning4j-data/deeplearning4j-datasets/src/main/java/org/deeplearning4j/datasets/fetchers/CacheableExtractableDataSetFetcher.java index 97574a99f..4a8d01aa4 100644 --- a/deeplearning4j/deeplearning4j-data/deeplearning4j-datasets/src/main/java/org/deeplearning4j/datasets/fetchers/CacheableExtractableDataSetFetcher.java +++ b/deeplearning4j/deeplearning4j-data/deeplearning4j-datasets/src/main/java/org/deeplearning4j/datasets/fetchers/CacheableExtractableDataSetFetcher.java @@ -86,7 +86,7 @@ public abstract class CacheableExtractableDataSetFetcher implements CacheableDat } try { - ArchiveUtils.unzipFileTo(tmpFile.getAbsolutePath(), localCacheDir.getAbsolutePath()); + ArchiveUtils.unzipFileTo(tmpFile.getAbsolutePath(), localCacheDir.getAbsolutePath(), false); } catch (Throwable t){ //Catch any errors during extraction, and delete the directory to avoid leaving the dir in an invalid state if(localCacheDir.exists()) diff --git a/deeplearning4j/deeplearning4j-data/deeplearning4j-datavec-iterators/src/main/java/org/deeplearning4j/datasets/datavec/RecordReaderDataSetIterator.java b/deeplearning4j/deeplearning4j-data/deeplearning4j-datavec-iterators/src/main/java/org/deeplearning4j/datasets/datavec/RecordReaderDataSetIterator.java index bda0f9c95..9f7813d5c 100644 --- a/deeplearning4j/deeplearning4j-data/deeplearning4j-datavec-iterators/src/main/java/org/deeplearning4j/datasets/datavec/RecordReaderDataSetIterator.java +++ b/deeplearning4j/deeplearning4j-data/deeplearning4j-datavec-iterators/src/main/java/org/deeplearning4j/datasets/datavec/RecordReaderDataSetIterator.java @@ -205,6 +205,7 @@ public class RecordReaderDataSetIterator implements DataSetIterator { this.numPossibleLabels = b.numPossibleLabels; this.regression = b.regression; this.preProcessor = b.preProcessor; + this.collectMetaData = b.collectMetaData; } /** diff --git a/deeplearning4j/deeplearning4j-nlp-parent/deeplearning4j-nlp-japanese/src/main/java/com/atilika/kuromoji/util/KuromojiBinFilesFetcher.java b/deeplearning4j/deeplearning4j-nlp-parent/deeplearning4j-nlp-japanese/src/main/java/com/atilika/kuromoji/util/KuromojiBinFilesFetcher.java index adcd87b5b..d2945cf3d 100644 --- a/deeplearning4j/deeplearning4j-nlp-parent/deeplearning4j-nlp-japanese/src/main/java/com/atilika/kuromoji/util/KuromojiBinFilesFetcher.java +++ b/deeplearning4j/deeplearning4j-nlp-parent/deeplearning4j-nlp-japanese/src/main/java/com/atilika/kuromoji/util/KuromojiBinFilesFetcher.java @@ -67,7 +67,7 @@ public class KuromojiBinFilesFetcher { new URL("https://dl4jdata.blob.core.windows.net/kuromoji/kuromoji_bin_files.tar.gz"), tarFile); } - ArchiveUtils.unzipFileTo(tarFile.getAbsolutePath(), rootDir.getAbsolutePath()); + ArchiveUtils.unzipFileTo(tarFile.getAbsolutePath(), rootDir.getAbsolutePath(), false); return rootDir.getAbsoluteFile(); } diff --git a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/graph/ComputationGraph.java b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/graph/ComputationGraph.java index 0a34fe95a..582b20a15 100755 --- a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/graph/ComputationGraph.java +++ b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/graph/ComputationGraph.java @@ -4170,6 +4170,7 @@ public class ComputationGraph implements Serializable, Model, NeuralNetwork { INDArray[] featuresMasks = next.getFeaturesMaskArrays(); INDArray[] labels = next.getLabels(); INDArray[] labelMasks = next.getLabelsMaskArrays(); + List meta = next.getExampleMetaData(); try (MemoryWorkspace ws = outputWs.notifyScopeEntered()) { INDArray[] out = outputOfLayersDetached(false, FwdPassType.STANDARD, getOutputLayerIndices(), features, featuresMasks, labelMasks, true, false, ws); @@ -4188,7 +4189,7 @@ public class ComputationGraph implements Serializable, Model, NeuralNetwork { try (MemoryWorkspace wsO = Nd4j.getWorkspaceManager().scopeOutOfWorkspaces()) { for (IEvaluation evaluation : evalsThisOutput) - evaluation.eval(currLabel, currOut, next.getLabelsMaskArray(i)); + evaluation.eval(currLabel, currOut, next.getLabelsMaskArray(i), meta); } } } diff --git a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/graph/util/ComputationGraphUtil.java b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/graph/util/ComputationGraphUtil.java index a27ce9a4c..4b9918203 100644 --- a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/graph/util/ComputationGraphUtil.java +++ b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/graph/util/ComputationGraphUtil.java @@ -23,6 +23,9 @@ import org.nd4j.linalg.dataset.api.MultiDataSet; import org.nd4j.linalg.dataset.api.iterator.DataSetIterator; import org.nd4j.linalg.dataset.api.iterator.MultiDataSetIterator; +import java.io.Serializable; +import java.util.List; + public class ComputationGraphUtil { private ComputationGraphUtil() {} @@ -33,13 +36,16 @@ public class ComputationGraphUtil { INDArray l = dataSet.getLabels(); INDArray fMask = dataSet.getFeaturesMaskArray(); INDArray lMask = dataSet.getLabelsMaskArray(); + List meta = dataSet.getExampleMetaData(); INDArray[] fNew = f == null ? null : new INDArray[] {f}; INDArray[] lNew = l == null ? null : new INDArray[] {l}; INDArray[] fMaskNew = (fMask != null ? new INDArray[] {fMask} : null); INDArray[] lMaskNew = (lMask != null ? new INDArray[] {lMask} : null); - return new org.nd4j.linalg.dataset.MultiDataSet(fNew, lNew, fMaskNew, lMaskNew); + org.nd4j.linalg.dataset.MultiDataSet mds = new org.nd4j.linalg.dataset.MultiDataSet(fNew, lNew, fMaskNew, lMaskNew); + mds.setExampleMetaData(meta); + return mds; } /** Convert a DataSetIterator to a MultiDataSetIterator, via an adaptor class */ diff --git a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/multilayer/MultiLayerNetwork.java b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/multilayer/MultiLayerNetwork.java index bce86b9ce..5cc536810 100755 --- a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/multilayer/MultiLayerNetwork.java +++ b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/nn/multilayer/MultiLayerNetwork.java @@ -25,14 +25,11 @@ import lombok.val; import org.apache.commons.lang3.ArrayUtils; import org.apache.commons.lang3.StringUtils; import org.bytedeco.javacpp.Pointer; -import org.nd4j.adapters.OutputAdapter; -import org.nd4j.linalg.dataset.AsyncDataSetIterator;; import org.deeplearning4j.datasets.iterator.MultiDataSetWrapperIterator; -import org.deeplearning4j.eval.RegressionEvaluation; import org.deeplearning4j.exception.DL4JException; import org.deeplearning4j.exception.DL4JInvalidInputException; -import org.deeplearning4j.nn.api.*; import org.deeplearning4j.nn.api.Updater; +import org.deeplearning4j.nn.api.*; import org.deeplearning4j.nn.api.layers.IOutputLayer; import org.deeplearning4j.nn.api.layers.RecurrentLayer; import org.deeplearning4j.nn.conf.*; @@ -44,8 +41,8 @@ import org.deeplearning4j.nn.gradient.Gradient; import org.deeplearning4j.nn.graph.ComputationGraph; import org.deeplearning4j.nn.layers.FrozenLayer; import org.deeplearning4j.nn.layers.FrozenLayerWithBackprop; -import org.deeplearning4j.nn.layers.recurrent.BidirectionalLayer; import org.deeplearning4j.nn.layers.LayerHelper; +import org.deeplearning4j.nn.layers.recurrent.BidirectionalLayer; import org.deeplearning4j.nn.layers.wrapper.BaseWrapperLayer; import org.deeplearning4j.nn.updater.UpdaterCreator; import org.deeplearning4j.nn.workspace.ArrayType; @@ -58,19 +55,23 @@ import org.deeplearning4j.util.CrashReportingUtil; import org.deeplearning4j.util.ModelSerializer; import org.deeplearning4j.util.NetworkUtils; import org.deeplearning4j.util.OutputLayerUtil; +import org.nd4j.adapters.OutputAdapter; import org.nd4j.base.Preconditions; import org.nd4j.evaluation.IEvaluation; import org.nd4j.evaluation.classification.Evaluation; import org.nd4j.evaluation.classification.ROC; import org.nd4j.evaluation.classification.ROCMultiClass; +import org.nd4j.evaluation.regression.RegressionEvaluation; import org.nd4j.linalg.api.buffer.DataType; import org.nd4j.linalg.api.memory.MemoryWorkspace; +import org.nd4j.linalg.api.memory.abstracts.DummyWorkspace; import org.nd4j.linalg.api.memory.conf.WorkspaceConfiguration; import org.nd4j.linalg.api.memory.enums.AllocationPolicy; import org.nd4j.linalg.api.memory.enums.LearningPolicy; import org.nd4j.linalg.api.memory.enums.ResetPolicy; import org.nd4j.linalg.api.memory.enums.SpillPolicy; import org.nd4j.linalg.api.ndarray.INDArray; +import org.nd4j.linalg.dataset.AsyncDataSetIterator; import org.nd4j.linalg.dataset.DataSet; import org.nd4j.linalg.dataset.api.MultiDataSet; import org.nd4j.linalg.dataset.api.iterator.DataSetIterator; @@ -84,7 +85,6 @@ import org.nd4j.linalg.heartbeat.reports.Task; import org.nd4j.linalg.heartbeat.utils.EnvironmentUtils; import org.nd4j.linalg.heartbeat.utils.TaskUtils; import org.nd4j.linalg.indexing.NDArrayIndex; -import org.nd4j.linalg.api.memory.abstracts.DummyWorkspace; import org.nd4j.linalg.primitives.Pair; import org.nd4j.linalg.primitives.Triple; import org.nd4j.linalg.schedule.ISchedule; @@ -96,6 +96,8 @@ import org.nd4j.util.OneTimeLogger; import java.io.*; import java.util.*; +; + /** * MultiLayerNetwork is a neural network with multiple layers in a stack, and usually an output layer.
@@ -3315,19 +3317,39 @@ public class MultiLayerNetwork implements Serializable, Classifier, Layer, Neura * @param iterator Iterator to evaluate on * @return Evaluation object; results of evaluation on all examples in the data set */ - public T evaluate(DataSetIterator iterator) { + public T evaluate(@NonNull DataSetIterator iterator) { return (T)evaluate(iterator, null); } + /** + * Evaluate the network (classification performance). + * Can only be used with MultiDataSetIterator instances with a single input/output array + * + * @param iterator Iterator to evaluate on + * @return Evaluation object; results of evaluation on all examples in the data set + */ + public Evaluation evaluate(@NonNull MultiDataSetIterator iterator) { + return evaluate(new MultiDataSetWrapperIterator(iterator)); + } + /** * Evaluate the network for regression performance * @param iterator Data to evaluate on - * @return + * @return Regression evaluation */ public T evaluateRegression(DataSetIterator iterator) { return (T)doEvaluation(iterator, new RegressionEvaluation(iterator.totalOutcomes()))[0]; } + /** + * Evaluate the network for regression performance + * Can only be used with MultiDataSetIterator instances with a single input/output array + * @param iterator Data to evaluate on + */ + public org.nd4j.evaluation.regression.RegressionEvaluation evaluateRegression(MultiDataSetIterator iterator) { + return evaluateRegression(new MultiDataSetWrapperIterator(iterator)); + } + /** * @deprecated To be removed - use {@link #evaluateROC(DataSetIterator, int)} to enforce selection of appropriate ROC/threshold configuration */ @@ -3424,6 +3446,7 @@ public class MultiLayerNetwork implements Serializable, Classifier, Layer, Neura INDArray labels = next.getLabels(); INDArray fMask = next.getFeaturesMaskArray(); INDArray lMask = next.getLabelsMaskArray(); + List meta = next.getExampleMetaData(); if (!useRnnSegments) { @@ -3433,7 +3456,7 @@ public class MultiLayerNetwork implements Serializable, Classifier, Layer, Neura try (MemoryWorkspace wsO = Nd4j.getWorkspaceManager().scopeOutOfWorkspaces()) { for (T evaluation : evaluations) - evaluation.eval(labels, out, lMask); + evaluation.eval(labels, out, lMask, meta); } } } else { diff --git a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/optimize/solvers/accumulation/encoding/threshold/AdaptiveThresholdAlgorithm.java b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/optimize/solvers/accumulation/encoding/threshold/AdaptiveThresholdAlgorithm.java index 299732287..7dd56815c 100644 --- a/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/optimize/solvers/accumulation/encoding/threshold/AdaptiveThresholdAlgorithm.java +++ b/deeplearning4j/deeplearning4j-nn/src/main/java/org/deeplearning4j/optimize/solvers/accumulation/encoding/threshold/AdaptiveThresholdAlgorithm.java @@ -222,8 +222,11 @@ public class AdaptiveThresholdAlgorithm implements ThresholdAlgorithm { if(a == null || Double.isNaN(a.lastThreshold)) return; + lastThresholdSum += a.lastThreshold; - lastSparsitySum += a.lastSparsity; + if (!Double.isNaN(a.lastSparsity)) { + lastSparsitySum += a.lastSparsity; + } count++; } diff --git a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/pom.xml b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/pom.xml index daf0dd9b7..1198ae733 100644 --- a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/pom.xml +++ b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/pom.xml @@ -38,16 +38,22 @@ nd4j-aeron ${nd4j.version} - - org.nd4j - nd4j-parameter-server-node_2.11 - ${nd4j.version} - org.deeplearning4j dl4j-spark_2.11 ${project.version} + + org.nd4j + nd4j-parameter-server-node_2.11 + ${nd4j.version} + + + net.jpountz.lz4 + lz4 + + + org.projectlombok lombok diff --git a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/main/java/org/deeplearning4j/spark/parameterserver/iterators/VirtualMultiDataSetIterator.java b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/main/java/org/deeplearning4j/spark/parameterserver/iterators/VirtualMultiDataSetIterator.java index 1de2d8636..a3c3b43a8 100644 --- a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/main/java/org/deeplearning4j/spark/parameterserver/iterators/VirtualMultiDataSetIterator.java +++ b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/main/java/org/deeplearning4j/spark/parameterserver/iterators/VirtualMultiDataSetIterator.java @@ -23,6 +23,7 @@ import org.nd4j.linalg.dataset.api.iterator.ParallelMultiDataSetIterator; import java.util.Iterator; import java.util.List; +import java.util.concurrent.atomic.AtomicInteger; /** * This MultiDataSetIterator implementation does accumulation of MultiDataSets from different Spark executors, wrt Thread/Device Affinity @@ -32,14 +33,16 @@ import java.util.List; public class VirtualMultiDataSetIterator implements ParallelMultiDataSetIterator { protected final List> iterators; + protected final AtomicInteger position; public VirtualMultiDataSetIterator(@NonNull List> iterators) { this.iterators = iterators; + this.position = new AtomicInteger(0); } @Override public MultiDataSet next(int num) { - return null; + return next(); } @Override @@ -59,27 +62,34 @@ public class VirtualMultiDataSetIterator implements ParallelMultiDataSetIterator @Override public boolean asyncSupported() { - return false; + return true; } @Override public void reset() { - + throw new UnsupportedOperationException(); } @Override public boolean hasNext() { - return false; + // just checking if that's not the last iterator, or if that's the last one - check if it has something + boolean ret = position.get() < iterators.size() - 1 + || (position.get() < iterators.size() && iterators.get(position.get()).hasNext()); + return ret; } @Override public MultiDataSet next() { - return null; + // TODO: this solution isn't ideal, it assumes non-empty iterators all the time. Would be nice to do something here + if (!iterators.get(position.get()).hasNext()) + position.getAndIncrement(); + + return iterators.get(position.get()).next(); } @Override public void remove() { - + // no-op } @Override diff --git a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/main/java/org/deeplearning4j/spark/parameterserver/pw/SharedTrainingWrapper.java b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/main/java/org/deeplearning4j/spark/parameterserver/pw/SharedTrainingWrapper.java index 81fce7fbf..b6a8bb81c 100644 --- a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/main/java/org/deeplearning4j/spark/parameterserver/pw/SharedTrainingWrapper.java +++ b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/main/java/org/deeplearning4j/spark/parameterserver/pw/SharedTrainingWrapper.java @@ -109,6 +109,7 @@ public class SharedTrainingWrapper { // now we're creating DataSetIterators, to feed ParallelWrapper iteratorDS = new VirtualDataSetIterator(iteratorsDS); + iteratorMDS = new VirtualMultiDataSetIterator(iteratorsMDS); } public static synchronized SharedTrainingWrapper getInstance(long id) { @@ -447,17 +448,19 @@ public class SharedTrainingWrapper { throw new DL4JInvalidConfigException("No iterators were defined for training"); try { - while((iteratorDS != null && iteratorDS.hasNext()) || (iteratorMDS != null && iteratorMDS.hasNext())) { + boolean dsNext; + boolean mdsNext; + while((dsNext = iteratorDS != null && iteratorDS.hasNext()) || (mdsNext = iteratorMDS != null && iteratorMDS.hasNext())) { //Loop as a guard against concurrent modifications and RCs if (wrapper != null) { - if (iteratorDS != null) + if (dsNext) wrapper.fit(iteratorDS); else wrapper.fit(iteratorMDS); } else { // if wrapper is null, we're fitting standalone model then - if (iteratorDS != null) { + if (dsNext) { if (model instanceof ComputationGraph) { ((ComputationGraph) originalModel).fit(iteratorDS); } else if (model instanceof MultiLayerNetwork) { @@ -472,7 +475,8 @@ public class SharedTrainingWrapper { } } - consumer.getUpdatesQueue().purge(); + if(consumer != null) + consumer.getUpdatesQueue().purge(); } } catch (Throwable t){ log.warn("Exception encountered during fit operation", t); diff --git a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/java/org/deeplearning4j/spark/parameterserver/BaseSparkTest.java b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/java/org/deeplearning4j/spark/parameterserver/BaseSparkTest.java index c97292a2c..50aa564c1 100644 --- a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/java/org/deeplearning4j/spark/parameterserver/BaseSparkTest.java +++ b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/java/org/deeplearning4j/spark/parameterserver/BaseSparkTest.java @@ -116,8 +116,7 @@ public abstract class BaseSparkTest extends BaseDL4JTest implements Serializable } protected int numExecutors() { - int numProc = Runtime.getRuntime().availableProcessors(); - return Math.min(4, numProc); + return 4; } protected MultiLayerConfiguration getBasicConf() { diff --git a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/java/org/deeplearning4j/spark/parameterserver/train/GradientSharingTrainingTest.java b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/java/org/deeplearning4j/spark/parameterserver/train/GradientSharingTrainingTest.java index 53a4b32b1..ab034604e 100644 --- a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/java/org/deeplearning4j/spark/parameterserver/train/GradientSharingTrainingTest.java +++ b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/java/org/deeplearning4j/spark/parameterserver/train/GradientSharingTrainingTest.java @@ -49,6 +49,7 @@ import org.junit.rules.TemporaryFolder; import org.nd4j.linalg.activations.Activation; import org.nd4j.linalg.api.ndarray.INDArray; import org.nd4j.linalg.dataset.DataSet; +import org.nd4j.linalg.dataset.api.MultiDataSet; import org.nd4j.linalg.dataset.api.iterator.DataSetIterator; import org.nd4j.linalg.indexing.NDArrayIndex; import org.nd4j.linalg.learning.config.AMSGrad; @@ -66,137 +67,170 @@ import java.util.concurrent.ConcurrentHashMap; import static org.junit.Assert.*; @Slf4j -@Ignore("AB 2019/05/21 - Failing - Issue #7657") +//@Ignore("AB 2019/05/21 - Failing - Issue #7657") public class GradientSharingTrainingTest extends BaseSparkTest { @Rule public TemporaryFolder testDir = new TemporaryFolder(); + @Override + public long getTimeoutMilliseconds() { + return 90000L; + } + @Test public void trainSanityCheck() throws Exception { - INDArray last = null; - INDArray lastDup = null; - for (String s : new String[]{"paths", "direct", "export"}) { - System.out.println("--------------------------------------------------------------------------------------------------------------"); - log.info("Starting: {}", s); - boolean isPaths = "paths".equals(s); + for(boolean mds : new boolean[]{false, true}) { + INDArray last = null; + INDArray lastDup = null; + for (String s : new String[]{"paths", "direct", "export"}) { + System.out.println("--------------------------------------------------------------------------------------------------------------"); + log.info("Starting: {} - {}", s, (mds ? "MultiDataSet" : "DataSet")); + boolean isPaths = "paths".equals(s); - RDDTrainingApproach rddTrainingApproach; - switch (s) { - case "direct": - rddTrainingApproach = RDDTrainingApproach.Direct; - break; - case "export": - rddTrainingApproach = RDDTrainingApproach.Export; - break; - case "paths": - rddTrainingApproach = RDDTrainingApproach.Direct; //Actualy not used for fitPaths - break; - default: - throw new RuntimeException(); - } - - File temp = testDir.newFolder(); - - - //TODO this probably won't work everywhere... - String controller = Inet4Address.getLocalHost().getHostAddress(); - String networkMask = controller.substring(0, controller.lastIndexOf('.')) + ".0" + "/16"; - - VoidConfiguration voidConfiguration = VoidConfiguration.builder() - .unicastPort(40123) // Should be open for IN/OUT communications on all Spark nodes - .networkMask(networkMask) // Local network mask - .controllerAddress(controller) - .meshBuildMode(MeshBuildMode.PLAIN) // everyone is connected to the master - .build(); - TrainingMaster tm = new SharedTrainingMaster.Builder(voidConfiguration, 2, new AdaptiveThresholdAlgorithm(1e-3), 16) - .rngSeed(12345) - .collectTrainingStats(false) - .batchSizePerWorker(16) // Minibatch size for each worker - .workersPerNode(2) // Workers per node - .rddTrainingApproach(rddTrainingApproach) - .exportDirectory("file:///" + temp.getAbsolutePath().replaceAll("\\\\", "/")) - .build(); - - - ComputationGraphConfiguration conf = new NeuralNetConfiguration.Builder() - .seed(12345) - .updater(new AMSGrad(0.1)) - .graphBuilder() - .addInputs("in") - .layer("out", new OutputLayer.Builder().nIn(784).nOut(10).activation(Activation.SOFTMAX) - .lossFunction(LossFunctions.LossFunction.MCXENT).build(), "in") - .setOutputs("out") - .build(); - - - SparkComputationGraph sparkNet = new SparkComputationGraph(sc, conf, tm); - sparkNet.setCollectTrainingStats(tm.getIsCollectTrainingStats()); - - System.out.println(Arrays.toString(sparkNet.getNetwork().params().get(NDArrayIndex.point(0), NDArrayIndex.interval(0, 256)).dup().data().asFloat())); - File f = testDir.newFolder(); - DataSetIterator iter = new MnistDataSetIterator(16, true, 12345); - int count = 0; - List paths = new ArrayList<>(); - List ds = new ArrayList<>(); - while (iter.hasNext() && count++ < 8) { - DataSet d = iter.next(); - if (isPaths) { - File out = new File(f, count + ".bin"); - d.save(out); - String path = "file:///" + out.getAbsolutePath().replaceAll("\\\\", "/"); - paths.add(path); - } - ds.add(d); - } - - int numIter = 1; - double[] acc = new double[numIter + 1]; - for (int i = 0; i < numIter; i++) { - //Check accuracy before: - DataSetIterator testIter = new EarlyTerminationDataSetIterator(new MnistDataSetIterator(32, false, 12345), 10); - Evaluation eBefore = sparkNet.getNetwork().evaluate(testIter); - - INDArray paramsBefore = sparkNet.getNetwork().params().dup(); - ComputationGraph after; + RDDTrainingApproach rddTrainingApproach; switch (s) { case "direct": + rddTrainingApproach = RDDTrainingApproach.Direct; + break; case "export": - JavaRDD dsRDD = sc.parallelize(ds); - after = sparkNet.fit(dsRDD); + rddTrainingApproach = RDDTrainingApproach.Export; break; case "paths": - JavaRDD pathRdd = sc.parallelize(paths); - after = sparkNet.fitPaths(pathRdd); + rddTrainingApproach = RDDTrainingApproach.Direct; //Actualy not used for fitPaths break; default: throw new RuntimeException(); } - INDArray paramsAfter = after.params(); - System.out.println(Arrays.toString(paramsBefore.get(NDArrayIndex.point(0), NDArrayIndex.interval(0, 256)).dup().data().asFloat())); - System.out.println(Arrays.toString(paramsAfter.get(NDArrayIndex.point(0), NDArrayIndex.interval(0, 256)).dup().data().asFloat())); - System.out.println(Arrays.toString( - Transforms.abs(paramsAfter.sub(paramsBefore)).get(NDArrayIndex.point(0), NDArrayIndex.interval(0, 256)).dup().data().asFloat())); - assertNotEquals(paramsBefore, paramsAfter); + File temp = testDir.newFolder(); - testIter = new EarlyTerminationDataSetIterator(new MnistDataSetIterator(32, false, 12345), 10); - Evaluation eAfter = after.evaluate(testIter); + //TODO this probably won't work everywhere... + String controller = Inet4Address.getLocalHost().getHostAddress(); + String networkMask = controller.substring(0, controller.lastIndexOf('.')) + ".0" + "/16"; - double accAfter = eAfter.accuracy(); - double accBefore = eBefore.accuracy(); - assertTrue("after: " + accAfter + ", before=" + accBefore, accAfter >= accBefore + 0.005); + VoidConfiguration voidConfiguration = VoidConfiguration.builder() + .unicastPort(40123) // Should be open for IN/OUT communications on all Spark nodes + .networkMask(networkMask) // Local network mask + .controllerAddress(controller) + .meshBuildMode(MeshBuildMode.PLAIN) // everyone is connected to the master + .build(); + TrainingMaster tm = new SharedTrainingMaster.Builder(voidConfiguration, 2, new AdaptiveThresholdAlgorithm(1e-3), 16) + .rngSeed(12345) + .collectTrainingStats(false) + .batchSizePerWorker(16) // Minibatch size for each worker + .workersPerNode(2) // Workers per node + .rddTrainingApproach(rddTrainingApproach) + .exportDirectory("file:///" + temp.getAbsolutePath().replaceAll("\\\\", "/")) + .build(); - if (i == 0) { - acc[0] = eBefore.accuracy(); + + ComputationGraphConfiguration conf = new NeuralNetConfiguration.Builder() + .seed(12345) + .updater(new AMSGrad(0.1)) + .graphBuilder() + .addInputs("in") + .layer("out", new OutputLayer.Builder().nIn(784).nOut(10).activation(Activation.SOFTMAX) + .lossFunction(LossFunctions.LossFunction.MCXENT).build(), "in") + .setOutputs("out") + .build(); + + + SparkComputationGraph sparkNet = new SparkComputationGraph(sc, conf, tm); + sparkNet.setCollectTrainingStats(tm.getIsCollectTrainingStats()); + + System.out.println(Arrays.toString(sparkNet.getNetwork().params().get(NDArrayIndex.point(0), NDArrayIndex.interval(0, 256)).dup().data().asFloat())); + File f = testDir.newFolder(); + DataSetIterator iter = new MnistDataSetIterator(16, true, 12345); + int count = 0; + List paths = new ArrayList<>(); + List ds = new ArrayList<>(); + while (iter.hasNext() && count++ < 8) { + DataSet d = iter.next(); + if (isPaths) { + File out = new File(f, count + ".bin"); + if(mds){ + d.toMultiDataSet().save(out); + } else { + d.save(out); + } + String path = "file:///" + out.getAbsolutePath().replaceAll("\\\\", "/"); + paths.add(path); + } + ds.add(d); } - acc[i + 1] = eAfter.accuracy(); + + int numIter = 1; + double[] acc = new double[numIter + 1]; + for (int i = 0; i < numIter; i++) { + //Check accuracy before: + DataSetIterator testIter = new EarlyTerminationDataSetIterator(new MnistDataSetIterator(32, false, 12345), 10); + Evaluation eBefore = sparkNet.getNetwork().evaluate(testIter); + + INDArray paramsBefore = sparkNet.getNetwork().params().dup(); + ComputationGraph after; + if(mds) { + //Fitting from MultiDataSet + List mdsList = new ArrayList<>(); + for(DataSet d : ds){ + mdsList.add(d.toMultiDataSet()); + } + switch (s) { + case "direct": + case "export": + JavaRDD dsRDD = sc.parallelize(mdsList); + after = sparkNet.fitMultiDataSet(dsRDD); + break; + case "paths": + JavaRDD pathRdd = sc.parallelize(paths); + after = sparkNet.fitPathsMultiDataSet(pathRdd); + break; + default: + throw new RuntimeException(); + } + } else { + //Fitting from DataSet + switch (s) { + case "direct": + case "export": + JavaRDD dsRDD = sc.parallelize(ds); + after = sparkNet.fit(dsRDD); + break; + case "paths": + JavaRDD pathRdd = sc.parallelize(paths); + after = sparkNet.fitPaths(pathRdd); + break; + default: + throw new RuntimeException(); + } + } + + INDArray paramsAfter = after.params(); + System.out.println(Arrays.toString(paramsBefore.get(NDArrayIndex.point(0), NDArrayIndex.interval(0, 256)).dup().data().asFloat())); + System.out.println(Arrays.toString(paramsAfter.get(NDArrayIndex.point(0), NDArrayIndex.interval(0, 256)).dup().data().asFloat())); + System.out.println(Arrays.toString( + Transforms.abs(paramsAfter.sub(paramsBefore)).get(NDArrayIndex.point(0), NDArrayIndex.interval(0, 256)).dup().data().asFloat())); + assertNotEquals(paramsBefore, paramsAfter); + + + testIter = new EarlyTerminationDataSetIterator(new MnistDataSetIterator(32, false, 12345), 10); + Evaluation eAfter = after.evaluate(testIter); + + double accAfter = eAfter.accuracy(); + double accBefore = eBefore.accuracy(); + assertTrue("after: " + accAfter + ", before=" + accBefore, accAfter >= accBefore + 0.005); + + if (i == 0) { + acc[0] = eBefore.accuracy(); + } + acc[i + 1] = eAfter.accuracy(); + } + log.info("Accuracies: {}", Arrays.toString(acc)); + last = sparkNet.getNetwork().params(); + lastDup = last.dup(); } - log.info("Accuracies: {}", Arrays.toString(acc)); - last = sparkNet.getNetwork().params(); - lastDup = last.dup(); } } @@ -289,7 +323,7 @@ public class GradientSharingTrainingTest extends BaseSparkTest { } - @Test + @Test @Ignore public void testEpochUpdating() throws Exception { //Ensure that epoch counter is incremented properly on the workers @@ -316,7 +350,7 @@ public class GradientSharingTrainingTest extends BaseSparkTest { ComputationGraphConfiguration conf = new NeuralNetConfiguration.Builder() .seed(12345) - .updater(new AMSGrad(0.1)) + .updater(new AMSGrad(0.001)) .graphBuilder() .addInputs("in") .layer("out", new OutputLayer.Builder().nIn(784).nOut(10).activation(Activation.SOFTMAX) diff --git a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/resources/log4j.properties b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/resources/log4j.properties index 5d1edb39f..4bee14770 100644 --- a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/resources/log4j.properties +++ b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/resources/log4j.properties @@ -20,12 +20,12 @@ log4j.appender.Console.layout=org.apache.log4j.PatternLayout log4j.appender.Console.layout.ConversionPattern=%d{ABSOLUTE} %-5p ~ %m%n log4j.appender.org.springframework=DEBUG -log4j.appender.org.deeplearning4j=DEBUG -log4j.appender.org.nd4j=DEBUG +log4j.appender.org.deeplearning4j=INFO +log4j.appender.org.nd4j=INFO log4j.logger.org.springframework=INFO -log4j.logger.org.deeplearning4j=DEBUG -log4j.logger.org.nd4j=DEBUG +log4j.logger.org.deeplearning4j=INFO +log4j.logger.org.nd4j=INFO log4j.logger.org.apache.spark=WARN diff --git a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/resources/logback.xml b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/resources/logback.xml index 4d94f2516..9605642db 100644 --- a/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/resources/logback.xml +++ b/deeplearning4j/deeplearning4j-scaleout/spark/dl4j-spark-parameterserver/src/test/resources/logback.xml @@ -35,7 +35,7 @@ - + diff --git a/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-ui-components/pom.xml b/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-ui-components/pom.xml index 4f2436e28..8f83b803e 100644 --- a/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-ui-components/pom.xml +++ b/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-ui-components/pom.xml @@ -25,10 +25,6 @@ deeplearning4j-ui-components - - 2.3.23 - - org.projectlombok diff --git a/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-ui-components/src/test/java/org/deeplearning4j/ui/TestStandAlone.java b/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-ui-components/src/test/java/org/deeplearning4j/ui/TestStandAlone.java index aaca2eb26..7ba9f9c36 100644 --- a/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-ui-components/src/test/java/org/deeplearning4j/ui/TestStandAlone.java +++ b/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-ui-components/src/test/java/org/deeplearning4j/ui/TestStandAlone.java @@ -24,6 +24,7 @@ import org.deeplearning4j.ui.components.chart.style.StyleChart; import org.deeplearning4j.ui.components.table.ComponentTable; import org.deeplearning4j.ui.components.table.style.StyleTable; import org.deeplearning4j.ui.standalone.StaticPageUtil; +import org.junit.Ignore; import org.junit.Test; import java.awt.*; diff --git a/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-vertx/pom.xml b/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-vertx/pom.xml index 4405d15f7..a66b85ece 100644 --- a/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-vertx/pom.xml +++ b/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-vertx/pom.xml @@ -60,7 +60,7 @@ org.freemarker freemarker - 2.3.29 + ${freemarker.version} diff --git a/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-vertx/src/main/java/org/deeplearning4j/ui/module/train/TrainModule.java b/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-vertx/src/main/java/org/deeplearning4j/ui/module/train/TrainModule.java index 5648de738..00e2c9422 100644 --- a/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-vertx/src/main/java/org/deeplearning4j/ui/module/train/TrainModule.java +++ b/deeplearning4j/deeplearning4j-ui-parent/deeplearning4j-vertx/src/main/java/org/deeplearning4j/ui/module/train/TrainModule.java @@ -200,6 +200,7 @@ public class TrainModule implements UIModule { })); r.add(new Route("/train/:sessionId/info", HttpMethod.GET, (path, rc) -> this.sessionInfoForSession(path.get(0), rc))); } else { + r.add(new Route("/train", HttpMethod.GET, (path, rc) -> rc.reroute("/train/overview"))); r.add(new Route("/train/sessions/current", HttpMethod.GET, (path, rc) -> rc.response().end(currentSessionID == null ? "" : currentSessionID))); r.add(new Route("/train/sessions/set/:to", HttpMethod.GET, (path, rc) -> this.setSession(path.get(0), rc))); r.add(new Route("/train/overview", HttpMethod.GET, (path, rc) -> this.renderFtl("TrainingOverview.html.ftl", rc))); diff --git a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/SDVariable.java b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/SDVariable.java index 6d9e34ed0..65416a659 100644 --- a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/SDVariable.java +++ b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/SDVariable.java @@ -1654,29 +1654,6 @@ public class SDVariable implements Serializable { return x; } - @Override - public boolean equals(Object o) { - if (this == o) { - return true; - } - if (!(o instanceof SDVariable)) { - return false; - } - - SDVariable that = (SDVariable) o; - - if (!Objects.equals(varName, that.varName)) { - return false; - } - if (variableType != that.variableType) { - return false; - } - if(sameDiff != that.sameDiff){ - return false; - } - return dataType == that.dataType; - } - @Override public int hashCode() { int result = super.hashCode(); @@ -1695,4 +1672,26 @@ public class SDVariable implements Serializable { v.sameDiff = sd; return v; } + + @Override + public boolean equals(Object o){ + if(o == this) return true; + if(!(o instanceof SDVariable)) + return false; + + SDVariable s = (SDVariable)o; + if(!varName.equals(s.varName)) + return false; + if(variableType != s.variableType) + return false; + if(dataType != s.dataType) + return false; + + if(variableType == VariableType.VARIABLE || variableType == VariableType.CONSTANT){ + INDArray a1 = getArr(); + INDArray a2 = s.getArr(); + return a1.equals(a2); + } + return true; + } } diff --git a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/SameDiff.java b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/SameDiff.java index 7ca809b2d..3411e2007 100644 --- a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/SameDiff.java +++ b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/SameDiff.java @@ -1234,13 +1234,14 @@ public class SameDiff extends SDBaseOps { @Override public boolean equals(Object o) { if (this == o) return true; - if (o == null || getClass() != o.getClass()) return false; + if (o == null || getClass() != o.getClass()) + return false; SameDiff sameDiff = (SameDiff) o; - if (variables != null ? !variables.equals(sameDiff.variables) : sameDiff.variables != null) - return false; - return sameDiffFunctionInstances != null ? sameDiffFunctionInstances.equals(sameDiff.sameDiffFunctionInstances) : sameDiff.sameDiffFunctionInstances == null; + boolean eqVars = variables.equals(sameDiff.variables); + boolean eqOps = ops.equals(sameDiff.ops); + return eqVars && eqOps; } /** @@ -5843,4 +5844,10 @@ public class SameDiff extends SDBaseOps { return base + "_" + inc; } + + + @Override + public String toString(){ + return "SameDiff(nVars=" + variables.size() + ",nOps=" + ops.size() + ")"; + } } diff --git a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/internal/Variable.java b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/internal/Variable.java index e8041955b..4e7c88a4b 100644 --- a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/internal/Variable.java +++ b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/autodiff/samediff/internal/Variable.java @@ -16,10 +16,7 @@ package org.nd4j.autodiff.samediff.internal; -import lombok.AllArgsConstructor; -import lombok.Builder; -import lombok.Data; -import lombok.NoArgsConstructor; +import lombok.*; import org.nd4j.autodiff.samediff.SDVariable; import java.util.List; @@ -28,6 +25,7 @@ import java.util.List; @NoArgsConstructor @Data //TODO immutable? @Builder +@EqualsAndHashCode(exclude = {"gradient", "variableIndex"}) public class Variable { protected String name; protected SDVariable variable; diff --git a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/classification/EvaluationBinary.java b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/classification/EvaluationBinary.java index 0d2f1fb62..fea9b7308 100644 --- a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/classification/EvaluationBinary.java +++ b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/classification/EvaluationBinary.java @@ -173,9 +173,6 @@ public class EvaluationBinary extends BaseEvaluation { @Override public void eval(INDArray labels, INDArray networkPredictions, INDArray maskArray, List recordMetaData) { - if(recordMetaData != null){ - throw new UnsupportedOperationException("Evaluation with record metadata not yet implemented for EvaluationBinary"); - } eval(labels, networkPredictions, maskArray); } diff --git a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/classification/EvaluationCalibration.java b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/classification/EvaluationCalibration.java index 0d137d0e9..1a0348324 100644 --- a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/classification/EvaluationCalibration.java +++ b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/classification/EvaluationCalibration.java @@ -325,7 +325,7 @@ public class EvaluationCalibration extends BaseEvaluation @Override public void eval(INDArray labels, INDArray networkPredictions, INDArray maskArray, List recordMetaData) { - throw new UnsupportedOperationException("Not yet implemented"); + eval(labels, networkPredictions, maskArray); } @Override diff --git a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/regression/RegressionEvaluation.java b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/regression/RegressionEvaluation.java index cc206f0df..b5fac0dd4 100644 --- a/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/regression/RegressionEvaluation.java +++ b/nd4j/nd4j-backends/nd4j-api-parent/nd4j-api/src/main/java/org/nd4j/evaluation/regression/RegressionEvaluation.java @@ -229,7 +229,7 @@ public class RegressionEvaluation extends BaseEvaluation { @Override public void eval(INDArray labels, INDArray networkPredictions, INDArray maskArray, List recordMetaData) { - throw new UnsupportedOperationException("Not yet implemented"); + eval(labels, networkPredictions, maskArray); } @Override diff --git a/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/autodiff/samediff/SameDiffTests.java b/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/autodiff/samediff/SameDiffTests.java index 409ac422a..73780538a 100644 --- a/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/autodiff/samediff/SameDiffTests.java +++ b/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/autodiff/samediff/SameDiffTests.java @@ -3556,4 +3556,52 @@ public class SameDiffTests extends BaseNd4jTest { assertTrue(msg, msg.contains("\"labels\"") && msg.contains("No array was provided")); } } + + + @Test + public void testEquals1(){ + + SameDiff sd1 = SameDiff.create(); + SameDiff sd2 = SameDiff.create(); + + assertEquals(sd1, sd2); + + SDVariable p1 = sd1.placeHolder("ph", DataType.FLOAT, -1, 10); + SDVariable p2 = sd2.placeHolder("ph", DataType.FLOAT, -1, 10); + + assertEquals(sd1, sd2); + + SDVariable w1 = sd1.constant("c1",1.0f); + SDVariable w2 = sd2.constant("c1",1.0f); + + assertEquals(sd1, sd2); + + SDVariable a1 = p1.add("add", w1); + SDVariable a2 = p2.add("add", w2); + + assertEquals(sd1, sd2); + + SDVariable w1a = sd1.constant("c2", 2.0f); + SDVariable w2a = sd2.constant("cX", 2.0f); + + assertNotEquals(sd1, sd2); + w2a.rename("c2"); + + assertEquals(sd1, sd2); + + sd2.createGradFunction("ph"); + + assertEquals(sd1, sd2); + + w2a.getArr().assign(3.0f); + + assertNotEquals(sd1, sd2); + + w1a.getArr().assign(3.0f); + assertEquals(sd1, sd2); + + SDVariable s1 = p1.sub("op", w1); + SDVariable s2 = p2.add("op", w1); + assertNotEquals(sd1, sd2); + } } diff --git a/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/linalg/nativ/OpsMappingTests.java b/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/linalg/nativ/OpsMappingTests.java index 454739496..03b469e70 100644 --- a/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/linalg/nativ/OpsMappingTests.java +++ b/nd4j/nd4j-backends/nd4j-tests/src/test/java/org/nd4j/linalg/nativ/OpsMappingTests.java @@ -61,7 +61,7 @@ public class OpsMappingTests extends BaseNd4jTest { @Override public long getTimeoutMilliseconds() { - return 90000L; + return 180000L; //Can be slow on some CI machines such as PPC } @Test diff --git a/nd4j/nd4j-common/src/main/java/org/nd4j/resources/Downloader.java b/nd4j/nd4j-common/src/main/java/org/nd4j/resources/Downloader.java index ecaf3ea7f..05c44c29e 100644 --- a/nd4j/nd4j-common/src/main/java/org/nd4j/resources/Downloader.java +++ b/nd4j/nd4j-common/src/main/java/org/nd4j/resources/Downloader.java @@ -95,7 +95,7 @@ public class Downloader { } // try extracting try{ - ArchiveUtils.unzipFileTo(f.getAbsolutePath(), extractToDir.getAbsolutePath()); + ArchiveUtils.unzipFileTo(f.getAbsolutePath(), extractToDir.getAbsolutePath(), false); } catch (Throwable t){ log.warn("Error extracting {} files from file {} - retrying...", name, f.getAbsolutePath(), t); f.delete(); diff --git a/nd4j/nd4j-common/src/main/java/org/nd4j/util/ArchiveUtils.java b/nd4j/nd4j-common/src/main/java/org/nd4j/util/ArchiveUtils.java index f0c6ef318..d51d9ca9b 100644 --- a/nd4j/nd4j-common/src/main/java/org/nd4j/util/ArchiveUtils.java +++ b/nd4j/nd4j-common/src/main/java/org/nd4j/util/ArchiveUtils.java @@ -51,6 +51,10 @@ public class ArchiveUtils { * @throws IOException */ public static void unzipFileTo(String file, String dest) throws IOException { + unzipFileTo(file, dest, true); + } + + public static void unzipFileTo(String file, String dest, boolean logFiles) throws IOException { File target = new File(file); if (!target.exists()) throw new IllegalArgumentException("Archive doesnt exist"); @@ -93,7 +97,9 @@ public class ArchiveUtils { fos.close(); ze = zis.getNextEntry(); - log.debug("File extracted: " + newFile.getAbsoluteFile()); + if(logFiles) { + log.info("File extracted: " + newFile.getAbsoluteFile()); + } } zis.closeEntry(); @@ -112,7 +118,9 @@ public class ArchiveUtils { TarArchiveEntry entry; /* Read the tar entries using the getNextEntry method **/ while ((entry = (TarArchiveEntry) tarIn.getNextEntry()) != null) { - log.info("Extracting: " + entry.getName()); + if(logFiles) { + log.info("Extracting: " + entry.getName()); + } /* If the entry is a directory, create the directory. */ if (entry.isDirectory()) { From 3900d9ff06e0e3c4289474c9aebf04a165be7eaa Mon Sep 17 00:00:00 2001 From: Susan Eraly Date: Mon, 30 Mar 2020 03:07:47 -0700 Subject: [PATCH 6/6] Simple fix in bidirectional lstm import (#293) * first pass Signed-off-by: eraly * cleanup Signed-off-by: eraly --- .../modelimport/keras/layers/wrappers/KerasBidirectional.java | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/wrappers/KerasBidirectional.java b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/wrappers/KerasBidirectional.java index d37ee399c..3b7cb1721 100644 --- a/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/wrappers/KerasBidirectional.java +++ b/deeplearning4j/deeplearning4j-modelimport/src/main/java/org/deeplearning4j/nn/modelimport/keras/layers/wrappers/KerasBidirectional.java @@ -190,7 +190,7 @@ public class KerasBidirectional extends KerasLayer { "Keras Bidirectional layer accepts only one input (received " + inputType.length + ")"); InputPreProcessor preProcessor = getInputPreprocessor(inputType); if (preProcessor != null) - return preProcessor.getOutputType(inputType[0]); + return this.getBidirectionalLayer().getOutputType(-1, preProcessor.getOutputType(inputType[0])); else return this.getBidirectionalLayer().getOutputType(-1, inputType[0]); }