v0.4.0 - Merge branch 'dev' into 'main'

Closes #32 and aidge_core#196

See merge request !61

include/aidge/backend/cuda.hpp

@@ -13,15 +13,18 @@
 #define AIDGE_BACKEND_CUDA_IMPORTS_H_
 
 #include "aidge/backend/cuda/data/TensorImpl.hpp"
+#include "aidge/backend/cuda/operator/OperatorImpl.hpp"
 #include "aidge/backend/cuda/operator/AddImpl.hpp"
 #include "aidge/backend/cuda/operator/AndImpl.hpp"
 #include "aidge/backend/cuda/operator/ArgMaxImpl.hpp"
 #include "aidge/backend/cuda/operator/AvgPoolingImpl.hpp"
 #include "aidge/backend/cuda/operator/BatchNormImpl.hpp"
 #include "aidge/backend/cuda/operator/ConvImpl.hpp"
+#include "aidge/backend/cuda/operator/ClipImpl.hpp"
 #include "aidge/backend/cuda/operator/DivImpl.hpp"
 #include "aidge/backend/cuda/operator/FCImpl.hpp"
 #include "aidge/backend/cuda/operator/GlobalAveragePoolingImpl.hpp"
+#include "aidge/backend/cuda/operator/LRNImpl.hpp"
 #include "aidge/backend/cuda/operator/LnImpl.hpp"
 #include "aidge/backend/cuda/operator/MaxPoolingImpl.hpp"
 #include "aidge/backend/cuda/operator/MulImpl.hpp"
@@ -30,6 +33,7 @@
 #include "aidge/backend/cuda/operator/ReduceMeanImpl.hpp"
 #include "aidge/backend/cuda/operator/ReduceSumImpl.hpp"
 #include "aidge/backend/cuda/operator/ReLUImpl.hpp"
+#include "aidge/backend/cuda/operator/RoundImpl.hpp"
 #include "aidge/backend/cuda/operator/ShiftMaxImpl.hpp"
 #include "aidge/backend/cuda/operator/ShiftGELUImpl.hpp"
 #include "aidge/backend/cuda/operator/ReshapeImpl.hpp"

include/aidge/backend/cuda/data/TensorImpl.hpp

@@ -94,7 +94,7 @@ public:
     }
 
     void copy(const void *src, NbElts_t length, NbElts_t offset = 0) override {
-        AIDGE_ASSERT(length <= mData.size() || length <= mNbElts, "TensorImpl_cuda<{}>::copy(): copy length ({}) is above capacity ({})", typeid(T).name(), length, mNbElts);
+        AIDGE_ASSERT(offset + length <= mNbElts, "TensorImpl_cuda<{}>::copy(): copy offset ({}) + length ({}) is above capacity ({})", typeid(T).name(), offset, length, mNbElts);
         const T* srcT = static_cast<const T *>(src);
         T* dstT = static_cast<T *>(rawPtr(offset));
 
@@ -107,7 +107,7 @@ public:
             return;
         }
 
-        AIDGE_ASSERT(length <= mData.size() || length <= mNbElts, "TensorImpl_cuda<{}>::copyCast(): copy length ({}) is above capacity ({})", typeid(T).name(), length, mNbElts);
+        AIDGE_ASSERT(offset + length <= mNbElts, "TensorImpl_cuda<{}>::copyCast(): copy offset ({}) + length ({}) is above capacity ({})", typeid(T).name(), offset, length, mNbElts);
         switch (srcDt) {
         case DataType::Float64:
             thrust_copy(static_cast<const double*>(src),
@@ -171,17 +171,17 @@ public:
     }
 
     void copyFromDevice(const void *src, const std::pair<std::string, DeviceIdx_t>& device, NbElts_t length, NbElts_t offset = 0) override {
-        AIDGE_ASSERT(length <= mData.size() || length <= mNbElts, "TensorImpl_cuda<{}>::copyFromDevice(): copy length ({}) is above capacity ({})", typeid(T).name(), length, mNbElts);
+        AIDGE_ASSERT(offset + length <= mNbElts, "TensorImpl_cuda<{}>::copyFromDevice(): copy offset ({}) + length ({}) is above capacity ({})", typeid(T).name(), offset, length, mNbElts);
         CHECK_CUDA_STATUS(cudaMemcpy(rawPtr(offset), src, length * sizeof(T), cudaMemcpyDeviceToDevice));
     }
 
     void copyFromHost(const void *src, NbElts_t length, NbElts_t offset = 0) override {
-        AIDGE_ASSERT(length <= mData.size() || length <= mNbElts, "TensorImpl_cuda<{}>::copyFromHost(): copy length ({}) is above capacity ({})", typeid(T).name(), length, mNbElts);
+        AIDGE_ASSERT(offset + length <= mNbElts, "TensorImpl_cuda<{}>::copyFromHost(): copy offset ({}) + length ({}) is above capacity ({})", typeid(T).name(), offset, length, mNbElts);
         CHECK_CUDA_STATUS(cudaMemcpy(rawPtr(offset), src, length * sizeof(T), cudaMemcpyHostToDevice));
     }
 
     void copyToHost(void *dst, NbElts_t length, NbElts_t offset = 0) const override {
-        AIDGE_ASSERT(length <= mData.size() || length <= mNbElts, "TensorImpl_cuda<{}>::copyToHost(): copy length ({}) is above capacity ({})", typeid(T).name(), length, mNbElts);
+        AIDGE_ASSERT(offset + length <= mNbElts, "TensorImpl_cuda<{}>::copyToHost(): copy offset ({}) + length ({}) is above capacity ({})", typeid(T).name(), offset, length, mNbElts);
         CHECK_CUDA_STATUS(cudaMemcpy(dst, rawPtr(offset), length * sizeof(T), cudaMemcpyDeviceToHost));
     }
 

include/aidge/backend/cuda/operator/AddImpl.hpp

@@ -36,7 +36,7 @@ public:
         return std::make_unique<AddImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/AndImpl.hpp

@@ -36,7 +36,7 @@ public:
         return std::make_unique<AndImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/ArgMaxImpl.hpp

@@ -36,7 +36,7 @@ public:
         return std::make_unique<ArgMaxImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/AvgPoolingImpl.hpp

@@ -37,7 +37,7 @@ public:
         return std::make_unique<AvgPoolingImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/BatchNormImpl.hpp

@@ -37,7 +37,7 @@ public:
         return std::make_unique<BatchNormImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/ClipImpl.hpp

+/********************************************************************************
+ * Copyright (c) 2024 CEA-List
+ *
+ * This program and the accompanying materials are made available under the
+ * terms of the Eclipse Public License 2.0 which is available at
+ * http://www.eclipse.org/legal/epl-2.0.
+ *
+ * SPDX-License-Identifier: EPL-2.0
+ *
+ ********************************************************************************/
+
+#ifndef AIDGE_BACKEND_CUDA_OPERATOR_CLIPIMPL_H_
+#define AIDGE_BACKEND_CUDA_OPERATOR_CLIPIMPL_H_
+
+#include <array>
+#include <memory>
+#include <tuple>
+#include <vector>
+
+#include <cudnn.h>
+
+#include "aidge/backend/OperatorImpl.hpp"
+#include "aidge/operator/Clip.hpp"
+#include "aidge/utils/Registrar.hpp"
+#include "aidge/utils/Types.h"
+
+#include "aidge/backend/cuda/utils/CudaUtils.hpp"
+
+namespace Aidge {
+// Operator implementation entry point for the backend
+class ClipImpl_cuda : public OperatorImpl {
+public:
+    ClipImpl_cuda(const Clip_Op& op) : OperatorImpl(op, "cuda") {}
+
+    static std::unique_ptr<ClipImpl_cuda> create(const Clip_Op& op) {
+        return std::make_unique<ClipImpl_cuda>(op);
+    }
+
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
+        return {
+            {DataType::Float64},
+            {DataType::Float32},
+            {DataType::Float16},
+        };
+    }
+
+    void forward() override;
+    void backward() override;
+
+private:
+    template <class T> void forward_();
+    template <class T> void backward_(const Tensor& outGrad);
+};
+
+// Implementation entry point registration to Operator
+REGISTRAR(Clip_Op, "cuda", Aidge::ClipImpl_cuda::create);
+}  // namespace Aidge
+
+#endif /* AIDGE_BACKEND_CUDA_OPERATOR_CLIPIMPL_H_ */

include/aidge/backend/cuda/operator/ClipImpl_CUDA_kernels.hpp

+/********************************************************************************
+ * Copyright (c) 2024 CEA-List
+ *
+ * This program and the accompanying materials are made available under the
+ * terms of the Eclipse Public License 2.0 which is available at
+ * http://www.eclipse.org/legal/epl-2.0.
+ *
+ * SPDX-License-Identifier: EPL-2.0
+ *
+ ********************************************************************************/
+
+#ifndef AIDGE_CUDA_OPERATOR_CLIPIMPL_KERNELS_H_
+#define AIDGE_CUDA_OPERATOR_CLIPIMPL_KERNELS_H_
+
+#include <stdexcept>
+#include <cfloat>
+#include <cuda.h>
+#include <cuda_runtime_api.h>
+#include <cuda_fp16.h>
+
+#include "aidge/data/Data.hpp"
+#include "aidge/backend/cuda/utils/CudaUtils.hpp"
+#include "aidge/utils/Types.h"
+
+namespace Aidge {
+
+template <class T>
+void clipForward(const T* input, T* output,int size,T min_val, T max_val);
+
+
+}
+#endif /* AIDGE_CUDA_OPERATOR_CLIPIMPL_KERNELS_H_ */
+
+
+
+
+

include/aidge/backend/cuda/operator/ConvImpl.hpp

@@ -43,7 +43,7 @@ public:
         return std::make_unique<ConvImpl_cuda<DIM>>(op, true);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Any}
         };

include/aidge/backend/cuda/operator/DivImpl.hpp

@@ -36,7 +36,7 @@ public:
         return std::make_unique<DivImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/FCImpl.hpp

@@ -36,7 +36,7 @@ public:
         return std::make_unique<FCImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/GlobalAveragePoolingImpl.hpp

@@ -36,7 +36,7 @@ public:
         return std::make_unique<GlobalAveragePoolingImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Any}
         };

include/aidge/backend/cuda/operator/ILayerNormImpl.hpp

@@ -37,7 +37,7 @@ public:
         return std::make_unique<ILayerNormImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/LRNImpl.hpp

+/********************************************************************************
+ * Copyright (c) 2023 CEA-List
+ *
+ * This program and the accompanying materials are made available under the
+ * terms of the Eclipse Public License 2.0 which is available at
+ * http://www.eclipse.org/legal/epl-2.0.
+ *
+ * SPDX-License-Identifier: EPL-2.0
+ *
+ ********************************************************************************/
+
+#ifndef AIDGE_BACKEND_CUDA_OPERATOR_LRNIMPL_H_
+#define AIDGE_BACKEND_CUDA_OPERATOR_LRNIMPL_H_
+
+#include <array>
+#include <memory>
+#include <tuple>
+#include <vector>
+
+#include <cudnn.h>
+
+#include "aidge/backend/OperatorImpl.hpp"
+#include "aidge/operator/LRN.hpp"
+#include "aidge/utils/Registrar.hpp"
+#include "aidge/utils/Types.h"
+
+#include "aidge/backend/cuda/utils/CudaUtils.hpp"
+
+namespace Aidge {
+// Operator implementation entry point for the backend
+class LRNImpl_cuda : public OperatorImpl {
+public:
+    LRNImpl_cuda(const LRN_Op& op) : OperatorImpl(op, "cuda") {}
+
+    static std::unique_ptr<LRNImpl_cuda> create(const LRN_Op& op) {
+        return std::make_unique<LRNImpl_cuda>(op);
+    }
+
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
+        return {
+            {DataType::Any}
+        };
+    }
+
+    void forward() override;
+    void backward() override;
+    ~LRNImpl_cuda();
+
+private:
+    // CuDNN specific variables
+    cudnnLRNDescriptor_t mLRNDesc = nullptr;
+    std::shared_ptr<Tensor> mInputFallback;
+    std::shared_ptr<Tensor> mOutputGradFallback;
+
+    template <class T> void forward_(const Tensor& input);
+    template <class T> void backward_(const Tensor& output_grad);
+};
+
+// Implementation entry point registration to Operator
+REGISTRAR(LRN_Op, "cuda", Aidge::LRNImpl_cuda::create);
+}  // namespace Aidge
+
+#endif /* AIDGE_BACKEND_CUDA_OPERATOR_LRNIMPL_H_ */

include/aidge/backend/cuda/operator/LnImpl.hpp

@@ -36,7 +36,7 @@ public:
         return std::make_unique<LnImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/MaxPoolingImpl.hpp

@@ -37,7 +37,7 @@ public:
         return std::make_unique<MaxPoolingImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Any}
         };

include/aidge/backend/cuda/operator/MulImpl.hpp

@@ -36,7 +36,7 @@ public:
         return std::make_unique<MulImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},

include/aidge/backend/cuda/operator/OperatorImpl.hpp

+/********************************************************************************
+ * Copyright (c) 2023 CEA-List
+ *
+ * This program and the accompanying materials are made available under the
+ * terms of the Eclipse Public License 2.0 which is available at
+ * http://www.eclipse.org/legal/epl-2.0.
+ *
+ * SPDX-License-Identifier: EPL-2.0
+ *
+ ********************************************************************************/
+
+#ifndef AIDGE_CUDA_OPERATOR_IMPL_H_
+#define AIDGE_CUDA_OPERATOR_IMPL_H_
+
+#include <cstddef>  // std::size_t
+#include <memory>
+#include <tuple>    // std::tuple
+#include <vector>
+
+#include "aidge/backend/OperatorImpl.hpp"
+#include "aidge/utils/Registrar.hpp"
+#include "aidge/utils/Types.h"
+
+namespace Aidge {
+
+template <class Op, class FwdFunc, class BwdFunc = void()>
+class OperatorImpl_cuda : public OperatorImpl,
+    public Registrable<OperatorImpl_cuda<Op, FwdFunc, BwdFunc>, ImplSpec, Impl<FwdFunc, BwdFunc>>
+{
+public:
+    OperatorImpl_cuda(const Op& op) : OperatorImpl(op, "cuda") {}
+
+    static std::unique_ptr<OperatorImpl_cuda<Op, FwdFunc, BwdFunc>> create(const Op& op) {
+        return std::make_unique<OperatorImpl_cuda<Op, FwdFunc, BwdFunc>>(op);
+    }
+
+    virtual std::shared_ptr<ProdConso> getProdConso() const override {
+        const auto impl = Registrar<OperatorImpl_cuda>::create(getBestMatch(getRequiredSpec()));
+        return impl.prodConso(mOp);
+    }
+
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
+        // return Registrar<OperatorImpl_cuda>::getKeys(); // Note: cannot return set due to python binding 
+        std::set<ImplSpec> implSpecsSet = Registrar<OperatorImpl_cuda>::getKeys();
+        return std::vector<ImplSpec>(implSpecsSet.begin(), implSpecsSet.end());
+    }
+
+    void forward() override;
+    void backward() override;
+};
+}  // namespace Aidge
+
+#endif /* AIDGE_CUDA_OPERATOR_IMPL_H_ */

include/aidge/backend/cuda/operator/PadImpl.hpp

@@ -37,7 +37,7 @@ public:
         return std::make_unique<PadImpl_cuda>(op);
     }
 
-    virtual std::set<ImplSpec> getAvailableImplSpecs() const override {
+    virtual std::vector<ImplSpec> getAvailableImplSpecs() const override {
         return {
             {DataType::Float64},
             {DataType::Float32},