Merge branch 'develop' of https://github.com/PaddlePaddle/models into develop

.gitignore

@@ -4,3 +4,4 @@
 *.pyc
 *~
 *.vscode
+*.idea
\ No newline at end of file

PaddleCV/Paddle3D/PointNet++/.gitignore

+checkpoints*
+ext_op/src/*.o
+ext_op/src/*.so
+*.log*
+dataset/Indoor3DSemSeg/*
+!dataset/Indoor3DSemSeg/*.sh
+dataset/ModelNet40/*
+!dataset/ModelNet40/*.sh

PaddleCV/Paddle3D/PointNet++/README.md

+# PointNet++ 分类和语义分割模型
+
+---
+## 内容
+
+- [简介](#简介)
+- [快速开始](#快速开始)
+- [参考文献](#参考文献)
+- [版本更新](#版本更新)
+
+## 简介
+
+[PointNet++](https://arxiv.org/abs/1706.02413) 是 Charles R. Qi, Li Yi, Hao Su, Leonidas J. Guibas 等人提出的，针对3D数据进行分类和语义分割的模型。该模型基于PointNet进行了拓展, 使用分层点集特征学习来提取点云数据的特征，首先通过对输入point进行分组和采样提取局部区域模式，然后使用多层感知器来获取点特征。PointNet++ 还将点特征传播用于语义分割模型，采用基于距离插值和跨级跳转连接的分层传播策略，对点特征进行向上采样，获得所有原始点的点特征。
+
+
+网络结构如下所示：
+
+<p align="center">
+<img src="image/pointnet2.jpg" height=300 width=800 hspace='10'/> <br />
+用于点云分类和分割的 PointNet++ 网络结构
+</p>
+
+集合抽象层是网络的基本模块，每个集合抽象层由三个关键层构成:采样层、分组层和特征提取层。
+
+- **采样层**：采样层使用最远点采样(FPS)的方法，从输入点中选择一组点，它定义了局部区域的中心。与随机抽样的方法相比，在质心数目相同的情况下，FPS可以更好的覆盖整个点集。
+
+- **分组层**：分组层通过寻找中心体周围的“邻近”点来构造局部区域集。在度量空间采样的点集中，点的邻域由度量距离定义。这种方法被称为“query ball”，它使得局部区域的特征在空间上更加一般化。
+
+- **特征提取层**: 特征提取层使用 mini-PointNet 对分组层给出的各个区域进行特征提取，获得局部特征。
+
+
+
+**注意:** PointNet++ 模型构建依赖于自定义的 C++ 算子，目前仅支持GPU设备在Linux/Unix系统上进行编译，本模型**不能运行在Windows系统或CPU设备上**
+
+
+## 快速开始
+
+### 安装
+
+**安装 [PaddlePaddle](https://github.com/PaddlePaddle/Paddle):**
+
+在当前目录下运行样例代码需要 PaddelPaddle Fluid [develop每日版本](https://www.paddlepaddle.org.cn/install/doc/tables#多版本whl包列表-dev-11)或使用PaddlePaddle [develop分支](https://github.com/PaddlePaddle/Paddle/tree/develop)源码编译安装. 
+
+为了使自定义算子与paddle版本兼容，建议您**优先使用源码编译paddle**，源码编译方式请参考[编译安装](https://www.paddlepaddle.org.cn/install/doc/source/ubuntu)
+
+
+### 编译自定义OP
+
+请确认Paddle版本为PaddelPaddle Fluid develop每日版本或基于Paddle develop分支源码编译安装，**推荐使用源码编译安装的方式**。
+
+自定义OP编译方式如下：
+
+    进入 `ext_op/src` 目录，执行编译脚本
+    ```
+    cd ext_op/src
+    sh make.sh
+    ```
+
+    成功编译后，`ext_op/src` 目录下将会生成 `pointnet2_lib.so` 
+
+    执行下列操作，确保自定义算子编译正确：
+
+    ```
+    # 设置动态库的路径到 LD_LIBRARY_PATH 中
+    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+
+    # 回到 ext_op 目录，添加 PYTHONPATH
+    cd ..
+    export PYTHONPATH=$PYTHONPATH:`pwd`
+
+    # 运行单测 
+    python tests/test_farthest_point_sampling_op.py
+    python tests/test_gather_point_op.py
+    python tests/test_group_points_op.py
+    python tests/test_query_ball_op.py
+    python tests/test_three_interp_op.py
+    python tests/test_three_nn_op.py
+    ```
+    单测运行成功会输出提示信息，如下所示：
+
+    ```
+    .
+    ----------------------------------------------------------------------
+    Ran 1 test in 13.205s
+
+    OK
+    ```
+
+**说明：** 更多关于自定义OP的编译说明，请参考[自定义OP编译](./ext_op/README.md)
+
+
+### 数据准备
+
+**ModelNet40 数据集:**
+
+PointNet++ 分类模型在 [ModelNet40 数据集](https://shapenet.cs.stanford.edu/media/modelnet40_ply_hdf5_2048.zip)上进行训练，我们提供了数据集下载脚本：
+
+```
+cd dataset/ModelNet40
+sh download.sh
+```
+
+数据目录结构如下所示：
+
+```
+  dataset/ModelNet40/modelnet40_ply_hdf5_2048
+  ├── train_files.txt
+  ├── test_files.txt
+  ├── shape_names.txt
+  ├── ply_data_train0.h5
+  ├── ply_data_train_0_id2file.json
+  ├── ply_data_test0.h5
+  ├── ply_data_test_0_id2file.json
+  |   ...
+
+```
+
+**Indoor3DSemSeg 数据集:**
+
+PointNet++ 分割模型在 [Indoor3DSemSeg 数据集](https://shapenet.cs.stanford.edu/media/indoor3d_sem_seg_hdf5_data.zip)上进行训练，我们提供了数据集下载脚本：
+
+```
+cd dataset/Indoor3DSemSeg
+sh download.sh
+```
+
+数据目录结构如下所示：
+
+```
+  dataset/Indoor3DSemSeg/
+  ├── all_files.txt
+  ├── room_filelist.txt
+  ├── ply_data_all_0.h5
+  ├── ply_data_all_1.h5
+  |   ...
+
+```
+
+### 训练
+
+分类/分割模型默认使用单卡训练，在启动训练前请指定单卡GPU，并将动态库的路径添加到 LD_LIBRARY_PATH 中：
+
+```
+# 指定0号卡进行GPU训练
+export CUDA_VISIBLE_DEVICES=0
+
+# 设置动态库的路径到 LD_LIBRARY_PATH 中
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+
+```
+
+**分类模型:**
+
+可通过如下方式启动 PointNet++ 分类模型的训练：
+
+```
+# 开始训练
+python train_cls.py --model=MSG --batch_size=16 --save_dir=checkpoints_msg_cls
+```
+
+我们同时提供了训练分类模型的“快速开始”脚本：
+
+```
+sh scripts/train_cls.sh
+```
+
+**语义分割模型:**
+
+可通过如下方式启动 PointNet++ 语义分割模型的训练：
+
+```
+# 开始训练
+python train_seg.py --model=MSG --batch_size=32 --save_dir=checkpoints_msg_seg
+```
+
+我们同时提供了训练语义分割模型的“快速开始”脚本：
+
+```
+sh scripts/train_seg.sh
+```
+
+### 模型评估
+
+
+分类/分割模型默认使用单卡评估，首先指定单卡GPU，并将动态库的路径添加到 LD_LIBRARY_PATH 中：
+
+```
+# 指定0号卡进行GPU评估
+export CUDA_VISIBLE_DEVICES=0
+
+# 设置动态库的路径到 LD_LIBRARY_PATH 中
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+
+```
+
+**分类模型:**
+
+可通过如下方式启动 PointNet++ 分类模型的评估：
+
+```
+# 对给定权重进行评估
+python eval_cls.py --model=MSG --weights=checkpoints_cls/200
+```
+
+我们同时提供了评估分类模型的“快速开始”脚本：
+
+```
+sh scripts/eval_cls.sh
+```
+
+分类模型的评估结果如下所示：
+
+| model | Top-1 | download |
+| :----- | :---: | :---: |
+| SSG(Single-Scale Group) | 89.3 | [model](https://paddlemodels.bj.bcebos.com/Paddle3D/pointnet2_ssg_cls.tar) |
+| MSG(Multi-Scale Group)  | 90.0 | [model](https://paddlemodels.bj.bcebos.com/Paddle3D/pointnet2_msg_cls.tar) |
+
+**语义分割模型:**
+
+可通过如下方式启动 PointNet++ 语义分割模型的评估：
+
+```
+# 对给定权重进行评估
+python eval_seg.py --model=MSG --weights=checkpoints_seg/200
+```
+
+我们同时提供了评估语义分割模型的“快速开始”脚本：
+
+```
+sh scripts/eval_seg.sh
+```
+
+语义分割模型的评估结果如下所示：
+
+| model | Top-1 | download |
+| :----- | :---: | :---: |
+| SSG(Single-Scale Group) | 86.1 | [model](https://paddlemodels.bj.bcebos.com/Paddle3D/pointnet2_ssg_seg.tar) |
+| MSG(Multi-Scale Group)  | 86.6 | [model](https://paddlemodels.bj.bcebos.com/Paddle3D/pointnet2_msg_seg.tar) |
+
+## 参考文献
+
+- [PointNet++: Deep Hierarchical Feature Learning on Point Sets in a Metric Space](https://arxiv.org/abs/1706.02413), Charles R. Qi, Li Yi, Hao Su, Leonidas J. Guibas.
+- [PointNet: Deep Learning on Point Sets for 3D Classification and Segmentation](https://www.semanticscholar.org/paper/PointNet%3A-Deep-Learning-on-Point-Sets-for-3D-and-Qi-Su/d997beefc0922d97202789d2ac307c55c2c52fba), Charles Ruizhongtai Qi, Hao Su, Kaichun Mo, Leonidas J. Guibas.
+
+## 版本更新
+
+- 11/2019, 新增 PointNet++ 分类和语义分割模型。

PaddleCV/Paddle3D/PointNet++/README_en.md

+# PointNet++ classification and semantic segmentation model
+
+---
+## Table of Contents
+
+- [Introduction](#introduction)
+- [Quick Start](#quick-start)
+- [Reference](#reference)
+- [Update](#update)
+
+## Introduction
+
+[PointNet++](https://arxiv.org/abs/1706.02413) is a point classification and segmentation model for 3D data proposed by Charles R. Qi, Li Yi, Hao Su, Leonidas J. Guibas.
+This model is a extension work based on PointNet extract features of point clouds data with hierarchical point set feature learning, perform set abstractions by grouping and sampling points at first to extract
+local region patterns, then use multi-layer perceptron to get point features. PointNet++ also used point feature propagation for semantic segmentation model, adopt a hierarchical
+propagation strategy with distance based interpolation and across level skip links, point features was upsampled to obtain point features for all the original points.
+
+The network structure is shown as below.
+
+<p align="center">
+<img src="image/pointnet2.jpg" height=300 width=800 hspace='10'/> <br />
+PointNet++ architecture for Point set Segmentation and Classification
+</p>
+
+Set Abstraction layer is the basic module of the network, each set abstraction layer is made of three key layers：Sampling layer, Grouping layer and PointNet layer.
+
+- **Sample layer**: Sampling layer uses farthest point sampling(FPS) to select a set of points from input points, which defines the centroids of local regions. Compared with random sampling, it has better converage of the entire point set given the same number of centroids.
+
+- **Grouping layer**: Grouping layer constructs local region sets by finding "neighboring" points around the centroids. In a point set sampled from a metric space, the neighborhood of a point is defined by metric distance. This method is called "ball query", which make local region feature more generalizable across space.
+
+- **PointNet layer**: PointNet layer uses a mini-PointNet to encode local region patterns into feature vectors.
+
+
+**NOTE:** PointNet++ model builds base on custom C++ operations, which can only support GPU devices and compiled on Linux/Unix currently, this model **cannot run on Windows or CPU deivices**.
+
+
+## Quick Start
+
+### Installation
+
+**Install [PaddlePaddle](https://github.com/PaddlePaddle/Paddle):**
+
+Running sample code in this directory requires PaddelPaddle Fluid develop [daily version wheel](https://www.paddlepaddle.org.cn/install/doc/tables#多版本whl包列表-dev-11) or compiled from PaddlePaddle [develop branch](https://github.com/PaddlePaddle/Paddle/tree/develop).
+
+In order to make the custom OP compatible with the Paddle version, it is recommended to **compile from PaddlePaddle develop branch source code**. For source code compilation, please refer to [Compile and Install](https://www.paddlepaddle.org.cn/install/doc/source/ubuntu)
+
+### Compile custom operations
+
+Please make sure you are using PaddlePaddle Fluid develop daily version or compiled from PaddlePaddle develop branch.
+Custom operations can be compiled as follows:
+
+```
+cd ext_op/src
+sh make.sh
+```
+
+If the compilation is finished successfully, `pointnet2_lib.so` will be generated under `exr_op/src`.
+
+Make sure custom operations pass as follows:
+
+```
+# export paddle libs to LD_LIBRARY_PATH for custom op library
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+
+# back to ext_op and add PYTHONPATH
+cd ..
+export PYTHONPATH=$PYTHONPATH:`pwd`
+
+# Run unit tests
+python test/test_farthest_point_sampling_op.py
+python test/test_gather_point_op.py
+python test/test_group_points_op.py
+python test/test_query_ball_op.py
+python test/test_three_interp_op.py
+python test/test_three_nn_op.py
+```
+The prompt message for successful running is as follows:
+
+```
+.
+----------------------------------------------------------------------
+Ran 1 test in 13.205s
+
+OK
+```
+
+### Data preparation
+
+**ModelNet40 dataset:**
+
+PointNet++ classification models are trained on [ModelNet40 dataset](https://shapenet.cs.stanford.edu/media/modelnet40_ply_hdf5_2048.zip), we also provide download scripts as follows:
+
+```
+cd dataset/ModelNet40
+sh download.sh
+```
+
+The dataset catalog structure is as follows:
+
+```
+  dataset/ModelNet40/modelnet40_ply_hdf5_2048
+  ├── train_files.txt
+  ├── test_files.txt
+  ├── shape_names.txt
+  ├── ply_data_train0.h5
+  ├── ply_data_train_0_id2file.json
+  ├── ply_data_test0.h5
+  ├── ply_data_test_0_id2file.json
+  |   ...
+
+```
+
+**Indoor3DSemSeg dataset:**
+
+PointNet++ semantic segmentation models are trained on [Indoor3DSemSeg dataset](https://shapenet.cs.stanford.edu/media/indoor3d_sem_seg_hdf5_data.zip), we also provide download scripts as follows:
+
+```
+cd dataset/Indoor3DSemSeg
+sh download.sh
+```
+
+The dataset catalog structure is as follows:
+
+```
+  dataset/Indoor3DSemSeg/
+  ├── all_files.txt
+  ├── room_filelist.txt
+  ├── ply_data_all_0.h5
+  ├── ply_data_all_1.h5
+  |   ...
+
+```
+
+### Training
+
+**Classification Model:**
+
+For PointNet++ classification model, training can be start as follows:
+
+```
+# For single GPU deivces
+export CUDA_VISIBLE_DEVICES=0
+
+# enable gc to save GPU memory
+export FLAGS_fast_eager_deletion_mode=1
+export FLAGS_eager_delete_tensor_gb=0.0
+export FLAGS_fraction_of_gpu_memory_to_use=0.98
+
+# export paddle libs to LD_LIBRARY_PATH for custom op library
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+
+# start training
+python train_cls.py --model=MSG --batch_size=16 --save_dir=checkpoints_msg_cls
+```
+
+We also provided quick start script for training classification model as follows:
+
+```
+sh scripts/train_cls.sh
+```
+
+**Semantic Segmentation Model:**
+
+For PointNet++ semantic segmentation model, training can be start as follows:
+
+```
+# For single GPU deivces
+export CUDA_VISIBLE_DEVICES=0
+
+# enable gc to save GPU memory
+export FLAGS_fast_eager_deletion_mode=1
+export FLAGS_eager_delete_tensor_gb=0.0
+export FLAGS_fraction_of_gpu_memory_to_use=0.98
+
+# export paddle libs to LD_LIBRARY_PATH for custom op library
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+
+# start training
+python train_seg.py --model=MSG --batch_size=32 --save_dir=checkpoints_msg_seg
+```
+
+We also provided quick start scripts for training semantic segmentation model as follows:
+
+```
+sh scripts/train_seg.sh
+```
+
+### Evaluation
+
+**Classification Model:**
+
+For PointNet++ classification model, evaluation can be start as follows:
+
+```
+# For single GPU deivces
+export CUDA_VISIBLE_DEVICES=0
+
+# export paddle libs to LD_LIBRARY_PATH for custom op library
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+
+# start evaluation with given weights
+python eval_cls.py --model=MSG --weights=checkpoints_cls/200
+```
+
+We also provided quick start script for training classification model as follows:
+
+```
+sh scripts/eval_cls.sh
+```
+
+Classification model evaluation result is shown as below:
+
+| model | Top-1 | download |
+| :----- | :---: | :---: |
+| SSG(Single-Scale Group) | 89.3 | [model]() |
+| MSG(Multi-Scale Group)  | 90.0 | [model]() |
+
+**Semantic Segmentation Model:**
+
+For PointNet++ semantic segmentation model, evaluation can be start as follows:
+
+```
+# For single GPU deivces
+export CUDA_VISIBLE_DEVICES=0
+
+# export paddle libs to LD_LIBRARY_PATH for custom op library
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+
+# start evaluation with given weights 
+python eval_seg.py --model=MSG --weights=checkpoints_seg/200
+```
+
+We also provided quick start scripts for training semantic segmentation model as follows:
+
+```
+sh scripts/eval_seg.sh
+```
+
+Semantic segmentation model evaluation result is shown as below:
+
+| model | Top-1 | download |
+| :----- | :---: | :---: |
+| SSG(Single-Scale Group) | 86.1 | [model]() |
+| MSG(Multi-Scale Group)  | 86.8 | [model]() |
+
+## Reference
+
+- [PointNet++: Deep Hierarchical Feature Learning on Point Sets in a Metric Space](https://arxiv.org/abs/1706.02413), Charles R. Qi, Li Yi, Hao Su, Leonidas J. Guibas.
+- [PointNet: Deep Learning on Point Sets for 3D Classification and Segmentation](https://www.semanticscholar.org/paper/PointNet%3A-Deep-Learning-on-Point-Sets-for-3D-and-Qi-Su/d997beefc0922d97202789d2ac307c55c2c52fba), Charles Ruizhongtai Qi, Hao Su, Kaichun Mo, Leonidas J. Guibas.
+
+## Update
+
+- 11/2019, Add PointNet++ classification and semantic segmentation model.

PaddleCV/Paddle3D/PointNet++/data/__init__.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+
+from . import indoor3d_reader
+from . import modelnet40_reader
+from .indoor3d_reader import *
+from .modelnet40_reader import *
+
+__all__ = indoor3d_reader.__all__
+__all__ += modelnet40_reader.__all__

PaddleCV/Paddle3D/PointNet++/data/data_utils.py

+"""
+This code is based on https://github.com/erikwijmans/Pointnet2_PyTorch/blob/master/pointnet2/data/data_utils.py
+"""
+
+from __future__ import absolute_import
+from __future__ import division
+from __future__ import print_function
+from __future__ import unicode_literals
+
+import numpy as np
+
+def angle_axis(angle, axis):
+    """
+    Returns a 4x4 rotation matrix that performs a rotation around axis by angle
+    Parameters
+    ----------
+    angle : float
+        Angle to rotate by
+    axis: np.ndarray
+        Axis to rotate about
+    Returns
+    -------
+    Tensor
+        3x3 rotation matrix
+    """
+    u = axis / np.linalg.norm(axis)
+    cosval, sinval = np.cos(angle), np.sin(angle)
+
+    # yapf: disable
+    cross_prod_mat = np.array([[0.0, -u[2], u[1]],
+                                [u[2], 0.0, -u[0]],
+                                [-u[1], u[0], 0.0]])
+
+    R = np.array(
+        cosval * np.eye(3)
+        + sinval * cross_prod_mat
+        + (1.0 - cosval) * np.outer(u, u)).astype("float32")
+    return R
+
+class PointcloudScale(object):
+    def __init__(self, lo=0.8, hi=1.25):
+        self.lo, self.hi = lo, hi
+
+    def __call__(self, points):
+        scaler = np.random.uniform(self.lo, self.hi)
+        points[:, 0:3] *= scaler
+        return points
+
+class PointcloudRotate(object):
+    def __init__(self, axis=np.array([0.0, 1.0, 0.0])):
+        self.axis = axis
+
+    def __call__(self, points):
+        rotation_angle = np.random.uniform() * 2 * np.pi
+        rotation_matrix = angle_axis(rotation_angle, self.axis)
+
+        normals = points.shape[1] > 3
+        if not normals:
+            return np.matmul(points, rotation_matrix.T)
+        else:
+            pc_xyz = points[:, 0:3]
+            pc_normals = points[:, 3:]
+            points[:, 0:3] = np.matmul(pc_xyz, rotation_matrix.T)
+            points[:, 3:] = np.matmul(pc_normals, rotation_matrix.T)
+            return points
+
+class PointcloudTranslate(object):
+    def __init__(self, translate_range=0.1):
+        self.translate_range = translate_range
+
+    def __call__(self, points):
+        translation = np.random.uniform(-self.translate_range, self.translate_range)
+        points[:, 0:3] += translation
+        return points
+
+
+class PointcloudJitter(object):
+    def __init__(self, std=0.01, clip=0.05):
+        self.std, self.clip = std, clip
+
+    def __call__(self, points):
+        jittered_data = np.random.normal(loc=0,scale=self.std,size=(points.shape[0],3))
+        jittered_data = np.clip(jittered_data, -self.clip, self.clip)
+
+        points[:, 0:3] += jittered_data
+        return points
+
+class PointcloudRotatePerturbation(object):
+    def __init__(self, angle_sigma=0.06, angle_clip=0.18):
+        self.angle_sigma, self.angle_clip = angle_sigma, angle_clip
+
+    def _get_angles(self):
+        angles = np.clip(
+            self.angle_sigma * np.random.randn(3), -self.angle_clip, self.angle_clip
+        )
+        return angles
+    def __call__(self, points):
+        angles = self._get_angles()
+        Rx = angle_axis(angles[0], np.array([1.0, 0.0, 0.0]))
+        Ry = angle_axis(angles[1], np.array([0.0, 1.0, 0.0]))
+        Rz = angle_axis(angles[2], np.array([0.0, 0.0, 1.0]))
+
+        rotation_matrix = np.matmul(np.matmul(Rz, Ry), Rx)
+
+        normals = points.shape[1] > 3
+        if not normals:
+            return np.matmul(points, rotation_matrix.T)
+        else:
+            pc_xyz = points[:, 0:3]
+            pc_normals = points[:, 3:]
+            points[:, 0:3] = np.matmul(pc_xyz, rotation_matrix.T)
+            points[:, 3:] = np.matmul(pc_normals, rotation_matrix.T)
+            return points
+
+
+class PointcloudRandomInputDropout(object):
+    def __init__(self, max_dropout_ratio=0.875):
+        assert max_dropout_ratio >= 0 and max_dropout_ratio < 1
+        self.max_dropout_ratio = max_dropout_ratio
+
+    def __call__(self, points):
+        dropout_ratio = np.random.random() * self.max_dropout_ratio  # 0~0.875
+        drop_idx = np.where(np.random.random((points.shape[0])) <= dropout_ratio)[0]
+        if len(drop_idx) > 0:
+            points[drop_idx] = points[0]  # set to the first point
+
+        return points

PaddleCV/Paddle3D/PointNet++/data/indoor3d_reader.py

+# Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://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.
+
+from __future__ import absolute_import
+from __future__ import division
+from __future__ import print_function
+from __future__ import unicode_literals
+
+import os
+import os.path as osp
+import signal
+import numpy as np
+import h5py
+import random
+import logging
+
+__all__ = ["Indoor3DReader"]
+
+logger = logging.getLogger(__name__)
+
+
+class Indoor3DReader(object):
+    def __init__(self, data_dir, test_area="Area_5"):
+        self.data_dir = data_dir
+        self.test_area = test_area
+        self.load_data()
+
+    def _read_data_file(self, fname):
+        assert osp.isfile(fname), \
+                "{} is not a file".format(fname)
+        with open(fname) as f:
+            return [line.strip() for line in f]
+
+    def _load_h5_file(self, fname):
+        assert osp.isfile(fname), \
+                "{} is not a file".format(fname)
+        f = h5py.File(fname, mode='r')
+        return f['data'][:], f['label'][:]
+
+    def load_data(self):
+        logger.info("Loading Indoor3D dataset from {} ...".format(self.data_dir))
+        # read all_files.txt
+        all_files_fname = osp.join(self.data_dir, 'all_files.txt')
+        all_files = self._read_data_file(all_files_fname)
+
+        # read room_filelist.txt
+        room_fname = osp.join(self.data_dir, 'room_filelist.txt')
+        room_filelist = self._read_data_file(room_fname)
+
+        points, labels = [], []
+        for f in all_files:
+            h5_fname = osp.join(self.data_dir, osp.split(f)[-1]) 
+            point, label = self._load_h5_file(h5_fname)
+            points.append(point)
+            labels.append(label)
+        points = np.concatenate(points, 0)
+        labels = np.concatenate(labels, 0)
+
+        train_idxs, test_idxs = [], []
+        for i, room in enumerate(room_filelist):
+            if self.test_area in room:
+                test_idxs.append(i)
+            else:
+                train_idxs.append(i)
+
+        self.data = {}
+        self.data['train'] = {}
+        self.data['train']['points'] = points[train_idxs, ...]
+        self.data['train']['labels'] = labels[train_idxs, ...]
+        self.data['test'] = {}
+        self.data['test']['points'] = points[test_idxs, ...]
+        self.data['test']['labels'] = labels[test_idxs, ...]
+        logger.info("Load data finished")
+
+    def get_reader(self, batch_size, num_points, mode='train', shuffle=True):
+        assert mode in ['train', 'test'], \
+                "mode can only be 'train' or 'test'"
+        data = self.data[mode]
+        points = data['points']
+        labels = data['labels']
+
+        if mode == 'train' and shuffle:
+            idxs = np.arange(len(points))
+            np.random.shuffle(idxs)
+            points = points[idxs]
+            labels = labels[idxs]
+
+        def reader():
+            batch_out = []
+            for point, label in zip(points, labels):
+                # shuffle points
+                p = point.copy()
+                l = label.copy()
+                pt_idxs = np.arange(num_points)
+                np.random.shuffle(pt_idxs)
+                p = p[pt_idxs]
+                l = l[pt_idxs]
+
+                xyz = p[:, :3]
+                feature = p[:, 3:]
+                label = l[:, np.newaxis]
+                batch_out.append((xyz, feature, label))
+
+                if len(batch_out) == batch_size:
+                    yield batch_out
+                    batch_out = []
+
+        return reader
+
+
+def _term_reader(signum, frame):
+    logger.info('pid {} terminated, terminate reader process '
+                'group {}...'.format(os.getpid(), os.getpgrp()))
+    os.killpg(os.getpgid(os.getpid()), signal.SIGKILL)
+
+signal.signal(signal.SIGINT, _term_reader)
+signal.signal(signal.SIGTERM, _term_reader)
+

PaddleCV/Paddle3D/PointNet++/data/modelnet40_reader.py

+# Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://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.
+
+from __future__ import absolute_import
+from __future__ import division
+from __future__ import print_function
+from __future__ import unicode_literals
+
+import os
+import os.path as osp
+import signal
+import numpy as np
+import h5py
+import random
+import logging
+
+__all__ = ["ModelNet40ClsReader"]
+
+logger = logging.getLogger(__name__)
+
+
+class ModelNet40ClsReader(object):
+    def __init__(self, data_dir, mode='train', transforms=None):
+        assert mode in ['train', 'test'], \
+                "mode can only be 'train' or 'test'"
+        self.data_dir = data_dir
+        self.mode = mode 
+        self.transforms = transforms
+        self.load_data()
+
+    def _read_data_file(self, fname):
+        assert osp.isfile(fname), \
+                "{} is not a file".format(fname)
+        with open(fname) as f:
+            return [line.strip()[5:] for line in f]
+
+    def _load_h5_file(self, fname):
+        assert osp.isfile(fname), \
+                "{} is not a file".format(fname)
+        f = h5py.File(fname, mode='r')
+        return f['data'][:], f['label'][:]
+
+    def load_data(self):
+        logger.info("Loading ModelNet40 dataset {} split from {} "
+                    "...".format(self.mode, self.data_dir))
+        if self.mode == 'train':
+            files_fname = osp.join(self.data_dir, 'train_files.txt')
+            files = self._read_data_file(files_fname)
+        else:
+            files_fname = osp.join(self.data_dir, 'test_files.txt')
+            files = self._read_data_file(files_fname)
+
+        points, labels = [], []
+        for f in files:
+            h5_fname = osp.join(self.data_dir, osp.split(f)[-1])
+            point, label = self._load_h5_file(h5_fname)
+            points.append(point)
+            labels.append(label)
+        self.points = np.concatenate(points, 0)
+        self.labels = np.concatenate(labels, 0)
+        logger.info("Load {} data finished".format(self.mode))
+
+    def get_reader(self, batch_size, num_points, shuffle=True):
+        self.num_points = min(num_points, self.points.shape[1])
+        points = self.points
+        labels = self.labels
+        if shuffle and self.mode == 'train':
+            idxs = np.arange(len(self.points))
+            np.random.shuffle(idxs)
+            points = points[idxs]
+            labels = labels[idxs]
+
+        def reader():
+            batch_out = []
+            for point, label in zip(points, labels):
+                p = point.copy()
+                l = label.copy()
+                pt_idxs = np.arange(self.num_points)
+                if shuffle:
+                    np.random.shuffle(pt_idxs)
+                c_points = p[pt_idxs]
+                if self.transforms is not None:
+                    for trans in self.transforms:
+                        c_points = trans(c_points)
+                
+                xyz = c_points[:, :3]
+                # modelnet40 only have xyz features
+                # feature = c_points[:, 3:]
+                label = l[:, np.newaxis]
+                batch_out.append((xyz, label))
+
+                if len(batch_out) == batch_size:
+                    yield batch_out
+                    batch_out = []
+        return reader
+
+
+def _term_reader(signum, frame):
+    logger.info('pid {} terminated, terminate reader process '
+                'group {}...'.format(os.getpid(), os.getpgrp()))
+    os.killpg(os.getpgid(os.getpid()), signal.SIGKILL)
+
+signal.signal(signal.SIGINT, _term_reader)
+signal.signal(signal.SIGTERM, _term_reader)
+

PaddleCV/Paddle3D/PointNet++/dataset/Indoor3DSemSeg/download.sh

+DIR="$( cd "$(dirname "$0")" ; pwd -P  )"
+cd "$DIR"
+
+echo "Downloading https://shapenet.cs.stanford.edu/media/indoor3d_sem_seg_hdf5_data.zip"
+wget https://shapenet.cs.stanford.edu/media/indoor3d_sem_seg_hdf5_data.zip
+
+echo "Unzip indoor3d_sem_seg_hdf5_data.zip"
+unzip indoor3d_sem_seg_hdf5_data.zip

PaddleCV/Paddle3D/PointNet++/dataset/ModelNet40/download.sh

+DIR="$( cd "$(dirname "$0")" ; pwd -P  )"
+cd "$DIR"
+
+echo "Downloading https://shapenet.cs.stanford.edu/media/modelnet40_ply_hdf5_2048.zip"
+wget https://shapenet.cs.stanford.edu/media/modelnet40_ply_hdf5_2048.zip
+
+echo "Unzip modelnet40_ply_hdf5_2048.zip"
+unzip modelnet40_ply_hdf5_2048.zip

PaddleCV/Paddle3D/PointNet++/eval_cls.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+
+import os
+import sys
+import time
+import shutil
+import argparse
+import ast
+import logging
+import numpy as np
+import paddle.fluid as fluid
+
+from models import *
+from data.data_utils import *
+from data.modelnet40_reader import ModelNet40ClsReader 
+from utils import *
+
+logging.root.handlers = []
+FORMAT = '%(asctime)s-%(levelname)s: %(message)s'
+logging.basicConfig(level=logging.INFO, format=FORMAT, stream=sys.stdout)
+logger = logging.getLogger(__name__)
+
+np.random.seed(1024)
+
+
+def parse_args():
+    parser = argparse.ArgumentParser("PointNet++ semantic segmentation train script")
+    parser.add_argument(
+        '--model',
+        type=str,
+        default='MSG',
+        help='SSG or MSG model to train, default MSG')
+    parser.add_argument(
+        '--use_gpu',
+        type=ast.literal_eval,
+        default=True,
+        help='default use gpu.')
+    parser.add_argument(
+        '--batch_size',
+        type=int,
+        default=1,
+        help='evaluation batch size, default 1')
+    parser.add_argument(
+        '--num_points',
+        type=int,
+        default=4096,
+        help='number of points in a sample, default: 4096')
+    parser.add_argument(
+        '--num_classes',
+        type=int,
+        default=40,
+        help='number of classes in dataset, default: 13')
+    parser.add_argument(
+        '--weights',
+        type=str,
+        default='checkpoints/200',
+        help='directory name to save train snapshoot')
+    parser.add_argument(
+        '--data_dir',
+        type=str,
+        default='dataset/ModelNet40/modelnet40_ply_hdf5_2048',
+        help='dataset directory')
+    parser.add_argument(
+        '--log_interval',
+        type=int,
+        default=100,
+        help='mini-batch interval for logging.')
+    args = parser.parse_args()
+    return args
+
+
+def eval():
+    args = parse_args()
+    print_arguments(args)
+    # check whether the installed paddle is compiled with GPU
+    check_gpu(args.use_gpu)
+
+    assert args.model in ['MSG', 'SSG'], \
+            "--model can only be 'MSG' or 'SSG'"
+
+    # build model
+    startup = fluid.Program()
+    eval_prog = fluid.Program()
+    with fluid.program_guard(eval_prog, startup):
+        with fluid.unique_name.guard():
+            eval_model = PointNet2ClsMSG(args.num_classes, args.num_points) \
+                           if args.model == 'MSG' else \
+                         PointNet2ClsSSG(args.num_classes, args.num_points)
+            eval_model.build_model()
+            eval_feeds = eval_model.get_feeds()
+            eval_outputs = eval_model.get_outputs()
+            eval_pyreader = eval_model.get_pyreader()
+    eval_prog = eval_prog.clone(True)
+    eval_keys, eval_values = parse_outputs(eval_outputs)
+
+    place = fluid.CUDAPlace(0) if args.use_gpu else fluid.CPUPlace()
+    exe = fluid.Executor(place)
+    exe.run(startup)
+
+    assert os.path.exists(args.weights), "weights {} not exists.".format(args.weights)
+    def if_exist(var):
+        return os.path.exists(os.path.join(args.weights, var.name))
+    fluid.io.load_vars(exe, args.weights, eval_prog, predicate=if_exist)
+
+    eval_compile_prog = fluid.compiler.CompiledProgram(eval_prog)
+    
+    # get reader
+    modelnet_reader = ModelNet40ClsReader(args.data_dir, mode='test')
+    eval_reader = modelnet_reader.get_reader(args.batch_size, args.num_points)
+    eval_pyreader.decorate_sample_list_generator(eval_reader, place)
+
+    eval_stat = Stat()
+    try:
+        eval_pyreader.start()
+        eval_iter = 0
+        eval_periods = []
+        while True:
+            cur_time = time.time()
+            eval_outs = exe.run(eval_compile_prog, fetch_list=eval_values)
+            period = time.time() - cur_time
+            eval_periods.append(period)
+            eval_stat.update(eval_keys, eval_outs)
+            if eval_iter % args.log_interval == 0:
+                log_str = ""
+                for name, value in zip(eval_keys, eval_outs):
+                    log_str += "{}: {:.4f}, ".format(name, np.mean(value))
+                logger.info("[EVAL] batch {}: {}time: {:.2f}".format(eval_iter, log_str, period))
+            eval_iter += 1
+    except fluid.core.EOFException:
+        logger.info("[EVAL] Eval finished, {}average time: {:.2f}".format(eval_stat.get_mean_log(), np.mean(eval_periods[1:])))
+    finally:
+        eval_pyreader.reset()
+
+
+if __name__ == "__main__":
+    eval()

PaddleCV/Paddle3D/PointNet++/eval_seg.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+
+import os
+import sys
+import time
+import shutil
+import argparse
+import ast
+import logging
+import numpy as np
+import paddle.fluid as fluid
+
+from models import *
+from data.indoor3d_reader import Indoor3DReader
+from utils import *
+
+logging.root.handlers = []
+FORMAT = '%(asctime)s-%(levelname)s: %(message)s'
+logging.basicConfig(level=logging.INFO, format=FORMAT, stream=sys.stdout)
+logger = logging.getLogger(__name__)
+
+np.random.seed(1024)
+
+
+def parse_args():
+    parser = argparse.ArgumentParser("PointNet++ semantic segmentation train script")
+    parser.add_argument(
+        '--model',
+        type=str,
+        default='MSG',
+        help='SSG or MSG model to train, default MSG')
+    parser.add_argument(
+        '--use_gpu',
+        type=ast.literal_eval,
+        default=True,
+        help='default use gpu.')
+    parser.add_argument(
+        '--batch_size',
+        type=int,
+        default=1,
+        help='evaluation batch size, default 1')
+    parser.add_argument(
+        '--num_points',
+        type=int,
+        default=4096,
+        help='number of points in a sample, default: 4096')
+    parser.add_argument(
+        '--num_classes',
+        type=int,
+        default=13,
+        help='number of classes in dataset, default: 13')
+    parser.add_argument(
+        '--weights',
+        type=str,
+        default='checkpoints/200',
+        help='directory name to save train snapshoot')
+    parser.add_argument(
+        '--data_dir',
+        type=str,
+        default='dataset/Indoor3DSemSeg/indoor3d_sem_seg_hdf5_data',
+        help='dataset directory')
+    parser.add_argument(
+        '--log_interval',
+        type=int,
+        default=100,
+        help='mini-batch interval for logging.')
+    args = parser.parse_args()
+    return args
+
+
+def eval():
+    args = parse_args()
+    print_arguments(args)
+    # check whether the installed paddle is compiled with GPU
+    check_gpu(args.use_gpu)
+
+    assert args.model in ['MSG', 'SSG'], \
+            "--model can only be 'MSG' or 'SSG'"
+
+    # build model
+    startup = fluid.Program()
+    eval_prog = fluid.Program()
+    with fluid.program_guard(eval_prog, startup):
+        with fluid.unique_name.guard():
+            eval_model = PointNet2SemSegMSG(args.num_classes, args.num_points) \
+                           if args.model == 'MSG' else \
+                         PointNet2SemSegSSG(args.num_classes, args.num_points)
+            eval_model.build_model()
+            eval_feeds = eval_model.get_feeds()
+            eval_outputs = eval_model.get_outputs()
+            eval_pyreader = eval_model.get_pyreader()
+    eval_prog = eval_prog.clone(True)
+    eval_keys, eval_values = parse_outputs(eval_outputs)
+
+    place = fluid.CUDAPlace(0) if args.use_gpu else fluid.CPUPlace()
+    exe = fluid.Executor(place)
+    exe.run(startup)
+
+    assert os.path.exists(args.weights), "weights {} not exists.".format(args.weights)
+    def if_exist(var):
+        return os.path.exists(os.path.join(args.weights, var.name))
+    fluid.io.load_vars(exe, args.weights, eval_prog, predicate=if_exist)
+
+    eval_compile_prog = fluid.compiler.CompiledProgram(eval_prog)
+    
+    # get reader
+    indoor_reader = Indoor3DReader(args.data_dir)
+    eval_reader = indoor_reader.get_reader(args.batch_size, args.num_points, mode='test')
+    eval_pyreader.decorate_sample_list_generator(eval_reader, place)
+
+    eval_stat = Stat()
+    try:
+        eval_pyreader.start()
+        eval_iter = 0
+        eval_periods = []
+        while True:
+            cur_time = time.time()
+            eval_outs = exe.run(eval_compile_prog, fetch_list=eval_values)
+            period = time.time() - cur_time
+            eval_periods.append(period)
+            eval_stat.update(eval_keys, eval_outs)
+            if eval_iter % args.log_interval == 0:
+                log_str = ""
+                for name, value in zip(eval_keys, eval_outs):
+                    log_str += "{}: {:.4f}, ".format(name, np.mean(value))
+                logger.info("[EVAL] batch {}: {}time: {:.2f}".format(eval_iter, log_str, period))
+            eval_iter += 1
+    except fluid.core.EOFException:
+        logger.info("[EVAL] Eval finished, {}average time: {:.2f}".format(eval_stat.get_mean_log(), np.mean(eval_periods[1:])))
+    finally:
+        eval_pyreader.reset()
+
+
+if __name__ == "__main__":
+    eval()

PaddleCV/Paddle3D/PointNet++/ext_op/README.md

+# 自定义OP的编译过程
+
+## 代码结构
+
+  - src: 扩展OP C++/CUDA 源码
+  - pointnet_lib.py: Python封装
+  - tests: 各OP单测程序
+
+## 安装PaddlePaddle
+
+请通过如下方式安装PaddlePaddle：
+
+- 通过[Paddle develop分支](https://github.com/PaddlePaddle/Paddle/tree/develop)源码编译安装，编译方法如下:
+
+  1. [Ubuntu](https://www.paddlepaddle.org.cn/install/doc/source/ubuntu)
+  1. [CentOS](https://www.paddlepaddle.org.cn/install/doc/source/centos)
+  1. [MasOS](https://www.paddlepaddle.org.cn/install/doc/source/macos)
+  1. [Windows](https://www.paddlepaddle.org.cn/install/doc/source/windows)
+
+  **说明：** 推荐使用docker编译
+
+- 安装Paddle develop[每日版本whl包](https://www.paddlepaddle.org.cn/install/doc/tables#多版本whl包列表-dev-11)
+
+  **注意：** 编译自定义OP使用的gcc版本须与Paddle编译使用gcc版本一致，Paddle develop每日版本目前采用**gcc 4.8.2**版本编译，若使用每日版本，请使用**gcc 4.8.2**版本编译自定义OP，否则可能出现兼容性问题。
+
+## 编译自定义OP
+
+自定义op需要将实现的C++、CUDA代码编译成动态库，mask.sh中通过g++/nvcc编译，当然您也可以写Makefile或者CMake。
+
+编译需要include PaddlePaddle的相关头文件，链接PaddlePaddle的lib库。 头文件和lib库可通过下面命令获取到:
+
+```
+# python
+>>> import paddle
+>>> print(paddle.sysconfig.get_include())
+/paddle/pyenv/local/lib/python2.7/site-packages/paddle/include
+>>> print(paddle.sysconfig.get_lib())
+/paddle/pyenv/local/lib/python2.7/site-packages/paddle/libs
+```
+
+我们提供动态库编译脚本如下：
+
+```
+cd src
+sh make.sh
+```
+
+最终编译会产出`pointnet_lib.so`
+
+**说明：** 若使用源码编译安装PaddlePaddle的方式，编译过程中`cmake`未设置`WITH_MKLDNN`的方式，
+编译自定义OP时会报错找不到`mkldnn.h`等文件，可在`make.sh`中删除编译命令中的`-DPADDLE_WITH_MKLDNN`选项。
+
+## 设置环境变量
+
+需要将Paddle的核心库设置到`LD_LIBRARY_PATH`里, 先运行下面程序获取路径:
+
+```
+import paddle
+print(paddle.sysconfig.get_lib())
+```
+
+可通过如下方式添加动态库路径:
+
+```
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+```
+
+## 执行单测
+
+执行下列单测，确保自定义算子可在网络中正确使用：
+
+```
+# 回到 ext_op 目录，添加 PYTHONPATH
+cd ..
+export PYTHONPATH=$PYTHONPATH:`pwd`
+
+# 运行单测 
+python test/test_farthest_point_sampling_op.py
+python test/test_gather_point_op.py
+python test/test_group_points_op.py
+python test/test_query_ball_op.py
+python test/test_three_interp_op.py
+python test/test_three_nn_op.py
+```
+
+单测运行成功会输出提示信息，如下所示：
+
+```
+.
+----------------------------------------------------------------------
+Ran 1 test in 13.205s
+
+OK
+```
+
+更多关于如何在框架外部自定义 C++ OP，可阅读[官网说明文档](https://www.paddlepaddle.org.cn/documentation/docs/zh/advanced_usage/index_cn.html)

PaddleCV/Paddle3D/PointNet++/ext_op/__init__.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+
+from . import pointnet_lib
+from .pointnet_lib import *
+
+__all__ = pointnet_lib.__all__

PaddleCV/Paddle3D/PointNet++/ext_op/pointnet_lib.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+
+import os
+import paddle.fluid as fluid
+
+file_dir = os.path.dirname(os.path.abspath(__file__))
+fluid.load_op_library(os.path.join(file_dir, 'src/pointnet_lib.so'))
+
+from paddle.fluid.layer_helper import LayerHelper
+
+__all__ = ['three_nn', 'three_interp', 'query_ball', 'gather_point',
+            'farthest_point_sampling', 'group_points']
+
+
+def three_nn(input, known, eps=1e-10, name=None):
+    """
+    **Three Nearest Neighbor Layer**
+
+    This operator samples the top-3 nearest neighbor of each point
+    coordinates specified by Input(X) between known point coordinates
+    specified by Input(Known) and calcualte the distance between these
+    nearest neighbors.
+
+    Args:
+        input (Variable): The input tensor of three_nn operator. This
+                          is a 3-D tensor with shape of [B, N, 3].
+        known (Variable): The input tensor of known points of three_nn
+                          operator. This is a 3-D tensor with shape of
+                          [B, M, 3].
+        name(str|None): A name for this layer(optional). If set None, the layer
+                        will be named automatically.
+
+    Returns:
+        distance (Variable): The output distance tensor of three_nn operator.
+                             This is a 3-D tensor with shape of [B, N, 3].
+        idx (Variable): The output index tensor of three_nn operator.
+                             This is a 3-D tensor with shape of [B, N, 3].
+
+    Examples:
+
+        .. code-block:: python
+
+            import paddle.fluid as fluid
+            x = fluid.layers.data(name='x', shape=[16, 3], dtype='float32')
+            known = fluid.layers.data(name='known', shape=[32, 3], dtype='float32')
+            distance, idx = fluid.layers.three_nn(input, known)
+    """
+    helper = LayerHelper('three_nn', **locals())
+    dtype = helper.input_dtype()
+    dist = helper.create_variable_for_type_inference(dtype)
+    idx = helper.create_variable_for_type_inference(dtype)
+    helper.append_op(
+        type="three_nn",
+        inputs={"X": input,
+                "Known": known},
+        outputs={"Distance": dist,
+                 "Idx": idx},
+        attrs={'eps': eps})
+    return (dist, idx)
+
+
+def three_interp(input, weight, idx, name=None):
+    """
+    **Three Interpolate Layer**
+
+    This operator calculate interpolate results from input, weight and
+    index.
+
+    Args:
+        input (Variable): The input tensor of three_interp operator. This
+                          is a 3-D tensor with shape of [B, M, C].
+        weight (Variable): The weight tensor of three_interp operator. This
+                          is a 3-D tensor with shape of [B, N, 3].
+        idx (Variable): The index tensor of three_interp operator. This
+                          is a 3-D tensor with shape of [B, N, 3].
+        name(str|None): A name for this layer(optional). If set None, the layer
+                        will be named automatically.
+
+    Returns:
+        output (Variable): The output tensor of three_interp operator.
+                             This is a 3-D tensor with shape of [B, N, C].
+
+    Examples:
+
+        .. code-block:: python
+
+            import paddle.fluid as fluid
+            x = fluid.layers.data(name='x', shape=[16, 3], dtype='float32')
+            weight = fluid.layers.data(name='weight', shape=[32, 3], dtype='float32')
+            index = fluid.layers.data(name='index', shape=[32, 3], dtype='int32')
+            out = fluid.layers.three_interp(x, weight, index)
+    """
+    helper = LayerHelper('three_interp', **locals())
+    dtype = helper.input_dtype()
+    out = helper.create_variable_for_type_inference(dtype)
+    helper.append_op(
+        type="three_interp",
+        inputs={"X": input,
+                "Weight": weight,
+                "Idx": idx},
+        outputs={"Out": out, })
+    return out
+
+
+def query_ball(input, new_points, radius, n_sample):
+    """
+    **Query Ball Layer**
+
+    Output is a tensor with the indicies of the features that form the query balls.
+
+    Args:
+        input(Variable): XYZ coordinates of features with shape of [B,N,3].
+        new_points(Variable): Centers coordinates of the ball query with shape of [B,M,3].
+        radius(float|Variable): Radius of the balls.
+        n_sample(int|Variable): Maximum number of features in the balls.
+    Return:
+        output(Variable): Tensor with the indicies of the features that form the query balls,with shape of [B,M,n_sample]
+
+    Examples:
+        .. code-block::python
+
+            import paddle.fluid as fluid
+            x = fluid.layers.data(name='points',shape=[-1,5,3],dtype='float32')
+            new_points = fluid.layers.data(name='new_points', shape=[-1,2,3], dtype='float32')
+            output = fluid.layers.query_ball(x,new_points,radius=4.0,n_sample=5)
+
+
+
+    """
+    helper = LayerHelper('query_ball', **locals())
+    dtype = helper.input_dtype()
+    out = helper.create_variable_for_type_inference(dtype)
+    helper.append_op(
+        type="query_ball",
+        inputs={"Points": input,
+                "New_Points": new_points},
+        attrs={"N_sample": n_sample,
+               "Radius": radius},
+        outputs={"Output": out})
+    return out
+
+
+def farthest_point_sampling(input, sampled_point_num):
+    '''
+    Sampling point based on its max eucliden distance with other points. 
+    
+    Args:
+        input (Variable): input point cloud dataset with shape (B, N, 3)
+            B is batch size, N is points's nums, 3 is (x,y,z) coordinate
+        sampled_point_num (int): sampled points's nums
+
+    Retrun:
+        output (Variable): return sampled points with shape (B, M)
+            B is batch size, M is points's nums
+
+    Examples:
+        .. code-block:: python
+        x = fluid.layers.data(name='data', shape=(2,100,3), dtype='float32')
+        sampled_points = fluid.layers.farthest_point_sampling(
+            x, 50
+        )
+    '''
+
+    helper = LayerHelper('farthest_point_sampling', **locals())
+    dtype = input.dtype
+    op_out = helper.create_variable_for_type_inference(dtype)
+    helper.append_op(
+        type='farthest_point_sampling',
+        inputs={'X': input},
+        outputs={'Output': op_out},
+        attrs={'sampled_point_num': sampled_point_num})
+    return op_out
+
+
+def gather_point(input, index):
+    """
+    **Gather Point Layer**
+    Output is obtained by gathering entries of X indexed by `index` 
+    and concatenate them together.
+    .. math::
+        Out = X[Index]
+    .. code-block:: text
+        Given:
+        X = [[1, 2, 3],
+             [3, 4, 5],
+             [5, 6, 7]]
+        Index = [[1, 2]
+        Then:
+        Out = [[3, 4, 5],
+               [5, 6, 7]]
+    Args:
+        input (Variable): The source input with rank>=1, This
+                          is a 3-D tensor with shape of [B, N, 3].
+        index (Variable): The index input with shape of [B, M].
+      
+    Returns:
+        output (Variable): The output is a tensor with shape of [B,M].
+    Examples:
+        .. code-block:: python
+            import paddle.fluid as fluid
+            x = fluid.layers.data(name='x', shape=[-1, 5, 3], dtype='float32')
+            index = fluid.layers.data(name='index', shape=[-1, 1], dtype='int32')
+            output = fluid.layers.gather_point(x, index)
+    """
+
+    helper = LayerHelper('gather_point', **locals())
+    dtype = helper.input_dtype()
+    out = helper.create_variable_for_type_inference(dtype)
+    helper.append_op(
+        type="gather_point",
+        inputs={"X": input,
+                "Index": index},
+        outputs={"Output": out})
+    return out
+
+
+def group_points(input, idx, name=None):
+    """
+    **Group Points Layer**
+
+    This operator group input points with index.
+
+    Args:
+        input (Variable): The input tensor of three_interp operator. This
+                          is a 3-D tensor with shape of [B, N, C].
+        idx (Variable): The index tensor of three_interp operator. This
+                          is a 3-D tensor with shape of [B, M, S].
+        name(str|None): A name for this layer(optional). If set None, the layer
+                        will be named automatically.
+
+    Returns:
+        output (Variable): The output tensor of three_interp operator.
+                             This is a 4-D tensor with shape of [B, M, S, C].
+
+    Examples:
+
+        .. code-block:: python
+
+            import paddle.fluid as fluid
+            x = fluid.layers.data(name='x', shape=[16, 3], dtype='float32')
+            index = fluid.layers.data(name='index', shape=[32, 3], dtype='int32')
+            out  = fluid.layers.group_points(x, index)
+    """
+    helper = LayerHelper('group_points', **locals())
+    dtype = helper.input_dtype()
+    out = helper.create_variable_for_type_inference(dtype)
+    helper.append_op(
+        type="group_points",
+        inputs={"X": input,
+                "Idx": idx},
+        outputs={"Out": out, })
+    return out

PaddleCV/Paddle3D/PointNet++/ext_op/src/farthest_point_sampling_op.cc

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+    http://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. */
+
+#include <memory>
+#include <string>
+#include <vector>
+#include "paddle/fluid/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+
+class FarthestPointSamplingOpMaker : public framework::OpProtoAndCheckerMaker {
+public:
+  void Make() override {
+    AddInput("X",
+             "(Tensor)input point cloud dataset with shape (B, N, 3)"
+             "B is batch size, N is points's nums, 3 is (x,y,z) coordinate");
+    AddOutput("Output",
+              "(Tensor)return sampled points with shape (B, M)"
+              "B is batch size, M is points's nums");
+    AddAttr<int>("sampled_point_num", "sampling points's num")
+        .SetDefault(0)
+        .EqualGreaterThan(0);
+    AddComment(
+        R"Doc(
+            Sampling point based on 
+            its max eucliden distance with other points.)Doc");
+  }
+};
+
+class FarthestPointSamplingOp : public framework::OperatorWithKernel {
+public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+protected:
+  void InferShape(framework::InferShapeContext *ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) shoud not be null");
+    auto x_dims = ctx->GetInputDim("X");
+    PADDLE_ENFORCE(x_dims.size() == 3,
+                   "Input(X) of FathestPointSamplingOp should be 3-D Tensor");
+    const int m = ctx->Attrs().Get<int>("sampled_point_num");
+    ctx->SetOutputDim("Output", {x_dims[0], m});
+  }
+
+protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext &ctx) const override {
+    auto input_data_type = ctx.Input<Tensor>("X")->type();
+    return framework::OpKernelType(input_data_type, ctx.GetPlace());
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OPERATOR(farthest_point_sampling,
+                  ops::FarthestPointSamplingOp,
+                  ops::FarthestPointSamplingOpMaker);

PaddleCV/Paddle3D/PointNet++/ext_op/src/farthest_point_sampling_op.cu

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+    http://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. */
+
+#include "paddle/fluid/framework/eigen.h"
+#include "paddle/fluid/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+
+template <typename T, unsigned int block_size>
+__global__ void farthestpointsamplingKernel(int b,
+                                            int n,
+                                            int m,
+                                            const T *__restrict__ dataset,
+                                            T *__restrict__ temp,
+                                            int *__restrict__ idxs) {
+  // 1. add first point
+  // 2. add the point having farthest distance with first point's
+  // 3. make second point as first point, repeat 1,2
+  if (m <= 0) return;
+  const int BlockSize = block_size;
+  __shared__ float dists[BlockSize];
+  __shared__ int dists_i[BlockSize];
+  const int BufferSize = 3072;
+  __shared__ float buf[BufferSize * 3];
+
+  // one block one batch, n points
+  // one thread one point
+  for (int i = blockIdx.x; i < b; i += gridDim.x) {
+    // can select old point as first point randomly
+    int old = 0;
+    if (threadIdx.x == 0) idxs[i * m + 0] = old;
+
+    for (int j = threadIdx.x; j < n; j += blockDim.x) {
+      temp[blockIdx.x * n + j] = 1e38;
+    }
+    for (int j = threadIdx.x; j < min(BufferSize, n) * 3; j += blockDim.x) {
+      buf[j] = dataset[i * n * 3 + j];
+    }
+    // wait all threads do this in the same block
+    __syncthreads();
+
+    // out m points
+    for (int j = 1; j < m; j++) {
+      // Step 1.
+      // fatherest distance
+      int besti = 0;
+      float best = -1;
+      // first point in m points
+      float x1 = dataset[i * n * 3 + old * 3 + 0];
+      float y1 = dataset[i * n * 3 + old * 3 + 1];
+      float z1 = dataset[i * n * 3 + old * 3 + 2];
+
+      // Step 2.
+      // find farthest point of (x1, y1, z1)
+      for (int k = threadIdx.x; k < n; k += blockDim.x) {
+        float td = temp[blockIdx.x * n + k];
+        float x2, y2, z2;
+        if (k < BufferSize) {
+          x2 = buf[k * 3 + 0];
+          y2 = buf[k * 3 + 1];
+          z2 = buf[k * 3 + 2];
+        } else {
+          x2 = dataset[i * n * 3 + k * 3 + 0];
+          y2 = dataset[i * n * 3 + k * 3 + 1];
+          z2 = dataset[i * n * 3 + k * 3 + 2];
+        }
+        // compute eucliden distance
+        float d = (x2 - x1) * (x2 - x1) + (y2 - y1) * (y2 - y1) +
+                  (z2 - z1) * (z2 - z1);
+        float d2 = min(d, td);
+        if (d2 != td) temp[blockIdx.x * n + k] = d2;
+        if (d2 > best) {
+          best = d2;
+          besti = k;
+        }
+      }
+
+      // step 3.
+      dists[threadIdx.x] = best;
+      dists_i[threadIdx.x] = besti;
+      for (int u = 0; (1 << u) < blockDim.x; u++) {
+        __syncthreads();
+        if (threadIdx.x < (blockDim.x >> (u + 1))) {
+          int i1 = (threadIdx.x * 2) << u;
+          int i2 = (threadIdx.x * 2 + 1) << u;
+          if (dists[i1] < dists[i2]) {
+            dists[i1] = dists[i2];
+            dists_i[i1] = dists_i[i2];
+          }
+        }
+      }
+      __syncthreads();
+      // store the found node index
+      old = dists_i[0];
+      if (threadIdx.x == 0) idxs[i * m + j] = old;
+    }
+  }
+}
+
+template <typename T>
+class FarthestPointSamplingOpCUDAKernel : public framework::OpKernel<T> {
+public:
+  void Compute(const framework::ExecutionContext &ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "This kernel only runs on GPU device.");
+    auto *input = ctx.Input<Tensor>("X");
+    auto *output = ctx.Output<Tensor>("Output");
+    if (input->numel() == 0) return;
+    // allocate memory
+    auto *ptr_out_points_index = output->mutable_data<int>(ctx.GetPlace());
+
+    // b, n, m
+    int batch_size = input->dims()[0];
+    int in_n_points = input->dims()[1];
+    int out_m_points = ctx.Attr<int>("sampled_point_num");
+
+    const T *ptr_in_points = input->data<T>();
+
+    Tensor tmp;
+    auto *ptr_tmp_e =
+        tmp.mutable_data<T>({batch_size, in_n_points}, ctx.GetPlace());
+
+    // run fathest point sampling kernel
+    // P40 have max 512 thread
+    farthestpointsamplingKernel<T, 512><<<32, 512>>>(batch_size,
+                                                     in_n_points,
+                                                     out_m_points,
+                                                     ptr_in_points,
+                                                     ptr_tmp_e,
+                                                     ptr_out_points_index);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_CUDA_KERNEL(farthest_point_sampling,
+                        ops::FarthestPointSamplingOpCUDAKernel<float>,
+                        ops::FarthestPointSamplingOpCUDAKernel<double>);

PaddleCV/Paddle3D/PointNet++/ext_op/src/gather_point_op.cc

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+    http://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. */
+
+#include "paddle/fluid/framework/op_registry.h"
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+
+class GatherPointOp : public framework::OperatorWithKernel {
+public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) shoud not be null");
+    auto x_dims = ctx->GetInputDim("X");
+    PADDLE_ENFORCE(x_dims.size() == 3 && x_dims[2] == 3,
+                   "Input(X) of GatherPointOp should be 3-D Tensor, the last "
+                   "dimension must be 3");
+    auto index_dims = ctx->GetInputDim("Index");
+    PADDLE_ENFORCE(index_dims.size() == 2 && index_dims[0] == x_dims[0],
+                   "Index of GatherPointop should be 2-D Tensor");
+    ctx->SetOutputDim("Output", {x_dims[0], index_dims[1], 3});
+  }
+
+protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    auto input_data_type = ctx.Input<Tensor>("X")->type();
+    return framework::OpKernelType(input_data_type, ctx.GetPlace());
+  }
+};
+
+class GatherPointOpMaker : public framework::OpProtoAndCheckerMaker {
+public:
+  void Make() override {
+    AddInput("X",
+             "Input points with shape (batch, n, 3), n is input "
+             "points's num");
+    AddInput("Index",
+             "input index with shape (batch, m), m is output points's num");
+    AddOutput("Output", "output points with shape(batch, m, 3)");
+    AddComment(
+        R"Doc(
+        Gather Point Operator.
+        Out is obtained by gathering entries of X indexed by Index and 
+        concatenate them together.
+
+        Example:
+        X = [[1, 2, 3],
+             [3, 4, 5],
+             [5, 6, 7]]
+        Index = [[1, 2]]
+
+        Then:
+        Out = [[3, 4, 5],[5, 6, 7]])Doc");
+  }
+};
+
+class GatherPointOpGrad : public framework::OperatorWithKernel {
+public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("Index"), "Input(Index) should not be null");
+    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Output")),
+                   "Input(Output@GRAD) should not be null");
+    auto dim_x = ctx->GetInputDim("X");
+    if (ctx->HasOutput(framework::GradVarName("X"))) {
+      ctx->SetOutputDim(framework::GradVarName("X"), dim_x);
+    }
+  }
+
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::OpKernelType(
+        ctx.Input<Tensor>(framework::GradVarName("Output"))->type(),
+        ctx.GetPlace());
+  }
+};
+
+template <typename T>
+class GatherPointGradDescMaker : public framework::SingleGradOpMaker<T> {
+public:
+  using framework::SingleGradOpMaker<T>::SingleGradOpMaker;
+
+protected:
+  std::unique_ptr<T> Apply() const override {
+    auto* op = new T();
+    op->SetType("gather_point_grad");
+    op->SetInput("X", this->Input("X"));
+    op->SetInput("Index", this->Input("Index"));
+    op->SetInput(framework::GradVarName("Output"), this->OutputGrad("Output"));
+    op->SetOutput(framework::GradVarName("X"), this->InputGrad("X"));
+    op->SetAttrMap(this->Attrs());
+    return std::unique_ptr<T>(op);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OPERATOR(gather_point,
+                  ops::GatherPointOp,
+                  ops::GatherPointOpMaker,
+                  ops::GatherPointGradDescMaker<paddle::framework::OpDesc>,
+                  ops::GatherPointGradDescMaker<paddle::imperative::OpBase>);
+REGISTER_OPERATOR(gather_point_grad, ops::GatherPointOpGrad);

PaddleCV/Paddle3D/PointNet++/ext_op/src/gather_point_op.cu

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+    http://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. */
+
+#include "paddle/fluid/framework/eigen.h"
+#include "paddle/fluid/framework/op_registry.h"
+#include "paddle/fluid/platform/cuda_primitives.h"
+
+#include "util.cu.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+
+template <typename T>
+__global__ void GatherPointKernel(int b,
+                                  int n,
+                                  int m,
+                                  const T *__restrict__ inp,
+                                  const int *__restrict__ idx,
+                                  T *__restrict__ out) {
+  for (int i = blockIdx.x; i < b; i += gridDim.x) {
+    for (int j = blockIdx.y * blockDim.x + threadIdx.x; j < m;
+         j += blockDim.x * gridDim.y) {
+      int a = idx[i * m + j];
+      for (int k = 0; k < 3; k++) {
+        out[(i * m + j) * 3 + k] = inp[(i * n + a) * 3 + k];
+      }
+    }
+  }
+}
+
+template <typename T>
+__global__ void GatherPointGradKernel(int b,
+                                      int n,
+                                      int m,
+                                      const T *__restrict__ out_grad,
+                                      const int *__restrict__ idx,
+                                      T *__restrict__ in_grad) {
+  for (int i = blockIdx.x; i < b; i += gridDim.x) {
+    for (int j = blockIdx.y * blockDim.x + threadIdx.x; j < m;
+         j += blockDim.x * gridDim.y) {
+      int a = idx[i * m + j];
+      const T *out_grad_pos = &out_grad[(i * m + j) * 3];
+      T *in_grad_pos = &in_grad[(i * n + a) * 3];
+      for (int k = 0; k < 3; k++) {
+        platform::CudaAtomicAdd(&in_grad_pos[k], out_grad_pos[k]);
+      }
+    }
+  }
+}
+
+template <typename T>
+class GatherPointOpCUDAKernel : public framework::OpKernel<T> {
+public:
+  void Compute(const framework::ExecutionContext &ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "This kernel only runs on GPU device.");
+    auto *points = ctx.Input<Tensor>("X");
+    auto *index = ctx.Input<Tensor>("Index");
+    auto *output = ctx.Output<Tensor>("Output");
+
+    if (points->numel() == 0) return;
+
+    const T *p_points = points->data<T>();
+    const int *p_index = index->data<int>();
+    T *p_out_points = output->mutable_data<T>(ctx.GetPlace());
+
+    int batch_size = points->dims()[0];
+    int n_points = points->dims()[1];
+    int m_points = index->dims()[1];
+
+    GatherPointKernel<<<dim3(2, 8, 1), 512>>>(
+        batch_size, n_points, m_points, p_points, p_index, p_out_points);
+  }
+};
+
+template <typename T>
+class GatherPointGradOpCUDAKernel : public framework::OpKernel<T> {
+public:
+  void Compute(const framework::ExecutionContext &ctx) const override {
+    auto *points = ctx.Input<Tensor>("X");
+    auto *index = ctx.Input<Tensor>("Index");
+    auto *output_grad = ctx.Input<Tensor>(framework::GradVarName("Output"));
+    auto *points_grad = ctx.Output<Tensor>(framework::GradVarName("X"));
+
+    if (points->numel() == 0) return;
+
+    const T *p_output_grad = output_grad->data<T>();
+    const int *p_index = index->data<int>();
+    T *p_points_grad = points_grad->mutable_data<T>(ctx.GetPlace());
+    int pnum = points_grad->numel();
+
+    auto &dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
+    Zero<<<(pnum + 512 - 1) / 512, 512, 0, dev_ctx.stream()>>>(p_points_grad,
+                                                               pnum);
+
+    int batch_size = points->dims()[0];
+    int n_points = points->dims()[1];
+    int m_points = index->dims()[1];
+
+    GatherPointGradKernel<<<dim3(2, 8, 1), 512, 0, dev_ctx.stream()>>>(
+        batch_size, n_points, m_points, p_output_grad, p_index, p_points_grad);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_CUDA_KERNEL(gather_point,
+                        ops::GatherPointOpCUDAKernel<float>,
+                        ops::GatherPointOpCUDAKernel<double>,
+                        ops::GatherPointOpCUDAKernel<int>);
+REGISTER_OP_CUDA_KERNEL(gather_point_grad,
+                        ops::GatherPointGradOpCUDAKernel<float>,
+                        ops::GatherPointGradOpCUDAKernel<double>,
+                        ops::GatherPointGradOpCUDAKernel<int>);

PaddleCV/Paddle3D/PointNet++/ext_op/src/group_points_op.cc

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+   http://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. */
+
+#include <memory>
+#include <string>
+#include <vector>
+#include "paddle/fluid/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+using framework::Tensor;
+
+class GroupPointsOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"),
+                   "Input(X) of GroupPointsOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("Idx"),
+                   "Input(Idx) of GroupPointsOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasOutput("Out"),
+                   "Output(Out) of GroupPointsOp should not be null.");
+
+    auto dim_x = ctx->GetInputDim("X");  // [B, C, N]
+    PADDLE_ENFORCE_EQ(dim_x.size(), 3, "X's dimension must be 3");
+
+    auto dim_idx = ctx->GetInputDim("Idx");  // [B, npoints, nsample]
+    PADDLE_ENFORCE_EQ(dim_idx.size(), 3, "Idx's dimension must be 3");
+
+    PADDLE_ENFORCE_EQ(dim_x[0], dim_idx[0],
+                      "X and Idx dim[0] should be equal.");
+
+    // output: [B, C, M, S]
+    std::vector<int64_t> dim_out({dim_x[0], dim_x[1], dim_idx[1], dim_idx[2]});
+    ctx->SetOutputDim("Out", framework::make_ddim(dim_out));
+  }
+
+ protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::OpKernelType(ctx.Input<Tensor>("X")->type(),
+                                   ctx.GetPlace());
+  }
+};
+
+class GroupPointsOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  void Make() override {
+    AddInput("X",
+             "The input tensor of group_points operator. "
+             "This is a 3-D tensor with shape of [B, C, N].");
+    AddInput("Idx",
+             "The input tensor of nearest neighbor index of group_points "
+             "operator. This is a 3-D tensor with shape of [B, M, S].");
+    AddOutput("Out",
+              "The output tensor of group_points operator. "
+              "This is a 4-D tensor with shape of [B, C, M, S].");
+
+    AddComment(R"DOC(
+          This operator group input points with index.
+         )DOC");
+  }
+};
+
+class GroupPointsOpGrad : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("Idx"), "Input(Idx) should not be null");
+    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")),
+                   "Input(Out@GRAD) should not be null");
+    auto dim_x = ctx->GetInputDim("X");
+    if (ctx->HasOutput(framework::GradVarName("X"))) {
+      ctx->SetOutputDim(framework::GradVarName("X"), dim_x);
+    }
+  }
+
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::OpKernelType(
+        ctx.Input<Tensor>(framework::GradVarName("Out"))->type(),
+        ctx.GetPlace());
+  }
+};
+
+template <typename T>
+class GroupPointsGradDescMaker : public framework::SingleGradOpMaker<T> {
+ public:
+  using framework::SingleGradOpMaker<T>::SingleGradOpMaker;
+
+ protected:
+  std::unique_ptr<T> Apply() const override {
+    auto* op = new T();
+    op->SetType("group_points_grad");
+    op->SetInput("X", this->Input("X"));
+    op->SetInput("Idx", this->Input("Idx"));
+    op->SetInput(framework::GradVarName("Out"), this->OutputGrad("Out"));
+    op->SetOutput(framework::GradVarName("X"), this->InputGrad("X"));
+    op->SetAttrMap(this->Attrs());
+    return std::unique_ptr<T>(op);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OPERATOR(group_points, ops::GroupPointsOp, ops::GroupPointsOpMaker,
+                  ops::GroupPointsGradDescMaker<paddle::framework::OpDesc>,
+                  ops::GroupPointsGradDescMaker<paddle::imperative::OpBase>);
+REGISTER_OPERATOR(group_points_grad, ops::GroupPointsOpGrad);

PaddleCV/Paddle3D/PointNet++/ext_op/src/group_points_op.cu

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+   http://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. */
+
+#include "paddle/fluid/framework/op_registry.h"
+#include "paddle/fluid/platform/cuda_primitives.h"
+
+#include "util.cu.h"
+
+#define TOTAL_THREADS 1024
+#define THREADS_PER_BLOCK 256
+#define DIVUP(m, n) ((m) / (n) + ((m) % (n) > 0))
+
+namespace paddle {
+namespace operators {
+
+using framework::Tensor;
+
+template <typename T>
+__global__ void KeGroupPointsFW(int b, int c, int n, int npoints, int nsample,
+                                const T* __restrict__ points,
+                                const int* __restrict__ idx,
+                                T* __restrict__ out) {
+  // points: (B, C, N)
+  // idx: (B, npoints, nsample)
+  // output:
+  //      out: (B, C, npoints, nsample)
+  int bs_idx = blockIdx.z;
+  int c_idx = blockIdx.y;
+  int index = blockIdx.x * blockDim.x + threadIdx.x;
+  int pt_idx = index / nsample;
+  if (bs_idx >= b || c_idx >= c || pt_idx >= npoints) return;
+
+  int sample_idx = index % nsample;
+
+  idx += bs_idx * npoints * nsample + pt_idx * nsample + sample_idx;
+  int in_idx = bs_idx * c * n + c_idx * n + idx[0];
+  int out_idx = bs_idx * c * npoints * nsample + c_idx * npoints * nsample +
+                pt_idx * nsample + sample_idx;
+
+  out[out_idx] = points[in_idx];
+}
+
+template <typename T>
+
+__global__ void KeGroupPointsBW(int b, int c, int n, int npoints, int nsample,
+                                const T* __restrict__ grad_out,
+                                const int* __restrict__ idx,
+                                T* __restrict__ grad_points) {
+  // grad_out: (B, C, npoints, nsample)
+  // idx: (B, npoints, nsample)
+  // output:
+  //      grad_points: (B, C, N)
+  int bs_idx = blockIdx.z;
+  int c_idx = blockIdx.y;
+  int index = blockIdx.x * blockDim.x + threadIdx.x;
+  int pt_idx = index / nsample;
+  if (bs_idx >= b || c_idx >= c || pt_idx >= npoints) return;
+
+  int sample_idx = index % nsample;
+  grad_out += bs_idx * c * npoints * nsample + c_idx * npoints * nsample +
+              pt_idx * nsample + sample_idx;
+  idx += bs_idx * npoints * nsample + pt_idx * nsample + sample_idx;
+
+  platform::CudaAtomicAdd(grad_points + bs_idx * c * n + c_idx * n + idx[0],
+                          grad_out[0]);
+}
+
+template <typename T>
+class GroupPointsOpCUDAKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "This kernel only runs on GPU device.");
+    auto* input = ctx.Input<Tensor>("X");
+    auto* idx = ctx.Input<Tensor>("Idx");
+    auto* output = ctx.Output<Tensor>("Out");
+    auto* input_data = input->data<T>();
+    auto* idx_data = idx->data<int>();
+
+    const int b = input->dims()[0];
+    const int c = input->dims()[1];
+    const int n = input->dims()[2];
+    const int m = idx->dims()[1];
+    const int s = idx->dims()[2];
+
+    auto* output_data = output->mutable_data<T>({b, c, m, s}, ctx.GetPlace());
+
+    dim3 blocks(DIVUP(m * s, THREADS_PER_BLOCK), c, b);
+    dim3 threads(THREADS_PER_BLOCK);
+    KeGroupPointsFW<
+        T><<<blocks, threads, 0, ctx.cuda_device_context().stream()>>>(
+        b, c, n, m, s, input_data, idx_data, output_data);
+  }
+};
+
+template <typename T>
+class GroupPointsGradOpCUDAKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    auto* input = ctx.Input<Tensor>("X");
+    auto* idx = ctx.Input<Tensor>("Idx");
+    auto* output_grad = ctx.Input<Tensor>(framework::GradVarName("Out"));
+    auto* input_grad = ctx.Output<Tensor>(framework::GradVarName("X"));
+    auto* idx_data = idx->data<int>();
+    auto output_grad_data = output_grad->data<T>();
+
+    const int b = input->dims()[0];
+    const int c = input->dims()[1];
+    const int n = input->dims()[2];
+    const int m = idx->dims()[1];
+    const int s = idx->dims()[2];
+
+    auto* input_grad_data =
+        input_grad->mutable_data<T>({b, c, n}, ctx.GetPlace());
+    auto& dev_ctx =
+        ctx.template device_context<platform::CUDADeviceContext>();
+    int pnum = input_grad->numel();
+    Zero<<<(pnum + 512 - 1) / 512, 512, 0, dev_ctx.stream()>>>(input_grad_data,
+                                                               pnum);
+
+    dim3 blocks(DIVUP(m * s, THREADS_PER_BLOCK), c, b);
+    dim3 threads(THREADS_PER_BLOCK);
+
+    KeGroupPointsBW<<<blocks, threads, 0, dev_ctx.stream()>>>(
+        b, c, n, m, s, output_grad_data, idx_data, input_grad_data);
+  }
+};
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_CUDA_KERNEL(group_points, ops::GroupPointsOpCUDAKernel<float>,
+                        ops::GroupPointsOpCUDAKernel<double>);
+REGISTER_OP_CUDA_KERNEL(group_points_grad,
+                        ops::GroupPointsGradOpCUDAKernel<float>,
+                        ops::GroupPointsGradOpCUDAKernel<double>);

PaddleCV/Paddle3D/PointNet++/ext_op/src/make.sh

+include_dir=$( python -c 'import paddle; print(paddle.sysconfig.get_include())' )
+lib_dir=$( python -c 'import paddle; print(paddle.sysconfig.get_lib())' )
+
+echo $include_dir
+echo $lib_dir
+
+OPS='farthest_point_sampling_op gather_point_op group_points_op query_ball_op three_interp_op three_nn_op'
+for op in ${OPS}
+do
+nvcc ${op}.cu -c -o ${op}.cu.o -ccbin cc -DPADDLE_WITH_CUDA -DEIGEN_USE_GPU -DPADDLE_USE_DSO -DPADDLE_WITH_MKLDNN -Xcompiler -fPIC -std=c++11 -Xcompiler -fPIC -w --expt-relaxed-constexpr -O0 -g -DNVCC \
+    -I ${include_dir}/third_party/ \
+    -I ${include_dir}
+done
+
+g++ farthest_point_sampling_op.cc farthest_point_sampling_op.cu.o gather_point_op.cc gather_point_op.cu.o group_points_op.cc group_points_op.cu.o query_ball_op.cu.o query_ball_op.cc three_interp_op.cu.o three_interp_op.cc three_nn_op.cu.o three_nn_op.cc -o pointnet_lib.so -DPADDLE_WITH_MKLDNN -shared -fPIC -std=c++11 -O0 -g \
+  -I ${include_dir}/third_party/ \
+  -I ${include_dir} \
+  -L ${lib_dir} \
+  -L /usr/local/cuda/lib64 -lpaddle_framework -lcudart
+
+rm *.cu.o

PaddleCV/Paddle3D/PointNet++/ext_op/src/query_ball_op.cc

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+    http://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. */
+
+#include "paddle/fluid/framework/op_registry.h"
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+
+class QueryBallOp : public framework::OperatorWithKernel {
+public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+  void InferShape(framework::InferShapeContext *ctx) const override {
+    // points: [b,n,3]
+    PADDLE_ENFORCE(ctx->HasInput("Points"), "Input(Points) shoud not be null");
+    auto p_dims = ctx->GetInputDim("Points");
+    PADDLE_ENFORCE(p_dims.size() == 3 && p_dims[2] == 3,
+                   "Input(Points) of QueryBallOp should be 3-D Tensor, the "
+                   "last dimension must be 3");
+    // new_points: [b,m,3]
+    PADDLE_ENFORCE(ctx->HasInput("New_Points"),
+                   "Input(New_Points) shoud not be null");
+    auto np_dims = ctx->GetInputDim("New_Points");
+    PADDLE_ENFORCE(np_dims.size() == 3 && np_dims[2] == 3,
+                   "Input(New_Points) of QueryBallOp should be 3-D Tensor, the "
+                   "last dimension must be 3");
+    int n_sample = ctx->Attrs().Get<int>("N_sample");
+    PADDLE_ENFORCE(n_sample >= 0,
+                   "The n_sample should be greater than or equal to 0.");
+    float radius = ctx->Attrs().Get<float>("Radius");
+    PADDLE_ENFORCE(radius >= 0,
+                   "The radius should be greater than or equal to 0.");
+    // output: [b,m,nsample]
+    std::vector<int64_t> dim_out({p_dims[0], np_dims[1], n_sample});
+    ctx->SetOutputDim("Output", framework::make_ddim(dim_out));
+  }
+
+protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext &ctx) const override {
+    auto input_data_type = ctx.Input<Tensor>("Points")->type();
+    return framework::OpKernelType(input_data_type, ctx.GetPlace());
+  }
+};
+
+class QueryBallOpMaker : public framework::OpProtoAndCheckerMaker {
+public:
+  void Make() override {
+    AddInput("Points",
+             "Input points with shape (batch, n, 3), n is input "
+             "points's num");
+    AddInput("New_Points",
+             "Query points with shape (batch, m, 3), m is query points's num");
+    AddOutput("Output", "output points with shape(batch, m, nsample)");
+    AddAttr<int>("N_sample",
+                 R"Doc(Number of points selected in each ball region")Doc")
+        .SetDefault(0)
+        .EqualGreaterThan(0);
+    AddAttr<float>("Radius",
+                   R"Doc(Ball search radius with shape(1))Doc")
+        .SetDefault(0)
+        .EqualGreaterThan(0);
+
+    AddComment(
+        R"Doc(Query Ball Points)Doc");
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OPERATOR(query_ball, ops::QueryBallOp, ops::QueryBallOpMaker);

PaddleCV/Paddle3D/PointNet++/ext_op/src/query_ball_op.cu

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+    http://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. */
+
+#include "paddle/fluid/framework/eigen.h"
+#include "paddle/fluid/framework/op_registry.h"
+
+#include "util.cu.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+
+template <typename T>
+// input: radius (1), nsample (1), points (b,n,3), new_points (b,m,3)
+// output: idx (b,m,nsample)
+__global__ void QueryBall(int b,
+                          int n,
+                          int m,
+                          T radius,
+                          int nsample,
+                          const T *points,
+                          const T *new_points,
+                          int *idx) {
+  int batch_index = blockIdx.x;
+  points += n * 3 * batch_index;
+  new_points += m * 3 * batch_index;
+  idx += m * nsample * batch_index;
+
+  int index = threadIdx.x;
+  int stride = blockDim.x;
+
+  for (int j = index; j < m; j += stride) {
+    int cnt = 0;
+    for (int k = 0; k < n; ++k) {
+      if (cnt == nsample)
+        break;  // only pick the FIRST nsample points in the ball
+      float x2 = new_points[j * 3 + 0];
+      float y2 = new_points[j * 3 + 1];
+      float z2 = new_points[j * 3 + 2];
+      float x1 = points[k * 3 + 0];
+      float y1 = points[k * 3 + 1];
+      float z1 = points[k * 3 + 2];
+      float d =
+          (x2 - x1) * (x2 - x1) + (y2 - y1) * (y2 - y1) + (z2 - z1) * (z2 - z1);
+      if (d < radius * radius) {
+        if (cnt == 0) {  // set ALL indices to k, s.t. if there are less points
+                         // in ball than nsample, we still have valid
+                         // (repeating) indices
+          for (int l = 0; l < nsample; ++l) idx[j * nsample + l] = k;
+        }
+        idx[j * nsample + cnt] = k;
+        cnt += 1;
+      }
+    }
+  }
+}
+
+template <typename T>
+class QueryBallOpCUDAKernel : public framework::OpKernel<T> {
+public:
+  void Compute(const framework::ExecutionContext &ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "This kernel only runs on GPU device.");
+    // input: radius (1), nsample (1), points (b,n,3), new_points (b,m,3)
+    // output: idx (b,m,nsample)
+    auto *points = ctx.Input<Tensor>("Points");
+    auto *new_points = ctx.Input<Tensor>("New_Points");
+    auto *output = ctx.Output<Tensor>("Output");
+
+    float radius = ctx.Attr<T>("Radius");
+    int nsample = ctx.Attr<int>("N_sample");
+
+    if (points->numel() == 0 || new_points->numel() == 0) return;
+
+    int batch_size = points->dims()[0];
+    int n = points->dims()[1];
+    int m = new_points->dims()[1];
+    // allocate memory
+    int* p_out_points = output->mutable_data<int>({batch_size, m, nsample}, ctx.GetPlace());
+
+    auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
+    int pnum = output->numel();
+    Zero<int><<<(pnum + 512 - 1) / 512, 512, 0, dev_ctx.stream()>>>(p_out_points,
+                                                               pnum);
+
+    const T *p_points = points->data<T>();
+    const T *p_new_points = new_points->data<T>();
+
+    QueryBall<<<batch_size, 256>>>(batch_size,
+                                   n,
+                                   m,
+                                   radius,
+                                   nsample,
+                                   p_points,
+                                   p_new_points,
+                                   p_out_points);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_CUDA_KERNEL(query_ball, ops::QueryBallOpCUDAKernel<float>);

PaddleCV/Paddle3D/PointNet++/ext_op/src/three_interp_op.cc

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+   http://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. */
+
+#include <memory>
+#include <string>
+#include <vector>
+#include "paddle/fluid/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+using framework::Tensor;
+
+class ThreeInterpOp : public framework::OperatorWithKernel {
+public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"),
+                   "Input(X) of ThreeInterpOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("Weight"),
+                   "Input(Weight) of ThreeInterpOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("Idx"),
+                   "Input(Idx) of ThreeInterpOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasOutput("Out"),
+                   "Output(Out) of ThreeInterpOp should not be null.");
+
+    auto dim_x = ctx->GetInputDim("X");  // [B, M, C]
+    PADDLE_ENFORCE_EQ(dim_x.size(), 3, "X's dimension must be 3");
+
+    auto dim_weight = ctx->GetInputDim("Weight");  // [B, N, 3]
+    PADDLE_ENFORCE_EQ(dim_weight.size(), 3, "Weight's dimension must be 3");
+
+    PADDLE_ENFORCE_EQ(
+        dim_x[0], dim_weight[0], "X and Weight dim[0] should be equal.");
+
+    auto dim_idx = ctx->GetInputDim("Idx");  // [B, N, 3]
+    PADDLE_ENFORCE_EQ(dim_idx.size(), 3, "Idx's dimension must be 3");
+
+    for (int i = 0; i < 3; i++) {
+      PADDLE_ENFORCE_EQ(
+          dim_weight[i], dim_idx[i], "Weight and Idx shape should be same.");
+    }
+
+    // output: [B, N, C]
+    std::vector<int64_t> dim_out({dim_x[0], dim_idx[1], dim_x[2]});
+    ctx->SetOutputDim("Out", framework::make_ddim(dim_out));
+  }
+
+protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::OpKernelType(ctx.Input<Tensor>("X")->type(),
+                                   ctx.GetPlace());
+  }
+};
+
+class ThreeInterpOpMaker : public framework::OpProtoAndCheckerMaker {
+public:
+  void Make() override {
+    AddInput("X",
+             "The input tensor of three_interp operator. "
+             "This is a 3-D tensor with shape of [B, M, C].");
+    AddInput("Weight",
+             "The input tensor of point weight of three_interp operator. "
+             "This is a 3-D tensor with shape of [B, N, 3].");
+    AddInput("Idx",
+             "The input tensor of nearest neighbor index of three_interp "
+             "operator. This is a 3-D tensor with shape of [B, N, 3].");
+    AddOutput("Out",
+              "The output tensor of three_interp operator. "
+              "This is a 3-D tensor with shape of [B, N, 3].");
+
+    AddComment(R"DOC(
+          This operator calculate interpolate results from input, weight and
+          index.
+         )DOC");
+  }
+};
+
+class ThreeInterpOpGrad : public framework::OperatorWithKernel {
+public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("Weight"), "Input(Weight) should not be null");
+    PADDLE_ENFORCE(ctx->HasInput("Idx"), "Input(Idx) should not be null");
+    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")),
+                   "Input(Out@GRAD) should not be null");
+    auto dim_x = ctx->GetInputDim("X");
+    if (ctx->HasOutput(framework::GradVarName("X"))) {
+      ctx->SetOutputDim(framework::GradVarName("X"), dim_x);
+    }
+  }
+
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::OpKernelType(
+        ctx.Input<Tensor>(framework::GradVarName("Out"))->type(),
+        ctx.GetPlace());
+  }
+};
+
+template <typename T>
+class ThreeInterpGradDescMaker : public framework::SingleGradOpMaker<T> {
+public:
+  using framework::SingleGradOpMaker<T>::SingleGradOpMaker;
+
+protected:
+  std::unique_ptr<T> Apply() const override {
+    auto* op = new T();
+    op->SetType("three_interp_grad");
+    op->SetInput("X", this->Input("X"));
+    op->SetInput("Weight", this->Input("Weight"));
+    op->SetInput("Idx", this->Input("Idx"));
+    op->SetInput(framework::GradVarName("Out"), this->OutputGrad("Out"));
+    op->SetOutput(framework::GradVarName("X"), this->InputGrad("X"));
+    op->SetAttrMap(this->Attrs());
+    return std::unique_ptr<T>(op);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OPERATOR(three_interp,
+                  ops::ThreeInterpOp,
+                  ops::ThreeInterpOpMaker,
+                  ops::ThreeInterpGradDescMaker<paddle::framework::OpDesc>,
+                  ops::ThreeInterpGradDescMaker<paddle::imperative::OpBase>);
+REGISTER_OPERATOR(three_interp_grad, ops::ThreeInterpOpGrad);

PaddleCV/Paddle3D/PointNet++/ext_op/src/three_interp_op.cu

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+   http://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. */
+
+#include "paddle/fluid/framework/op_registry.h"
+#include "paddle/fluid/platform/cuda_primitives.h"
+
+#include "util.cu.h"
+
+namespace paddle {
+namespace operators {
+
+using framework::Tensor;
+
+template <typename T>
+__global__ void KeThreeInterpFw(T* output,
+                                const T* input,
+                                const T* weight,
+                                const int* idx,
+                                const int b,
+                                const int m,
+                                const int c,
+                                const int n) {
+  int nthreads = b * n * c;
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (; tid < nthreads; tid += stride) {
+    int bi = tid / n / c;
+    int ni = (tid % (n * c)) / c;
+    int ci = tid % c;
+
+    int input_base_idx = bi * m * c;
+    int w_idx = bi * n * 3 + ni * 3;
+    output[tid] =
+        input[input_base_idx + idx[w_idx] * c + ci] * weight[w_idx] +
+        input[input_base_idx + idx[w_idx + 1] * c + ci] * weight[w_idx + 1] +
+        input[input_base_idx + idx[w_idx + 2] * c + ci] * weight[w_idx + 2];
+  }
+}
+
+template <typename T>
+__global__ void KeThreeInterpBw(T* input_grad,
+                                const T* output_grad,
+                                const T* weight,
+                                const int* idx,
+                                const int b,
+                                const int m,
+                                const int c,
+                                const int n) {
+  int nthreads = b * n * c;
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (; tid < nthreads; tid += stride) {
+    int bi = tid / n / c;
+    int ni = (tid % (c * n)) / c;
+    int ci = tid % c;
+
+    int input_base_idx = bi * m * c;
+    int w_idx = bi * n * 3 + ni * 3;
+    platform::CudaAtomicAdd(&input_grad[input_base_idx + idx[w_idx] * c + ci],
+                            output_grad[tid] * weight[w_idx]);
+    platform::CudaAtomicAdd(
+        &input_grad[input_base_idx + idx[w_idx + 1] * c + ci],
+        output_grad[tid] * weight[w_idx + 1]);
+    platform::CudaAtomicAdd(
+        &input_grad[input_base_idx + idx[w_idx + 2] * c + ci],
+        output_grad[tid] * weight[w_idx + 2]);
+  }
+}
+
+template <typename T>
+class ThreeInterpOpCUDAKernel : public framework::OpKernel<T> {
+public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "This kernel only runs on GPU device.");
+    auto* input = ctx.Input<Tensor>("X");
+    auto* weight = ctx.Input<Tensor>("Weight");
+    auto* idx = ctx.Input<Tensor>("Idx");
+    auto* output = ctx.Output<Tensor>("Out");
+    auto* input_data = input->data<T>();
+    auto* weight_data = weight->data<T>();
+    auto* idx_data = idx->data<int>();
+
+    const int b = input->dims()[0];
+    const int m = input->dims()[1];
+    const int c = input->dims()[2];
+    const int n = weight->dims()[1];
+
+    auto* output_data = output->mutable_data<T>({b, n, c}, ctx.GetPlace());
+
+    int pixelNum = b * n * c;
+    int grid_dim = (pixelNum + 512 - 1) / 512;
+    grid_dim = grid_dim > 8 ? 8 : grid_dim;
+
+    KeThreeInterpFw<
+        T><<<grid_dim, 512, 0, ctx.cuda_device_context().stream()>>>(
+        output_data, input_data, weight_data, idx_data, b, m, c, n);
+  }
+};
+
+template <typename T>
+class ThreeInterpGradOpCUDAKernel : public framework::OpKernel<T> {
+public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    auto* input = ctx.Input<Tensor>("X");
+    auto* weight = ctx.Input<Tensor>("Weight");
+    auto* idx = ctx.Input<Tensor>("Idx");
+    auto* output_grad = ctx.Input<Tensor>(framework::GradVarName("Out"));
+    auto* input_grad = ctx.Output<Tensor>(framework::GradVarName("X"));
+    auto* weight_data = weight->data<T>();
+    auto* idx_data = idx->data<int>();
+    auto output_grad_data = output_grad->data<T>();
+
+    const int b = input->dims()[0];
+    const int m = input->dims()[1];
+    const int c = input->dims()[2];
+    const int n = weight->dims()[1];
+
+    auto* input_grad_data =
+        input_grad->mutable_data<T>({b, m, c}, ctx.GetPlace());
+    auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
+    int pnum = input_grad->numel();
+    Zero<<<(pnum + 512 - 1) / 512, 512, 0, dev_ctx.stream()>>>(input_grad_data,
+                                                               pnum);
+
+    int pixelNum = b * n * c;
+    int grid_dim = (pixelNum + 512 - 1) / 512;
+    grid_dim = grid_dim > 8 ? 8 : grid_dim;
+
+    KeThreeInterpBw<
+        T><<<grid_dim, 512, 0, ctx.cuda_device_context().stream()>>>(
+        input_grad_data, output_grad_data, weight_data, idx_data, b, m, c, n);
+  }
+};
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_CUDA_KERNEL(three_interp,
+                        ops::ThreeInterpOpCUDAKernel<float>,
+                        ops::ThreeInterpOpCUDAKernel<double>);
+REGISTER_OP_CUDA_KERNEL(three_interp_grad,
+                        ops::ThreeInterpGradOpCUDAKernel<float>,
+                        ops::ThreeInterpGradOpCUDAKernel<double>);

PaddleCV/Paddle3D/PointNet++/ext_op/src/three_nn_op.cc

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+   http://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. */
+
+#include <memory>
+#include <string>
+#include <vector>
+#include "paddle/fluid/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+using framework::Tensor;
+
+class ThreeNNOp : public framework::OperatorWithKernel {
+public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"),
+                   "Input(X) of ThreeNNOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("Known"),
+                   "Input(Known) of ThreeNNOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasOutput("Distance"),
+                   "Output(Distance) of ThreeNNOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasOutput("Idx"),
+                   "Output(Idx) of ThreeNNOp should not be null.");
+
+    auto dim_x = ctx->GetInputDim("X");  // [B, N, 3]
+    PADDLE_ENFORCE_EQ(dim_x.size(), 3, "X's dimension must be 3");
+    PADDLE_ENFORCE_EQ(dim_x[2], 3, "X dim[2] must be 3");
+
+    auto dim_known = ctx->GetInputDim("Known");  // [B, M, 3]
+    PADDLE_ENFORCE_EQ(dim_known.size(), 3, "Known's dimension must be 3");
+    PADDLE_ENFORCE_EQ(dim_known[2], 3, "Known dim[2] must be 3");
+
+    PADDLE_ENFORCE_EQ(
+        dim_x[0], dim_known[0], "X and Known dim[0] should be equal.");
+    PADDLE_ENFORCE_GE(
+        dim_known[1], 3, "Known dim[1] shoule be greater or euqal than 3.");
+
+    ctx->SetOutputDim("Distance", dim_x);
+    ctx->SetOutputDim("Idx", dim_x);
+  }
+
+protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::OpKernelType(ctx.Input<Tensor>("X")->type(),
+                                   ctx.GetPlace());
+  }
+};
+
+class ThreeNNOpMaker : public framework::OpProtoAndCheckerMaker {
+public:
+  void Make() override {
+    AddInput("X",
+             "The input tensor of three_nn operator. "
+             "This is a 3-D tensor with shape of [B, N, 3].");
+    AddInput("Known",
+             "The input tensor of known points of three_nn operator. "
+             "This is a 3-D tensor with shape of [B, M, 3].");
+    AddOutput("Distance",
+              "The output distance tensor of three_nn operator. "
+              "This is a 3-D tensor with shape of [B, N, 3].");
+    AddOutput("Idx",
+              "The output index tensor of three_nn operator. "
+              "This is a 3-D tensor with shape of [B, N, 3].");
+
+    AddAttr<float>("eps", "minimum value of distance.").SetDefault(1e-10);
+
+    AddComment(R"DOC(
+          This operator samples the top-3 nearest neighbor of each point
+          coordinates specified by Input(X) between known point coordinates
+          specified by Input(Known) and calcualte the distance between these
+          nearest neighbors.
+         )DOC");
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OPERATOR(three_nn, ops::ThreeNNOp, ops::ThreeNNOpMaker);

PaddleCV/Paddle3D/PointNet++/ext_op/src/three_nn_op.cu

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+   http://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. */
+
+#include "paddle/fluid/framework/op_registry.h"
+#include "paddle/fluid/platform/cuda_primitives.h"
+
+namespace paddle {
+namespace operators {
+
+using framework::Tensor;
+
+template <typename T>
+__global__ void KeThreeNNFw(T* distance,
+                            int* idx,
+                            const T* input,
+                            const T* known,
+                            const float eps,
+                            const int b,
+                            const int n,
+                            const int m) {
+  int nthreads = b * n;
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (; tid < nthreads; tid += stride) {
+    int bi = tid / n;
+    int ni = tid % n;
+
+    int input_idx = tid * 3;
+    T x1 = input[input_idx];
+    T y1 = input[input_idx + 1];
+    T z1 = input[input_idx + 2];
+
+    distance[input_idx] = 1e40;
+    distance[input_idx + 1] = 1e40;
+    distance[input_idx + 2] = 1e40;
+    idx[input_idx] = 0;
+    idx[input_idx + 1] = 0;
+    idx[input_idx + 2] = 0;
+    for (int i = 0; i < m; i++) {
+      int known_idx = bi * m * 3 + i * 3;
+      double dist = (x1 - known[known_idx]) * (x1 - known[known_idx]) +
+                    (y1 - known[known_idx + 1]) * (y1 - known[known_idx + 1]) +
+                    (z1 - known[known_idx + 2]) * (z1 - known[known_idx + 2]);
+      T valid_dist = dist > eps ? static_cast<T>(dist) : eps;
+      if (dist < distance[input_idx]) {
+        distance[input_idx + 2] = distance[input_idx + 1];
+        idx[input_idx + 2] = idx[input_idx + 1];
+        distance[input_idx + 1] = distance[input_idx];
+        idx[input_idx + 1] = idx[input_idx];
+        distance[input_idx] = dist;
+        idx[input_idx] = i;
+      } else if (dist < distance[input_idx + 1]) {
+        distance[input_idx + 2] = distance[input_idx + 1];
+        idx[input_idx + 2] = idx[input_idx + 1];
+        distance[input_idx + 1] = dist;
+        idx[input_idx + 1] = i;
+      } else if (dist < distance[input_idx + 2]) {
+        distance[input_idx + 2] = dist;
+        idx[input_idx + 2] = i;
+      }
+    }
+  }
+}
+
+template <typename T>
+class ThreeNNOpCUDAKernel : public framework::OpKernel<T> {
+public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "This kernel only runs on GPU device.");
+    auto* input = ctx.Input<Tensor>("X");
+    auto* known = ctx.Input<Tensor>("Known");
+    auto* distance = ctx.Output<Tensor>("Distance");
+    auto* idx = ctx.Output<Tensor>("Idx");
+    auto* input_data = input->data<T>();
+    auto* known_data = known->data<T>();
+
+    const float eps = ctx.Attr<float>("eps");
+
+    const int b = input->dims()[0];
+    const int n = input->dims()[1];
+    const int m = known->dims()[1];
+
+    auto* idx_data = idx->mutable_data<int>({b, n, 3}, ctx.GetPlace());
+    auto* distance_data = distance->mutable_data<T>({b, n, 3}, ctx.GetPlace());
+
+    int pixelNum = b * n;
+    int grid_dim = (pixelNum + 512 - 1) / 512;
+    grid_dim = grid_dim > 8 ? 8 : grid_dim;
+
+    KeThreeNNFw<T><<<grid_dim, 512, 0, ctx.cuda_device_context().stream()>>>(
+        distance_data, idx_data, input_data, known_data, eps, b, n, m);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_CUDA_KERNEL(three_nn,
+                        ops::ThreeNNOpCUDAKernel<float>,
+                        ops::ThreeNNOpCUDAKernel<double>);

PaddleCV/Paddle3D/PointNet++/ext_op/src/util.cu.h

+/* Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+   http://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. */
+
+template <typename T>
+__global__ void Zero(T* x, int num) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < num;
+       i += blockDim.x * gridDim.x) {
+    x[i] = static_cast<T>(0);
+  }
+}

PaddleCV/Paddle3D/PointNet++/ext_op/tests/test_farthest_point_sampling_op.py

+# Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://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.
+
+from __future__ import print_function
+
+import unittest
+import numpy as np
+import paddle.fluid as fluid
+import pointnet_lib
+
+
+def farthest_point_sampling_np(xyz, npoint):
+    B, N, C = xyz.shape
+    S = npoint
+
+    centroids = np.zeros((B, S))
+    distance = np.ones((B, N)) * 1e10
+    farthest = 0
+    batch_indices = np.arange(B).astype('int32')
+    for i in range(S):
+        centroids[:, i] = farthest
+        centroid = xyz[batch_indices, farthest, :].reshape((B, 1, 3))
+        dist = np.sum((xyz - centroid)**2, -1)
+        mask = dist < distance
+        distance[mask] = dist[mask]
+        farthest = np.argmax(distance, -1)
+    return centroids.astype('int32')
+
+
+class TestFarthestPointSamplingOp(unittest.TestCase):
+    def test_check_output(self):
+        x_shape = (1, 512, 3)
+        x_type = 'float32'
+        sampled_point_num = 256
+
+        x = fluid.layers.data(
+            name='x', shape=x_shape, dtype=x_type, append_batch_size=False)
+        y = pointnet_lib.farthest_point_sampling(x, sampled_point_num)
+
+        x_np = np.random.randint(1, 100, (x_shape[0] * x_shape[1] *
+                                          3, )).reshape(x_shape).astype(x_type)
+        out_np = farthest_point_sampling_np(x_np, sampled_point_num)
+
+        place = fluid.CUDAPlace(0)
+        exe = fluid.Executor(place)
+        outs = exe.run(feed={'x': x_np}, fetch_list=[y])
+
+        self.assertTrue(np.allclose(outs[0], out_np))
+
+
+if __name__ == "__main__":
+    unittest.main()

PaddleCV/Paddle3D/PointNet++/ext_op/tests/test_gather_point_op.py

+#   Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://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.
+
+from __future__ import print_function
+
+import unittest
+import numpy as np
+import paddle.fluid as fluid
+import pointnet_lib
+
+
+def gather_point_np(points, index):
+    result = []
+    for i in range(len(index)):
+        a = points[i][index[i]]
+        result.append(a.tolist())
+    return result
+
+
+class TestGatherPointOp(unittest.TestCase):
+    def test_check_output(self):
+        x_shape = (1, 512, 3)
+        x_type = 'float32'
+        idx_shape = (1, 32)
+        idx_type = 'int32'
+
+        x = fluid.layers.data(
+            name='x', shape=x_shape, dtype=x_type, append_batch_size=False)
+        idx = fluid.layers.data(
+            name='idx', shape=idx_shape, dtype=idx_type, append_batch_size=False)
+        y = pointnet_lib.gather_point(x, idx)
+
+        x_np = np.random.uniform(-10, 10, x_shape).astype(x_type)
+        idx_np = np.random.randint(0, x_shape[1], idx_shape).astype(idx_type)
+        out_np = gather_point_np(x_np, idx_np)
+
+        place = fluid.CUDAPlace(0)
+        exe = fluid.Executor(place)
+        outs = exe.run(feed={'x': x_np, 'idx': idx_np}, fetch_list=[y])
+
+        self.assertTrue(np.allclose(outs[0], out_np))
+
+
+if __name__ == "__main__":
+    unittest.main()

PaddleCV/Paddle3D/PointNet++/ext_op/tests/test_group_points_op.py

+#   Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://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.
+
+from __future__ import print_function
+
+import unittest
+import numpy as np
+import paddle.fluid as fluid
+import pointnet_lib
+
+
+def group_points_np(x, idx):
+    b, m, s = idx.shape
+    _, c, n = x.shape
+
+    output = np.zeros((b, c, m, s)).astype(x.dtype)
+    for i in range(b):
+        for j in range(m):
+            for k in range(s):
+                output[i, :, j, k] = x[i, :, idx[i, j, k]]
+    return output
+
+
+class TestGroupPointsOp(unittest.TestCase):
+    def test_check_output(self):
+        x_shape = [8, 43, 29]
+        x_type = 'float32'
+        idx_shape = [8, 37, 41]
+        idx_type = 'int32'
+
+        x = fluid.layers.data(
+            name='x', shape=x_shape, dtype=x_type, append_batch_size=False)
+        idx = fluid.layers.data(
+            name='idx', shape=idx_shape, dtype=idx_type, append_batch_size=False)
+        y = pointnet_lib.group_points(x, idx)
+
+        x_np = np.random.uniform(-10, 10, x_shape).astype(x_type)
+        idx_np = np.random.randint(0, x_shape[2], idx_shape).astype(idx_type)
+        out_np = group_points_np(x_np, idx_np)
+
+        place = fluid.CUDAPlace(0)
+        exe = fluid.Executor(place)
+        outs = exe.run(feed={'x': x_np, 'idx': idx_np}, fetch_list=[y])
+
+        self.assertTrue(np.allclose(outs[0], out_np))
+
+
+if __name__ == "__main__":
+    unittest.main()

PaddleCV/Paddle3D/PointNet++/ext_op/tests/test_query_ball_op.py

+#   Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://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.
+
+from __future__ import print_function
+
+import unittest
+import numpy as np
+import paddle.fluid as fluid
+import pointnet_lib
+
+
+def query_ball_point_np(points, new_points, radius, nsample):
+    b, n, c = points.shape
+    _, m, _ = new_points.shape
+    out = np.zeros(shape=(b, m, nsample)).astype('int32')
+    radius_2 = radius * radius
+    for i in range(b):
+        for j in range(m):
+            cnt = 0
+            for k in range(n):
+                if (cnt == nsample):
+                    break
+                dist = np.sum(np.square(points[i][k] - new_points[i][j]))
+                if (dist < radius_2):
+                    if cnt == 0:
+                        out[i][j] = np.ones(shape=(nsample)) * k
+                    out[i][j][cnt] = k
+                    cnt += 1
+    return out
+
+
+class TestQueryBallOp(unittest.TestCase):
+    def test_check_output(self):
+        points_shape = [2, 5, 3]
+        new_points_shape = [2, 4, 3]
+        points_type = 'float32'
+        radius = 6
+        nsample = 5
+
+        points = fluid.layers.data(
+            name='points', shape=points_shape, dtype=points_type, append_batch_size=False)
+        new_points = fluid.layers.data(
+            name='new_points', shape=new_points_shape, dtype=points_type, append_batch_size=False)
+        y = pointnet_lib.query_ball(points, new_points, radius, nsample)
+
+        points_np = np.random.randint(1, 5, points_shape).astype(points_type)
+        new_points_np = np.random.randint(1, 5, new_points_shape).astype(points_type)
+        out_np = query_ball_point_np(points_np, new_points_np, radius, nsample)
+
+        place = fluid.CUDAPlace(0)
+        exe = fluid.Executor(place)
+        outs = exe.run(feed={'points': points_np, 'new_points': new_points_np}, fetch_list=[y])
+
+        self.assertTrue(np.allclose(outs[0], out_np))
+
+
+if __name__ == "__main__":
+    unittest.main()

PaddleCV/Paddle3D/PointNet++/ext_op/tests/test_three_interp_op.py

+#   Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://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.
+
+from __future__ import print_function
+
+import unittest
+import numpy as np
+import paddle.fluid as fluid
+import pointnet_lib
+
+
+def three_interp_np(x, weight, idx):
+    b, m, c = x.shape
+    n = weight.shape[1]
+
+    output = np.zeros((b, n, c)).astype('float32')
+    for i in range(b):
+        for j in range(n):
+            w1, w2, w3 = weight[i, j, :]
+            i1, i2, i3 = idx[i, j, :]
+            output[i, j, :] = w1 * x[i, i1, :] \
+                            + w2 * x[i, i2, :] \
+                            + w3 * x[i, i3, :]
+    return output
+
+
+class TestThreeInterpOp(unittest.TestCase):
+    def test_check_output(self):
+        input_shape = [8, 21, 29]
+        input_type = 'float32'
+        weight_shape = [8, 37, 3]
+        weight_type = 'float32'
+
+        x = fluid.layers.data(
+            name='x', shape=input_shape, dtype=input_type, append_batch_size=False)
+        weight = fluid.layers.data(
+            name='weight', shape=weight_shape, dtype=weight_type, append_batch_size=False)
+        idx = fluid.layers.data(
+            name='idx', shape=weight_shape, dtype="int32", append_batch_size=False)
+        y = pointnet_lib.three_interp(x, weight, idx)
+
+        x_np = np.random.random(input_shape).astype(input_type)
+        weight_np = np.random.random(weight_shape).astype(weight_type)
+        idx_np = np.random.uniform(0, input_shape[1], weight_shape).astype("int32")
+        out_np = three_interp_np(x_np, weight_np, idx_np)
+
+        place = fluid.CUDAPlace(0)
+        exe = fluid.Executor(place)
+        outs = exe.run(feed={'x': x_np, 'weight': weight_np, 'idx': idx_np}, fetch_list=[y])
+
+        self.assertTrue(np.allclose(outs[0], out_np))
+
+
+if __name__ == "__main__":
+    unittest.main()

PaddleCV/Paddle3D/PointNet++/ext_op/tests/test_three_nn_op.py

+#   Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://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.
+
+from __future__ import print_function
+
+import unittest
+import numpy as np
+import paddle.fluid as fluid
+import pointnet_lib
+
+
+def three_nn_np(x, known, eps=1e-10):
+    distance = np.ones_like(x).astype('float32') * 1e40
+    idx = np.zeros_like(x).astype('int32')
+
+    b, n, _ = x.shape
+    m = known.shape[1]
+    for i in range(b):
+        for j in range(n):
+            for k in range(m):
+                sub = x[i, j, :] - known[i, k, :]
+                d = float(np.sum(sub * sub))
+                valid_d = max(d, eps)
+                if d < distance[i, j, 0]:
+                    distance[i, j, 2] = distance[i, j, 1]
+                    idx[i, j, 2] = idx[i, j, 1]
+                    distance[i, j, 1] = distance[i, j, 0]
+                    idx[i, j, 1] = idx[i, j, 0]
+                    distance[i, j, 0] = valid_d
+                    idx[i, j, 0] = k
+                elif d < distance[i, j, 1]:
+                    distance[i, j, 2] = distance[i, j, 1]
+                    idx[i, j, 2] = idx[i, j, 1]
+                    distance[i, j, 1] = valid_d
+                    idx[i, j, 1] = k
+                elif d < distance[i, j, 2]:
+                    distance[i, j, 2] = valid_d
+                    idx[i, j, 2] = k
+    return distance, idx
+
+
+class TestThreeNNOp(unittest.TestCase):
+    def test_check_output(self):
+        input_shape = [16, 32, 3]
+        known_shape = [16, 8, 3]
+        input_type = 'float32'
+        eps = 1e-10
+
+        x = fluid.layers.data(
+            name='x', shape=input_shape, dtype=input_type, append_batch_size=False)
+        known = fluid.layers.data(
+            name='known', shape=known_shape, dtype=input_type, append_batch_size=False)
+        dist, idx = pointnet_lib.three_nn(x, known, eps)
+
+        x_np = np.random.random(input_shape).astype(input_type)
+        known_np = np.random.random(known_shape).astype(input_type)
+        dist_np, idx_np = three_nn_np(x_np, known_np, eps)
+
+        place = fluid.CUDAPlace(0)
+        exe = fluid.Executor(place)
+        outs = exe.run(feed={'x': x_np, 'known': known_np}, fetch_list=[dist, idx])
+
+        self.assertTrue(np.allclose(outs[0], dist_np))
+        self.assertTrue(np.allclose(outs[1], idx_np))
+
+
+if __name__ == "__main__":
+    unittest.main()

PaddleCV/Paddle3D/PointNet++/models/__init__.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+
+from __future__ import absolute_import
+
+from . import pointnet2_modules
+from . import pointnet2_seg
+from . import pointnet2_cls
+
+from .pointnet2_modules import *
+from .pointnet2_seg import *
+from .pointnet2_cls import *
+
+__all__ = pointnet2_modules.__all__
+__all__ += pointnet2_seg.__all__
+__all__ += pointnet2_cls.__all__

PaddleCV/Paddle3D/PointNet++/models/pointnet2_cls.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+"""
+Contains PointNet++ classification models
+"""
+
+from __future__ import absolute_import
+from __future__ import division
+from __future__ import print_function
+
+import numpy as np
+
+import paddle.fluid as fluid
+from paddle.fluid.param_attr import ParamAttr
+from paddle.fluid.initializer import Constant
+from .pointnet2_modules import *
+
+__all__ = ["PointNet2ClsSSG", "PointNet2ClsMSG"]
+
+
+class PointNet2Cls(object):
+    def __init__(self, num_classes, num_points, use_xyz=True):
+        self.num_classes = num_classes
+        self.num_points = num_points
+        self.use_xyz = use_xyz
+        self.out_feature = None
+        self.pyreader = None
+        self.model_config()
+
+    def model_config(self):
+        self.SA_confs = []
+
+    def build_input(self):
+        self.xyz = fluid.layers.data(name='xyz', shape=[self.num_points, 3], dtype='float32', lod_level=0)
+        self.label = fluid.layers.data(name='label', shape=[1], dtype='int64', lod_level=0)
+        self.pyreader = fluid.io.PyReader(
+                feed_list=[self.xyz, self.label],
+                capacity=64,
+                use_double_buffer=True,
+                iterable=False)
+        self.feed_vars = [self.xyz, self.label]
+
+    def build_model(self, bn_momentum=0.99):
+        self.build_input()
+
+        xyz, feature = self.xyz, None
+        for i, SA_conf in enumerate(self.SA_confs):
+            xyz, feature = pointnet_sa_module(
+                    xyz=xyz,
+                    feature=feature,
+                    bn_momentum=bn_momentum,
+                    use_xyz=self.use_xyz,
+                    name="sa_{}".format(i),
+                    **SA_conf)
+
+        out = fluid.layers.squeeze(feature, axes=[-1])
+        out = fc_bn(out,out_channels=512, bn=True, bn_momentum=bn_momentum, name="fc_1")
+        out = fluid.layers.dropout(out, 0.5, dropout_implementation="upscale_in_train")
+        out = fc_bn(out,out_channels=256, bn=True, bn_momentum=bn_momentum, name="fc_2")
+        out = fluid.layers.dropout(out, 0.5, dropout_implementation="upscale_in_train")
+        out = fc_bn(out,out_channels=self.num_classes, act=None, name="fc_3")
+        pred = fluid.layers.softmax(out)
+
+        # calc loss
+        self.loss = fluid.layers.cross_entropy(pred, self.label)
+        self.loss = fluid.layers.reduce_mean(self.loss)
+
+        # calc acc
+        pred = fluid.layers.reshape(pred, shape=[-1, self.num_classes])
+        label = fluid.layers.reshape(self.label, shape=[-1, 1])
+        self.acc1 = fluid.layers.accuracy(pred, label, k=1)
+
+    def get_feeds(self):
+        return self.feed_vars
+
+    def get_outputs(self):
+        return {"loss": self.loss, "accuracy": self.acc1}
+
+    def get_pyreader(self):
+        return self.pyreader
+
+
+class PointNet2ClsSSG(PointNet2Cls):
+    def __init__(self, num_classes, num_points, use_xyz=True):
+        super(PointNet2ClsSSG, self).__init__(num_classes, num_points, use_xyz)
+
+    def model_config(self):
+        self.SA_confs = [
+            {
+                "npoint": 512,
+                "radiuss": [0.2],
+                "nsamples": [64],
+                "mlps": [[64, 64, 128]],
+            },
+            {
+                "npoint": 128,
+                "radiuss": [0.4],
+                "nsamples": [64],
+                "mlps": [[128, 128, 256]],
+            },
+            {
+                "npoint":None,
+		"radiuss": [None],
+		"nsamples":[None],
+		"mlps": [[256, 512, 1024]],
+            },
+        ]
+
+
+class PointNet2ClsMSG(PointNet2Cls):
+    def __init__(self, num_classes, num_points, use_xyz=True):
+        super(PointNet2ClsMSG, self).__init__(num_classes, num_points, use_xyz)
+
+    def model_config(self):
+        self.SA_confs = [
+            {
+                "npoint": 512,
+                "radiuss": [0.1, 0.2, 0.4],
+                "nsamples": [16, 32, 128],
+                "mlps": [[32, 32, 64],
+                         [64, 64, 128],
+                         [64,96,128]],
+            },
+            {
+                "npoint": 128,
+                "radiuss": [0.2, 0.4, 0.8],
+                "nsamples": [32, 64, 128],
+                "mlps": [[64, 64, 128],
+                         [128, 128, 256],
+                         [128,128,256]],
+            },
+            {
+                "npoint":None,
+		"radiuss": [None],
+		"nsamples":[None],
+		"mlps": [[256, 512, 1024]],
+            },
+        ]
+
+

PaddleCV/Paddle3D/PointNet++/models/pointnet2_modules.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+"""
+Contains PointNet++  utility functions.
+"""
+
+from __future__ import absolute_import
+from __future__ import division
+from __future__ import print_function
+
+import numpy as np
+
+import paddle.fluid as fluid
+from paddle.fluid.param_attr import ParamAttr
+from paddle.fluid.initializer import Constant
+from ext_op import *
+
+__all__ = ["conv_bn", "pointnet_sa_module", "pointnet_fp_module","fc_bn"]
+
+
+def query_and_group(xyz, new_xyz, radius, nsample, features=None, use_xyz=True):
+    """
+    Perform query_ball and group_points
+
+    Args:
+        xyz (Variable): xyz coordiantes features with shape [B, N, 3]
+        new_xyz (Variable): centriods features with shape [B, npoint, 3]
+        radius (float32): radius of ball
+        nsample (int32): maximum number of gather features
+        features (Variable): features with shape [B, N, C]
+        use_xyz (bool): whether use xyz coordiantes features
+
+    Returns:
+        out (Variable): features with shape [B, npoint, nsample, C + 3]
+    """
+    idx = query_ball(xyz, new_xyz, radius, nsample)
+    idx.stop_gradient = True
+    xyz = fluid.layers.transpose(xyz,perm=[0, 2, 1])
+    grouped_xyz = group_points(xyz, idx)
+    expand_new_xyz = fluid.layers.unsqueeze(fluid.layers.transpose(new_xyz, perm=[0, 2, 1]), axes=[-1])
+    expand_new_xyz = fluid.layers.expand(expand_new_xyz, [1, 1, 1, grouped_xyz.shape[3]])
+    grouped_xyz -= expand_new_xyz
+
+    if features is not None:
+        grouped_features = group_points(features, idx)
+        return fluid.layers.concat([grouped_xyz, grouped_features], axis=1) \
+                if use_xyz else grouped_features
+    else:
+        assert use_xyz, "use_xyz should be True when features is None"
+        return grouped_xyz
+
+
+def group_all(xyz, features=None, use_xyz=True):
+    """
+    Group all xyz and features when npoint is None
+    See query_and_group
+    """
+    xyz = fluid.layers.transpose(xyz,perm=[0,2,1])
+    grouped_xyz = fluid.layers.unsqueeze(xyz, axes=[2])
+    if features is not None:
+        grouped_features = fluid.layers.unsqueeze(features, axes=[2])
+        return fluid.layers.concat([grouped_xyz, grouped_features], axis=1) if use_xyz else grouped_features
+    else:
+        return grouped_xyz
+
+
+def conv_bn(input, out_channels, bn=True, bn_momentum=0.99, act='relu', name=None):
+    param_attr = ParamAttr(name='{}_conv_weight'.format(name),)
+    bias_attr = ParamAttr(name='{}_conv_bias'.format(name)) \
+                                  if not bn else False
+    out = fluid.layers.conv2d(input,
+                              num_filters=out_channels,
+                              filter_size=1,
+                              stride=1,
+                              padding=0,
+                              dilation=1,
+                              param_attr=param_attr,
+                              bias_attr=bias_attr,
+			      act=act if not bn else None)
+    if bn:
+        bn_name = name + "_bn"
+        out = fluid.layers.batch_norm(out,
+                                      act=act,
+				      momentum=bn_momentum,
+                                      param_attr=ParamAttr(name=bn_name + "_scale"),
+                                      bias_attr=ParamAttr(name=bn_name + "_offset"),
+                                      moving_mean_name=bn_name + '_mean',
+                                      moving_variance_name=bn_name + '_var')
+
+    return out
+
+def fc_bn(input, out_channels, bn=False, bn_momentum=0.99, act='relu', name=None):
+    param_attr = ParamAttr(name='{}_fc_weight'.format(name))
+    if not bn:
+        bias_attr = ParamAttr(name='{}_fc_bias'.format(name))
+    else:
+        bias_attr = False
+    out = fluid.layers.fc(input,
+                          size=out_channels,
+			  param_attr=param_attr,
+			  bias_attr=bias_attr)
+    if bn:
+        bn_name = name + "_bn"
+        out = fluid.layers.batch_norm(out,
+                                      momentum=bn_momentum,
+                                      param_attr=ParamAttr(name=bn_name + "_scale"),
+                                      bias_attr=ParamAttr(name=bn_name + "_offset"),
+                                      moving_mean_name=bn_name + '_mean',
+                                      moving_variance_name=bn_name + '_var')
+    if act == "relu":
+        out = fluid.layers.relu(out)
+    return out
+
+def MLP(features, out_channels_list, bn=True, bn_momentum=0.99, act='relu', name=None):
+    out = features
+    for i, out_channels in enumerate(out_channels_list):
+        out = conv_bn(out, out_channels, bn=bn, act=act, bn_momentum=bn_momentum, name=name + "_{}".format(i))
+    return out
+        
+
+def pointnet_sa_module(xyz,
+                       npoint=None,
+                       radiuss=[],
+                       nsamples=[],
+                       mlps=[],
+                       feature=None,
+                       bn=True,
+		       bn_momentum=0.99,
+                       use_xyz=True,
+                       name=None):
+    """
+    PointNet MSG(Multi-Scale Group) Set Abstraction Module.
+    Call with radiuss, nsamples, mlps as single element list for 
+    SSG(Single-Scale Group).
+
+    Args:
+        xyz (Variable): xyz coordiantes features with shape [B, N, 3]
+        radiuss ([float32]): list of radius of ball
+        nsamples ([int32]): list of maximum number of gather features
+        mlps ([[int32]]): list of out_channels_list
+        feature (Variable): features with shape [B, C, N]
+        bn (bool): whether perform batch norm after conv2d
+	bn_momentum (float): momentum of batch norm
+        use_xyz (bool): whether use xyz coordiantes features
+
+    Returns:
+        new_xyz (Variable): centriods features with shape [B, npoint, 3]
+        out (Variable): features with shape [B, npoint, \sum_i{mlps[i][-1]}]
+    """
+    assert len(radiuss) == len(nsamples) == len(mlps), \
+            "radiuss, nsamples, mlps length should be same"
+
+    farthest_idx = farthest_point_sampling(xyz, npoint)
+    farthest_idx.stop_gradient = True
+    new_xyz = gather_point(xyz, farthest_idx) if npoint is not None else None
+
+    outs = []
+    for i, (radius, nsample, mlp) in enumerate(zip(radiuss, nsamples, mlps)):
+        out = query_and_group(xyz, new_xyz, radius, nsample, feature, use_xyz) if npoint is not None else group_all(xyz, feature, use_xyz)
+        out = MLP(out, mlp, bn=bn, bn_momentum=bn_momentum, name=name + '_mlp{}'.format(i))
+        out = fluid.layers.pool2d(out, pool_size=[1, out.shape[3]], pool_type='max')
+        out = fluid.layers.squeeze(out, axes=[-1])
+        outs.append(out)
+    out = fluid.layers.concat(outs, axis=1)
+
+    return (new_xyz, out)
+
+
+def pointnet_fp_module(unknown, known, unknown_feats, known_feats, mlp, bn=True, bn_momentum=0.99, name=None):
+    """
+    PointNet Feature Propagation Module
+
+    Args:
+        unknown (Variable): unknown xyz coordiantes features with shape [B, N, 3]
+        known (Variable): known xyz coordiantes features with shape [B, M, 3]
+        unknown_feats (Variable): unknown features with shape [B, N, C1] to be propagated to
+        known_feats (Variable): known features with shape [B, M, C2] to be propagated from
+        mlp ([int32]): out_channels_list
+        bn (bool): whether perform batch norm after conv2d
+        bn_momentum (float): momentum of batch norm
+
+    Returns:
+        new_features (Variable): new features with shape [B, N, mlp[-1]]
+    """
+    if known is None:
+        raise NotImplementedError("Not implement known as None currently.")
+    else:
+        dist, idx = three_nn(unknown, known, eps=0)
+        dist.stop_gradient = True
+        idx.stop_gradient = True
+        dist = fluid.layers.sqrt(dist)
+        ones = fluid.layers.fill_constant_batch_size_like(dist, dist.shape, dist.dtype, 1)
+        dist_recip = ones / (dist + 1e-8); # 1.0 / dist
+        norm = fluid.layers.reduce_sum(dist_recip, dim=-1, keep_dim=True)
+        weight = dist_recip / norm
+        weight.stop_gradient = True
+        interp_feats = three_interp(known_feats, weight, idx)
+
+    new_features = interp_feats if unknown_feats is None else \
+                    fluid.layers.concat([interp_feats, unknown_feats], axis=-1)
+    new_features = fluid.layers.transpose(new_features, perm=[0, 2, 1])
+    new_features = fluid.layers.unsqueeze(new_features, axes=[-1])
+    new_features = MLP(new_features, mlp, bn=bn, bn_momentum=bn_momentum, name=name + '_mlp')
+    new_features = fluid.layers.squeeze(new_features, axes=[-1])
+    new_features = fluid.layers.transpose(new_features, perm=[0, 2, 1])
+    
+    return new_features
+

PaddleCV/Paddle3D/PointNet++/models/pointnet2_seg.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+"""
+Contains PointNet++ SSG/MSG semantic segmentation models
+"""
+
+from __future__ import absolute_import
+from __future__ import division
+from __future__ import print_function
+
+import numpy as np
+
+import paddle.fluid as fluid
+from paddle.fluid.param_attr import ParamAttr
+from paddle.fluid.initializer import Constant
+from .pointnet2_modules import *
+
+__all__ = ["PointNet2SemSegSSG", "PointNet2SemSegMSG"]
+
+
+class PointNet2SemSeg(object):
+    def __init__(self, num_classes, num_points, use_xyz=True):
+        self.num_classes = num_classes
+        self.num_points = num_points
+        self.use_xyz = use_xyz
+        self.feed_vars = []
+        self.out_feature = None
+        self.pyreader = None
+        self.model_config()
+
+    def model_config(self):
+        self.SA_confs = []
+        self.FP_confs = []
+
+    def build_input(self):
+        self.xyz = fluid.layers.data(name='xyz', shape=[self.num_points, 3], dtype='float32', lod_level=0)
+        self.feature = fluid.layers.data(name='feature', shape=[self.num_points, 6], dtype='float32', lod_level=0)
+        self.label = fluid.layers.data(name='label', shape=[self.num_points, 1], dtype='int64', lod_level=0)
+        self.pyreader = fluid.io.PyReader(
+                feed_list=[self.xyz, self.feature, self.label],
+                capacity=64,
+                use_double_buffer=True,
+                iterable=False)
+        self.feed_vars = [self.xyz, self.feature, self.label]
+
+    def build_model(self, bn_momentum=0.99):
+        self.build_input()
+
+        xyzs, features = [self.xyz], [self.feature]
+        xyzi, featurei = xyzs[-1], fluid.layers.transpose(self.feature, perm=[0, 2, 1])
+        for i, SA_conf in enumerate(self.SA_confs):
+            xyzi, featurei = pointnet_sa_module(
+                    xyz=xyzi,
+                    feature=featurei,
+                    bn_momentum=bn_momentum,
+                    use_xyz=self.use_xyz,
+                    name="sa_{}".format(i),
+                    **SA_conf)
+            xyzs.append(xyzi)
+            features.append(fluid.layers.transpose(featurei, perm=[0, 2, 1]))
+        for i in range(-1, -(len(self.FP_confs) + 1), -1):
+            features[i - 1] = pointnet_fp_module(
+                    unknown=xyzs[i - 1],
+                    known=xyzs[i],
+                    unknown_feats=features[i - 1],
+                    known_feats=features[i],
+                    bn_momentum=bn_momentum,
+                    name="fp_{}".format(i+len(self.FP_confs)),
+                    **self.FP_confs[i])
+
+        out = fluid.layers.transpose(features[0], perm=[0, 2, 1])
+        out = fluid.layers.unsqueeze(out, axes=[-1])
+        out = conv_bn(out, out_channels=128, bn=True, bn_momentum=bn_momentum, name="output_1")
+        out = fluid.layers.dropout(out, 0.5, dropout_implementation="upscale_in_train")
+        out = conv_bn(out, out_channels=self.num_classes, bn=False, act=None, name="output_2")
+        out = fluid.layers.squeeze(out, axes=[-1])
+        out = fluid.layers.transpose(out, perm=[0, 2, 1])
+        pred = fluid.layers.softmax(out)
+
+        # calc loss
+        self.loss = fluid.layers.cross_entropy(pred, self.label)
+        self.loss = fluid.layers.reduce_mean(self.loss)
+
+        # calc acc
+        pred = fluid.layers.reshape(pred, shape=[-1, self.num_classes])
+        label = fluid.layers.reshape(self.label, shape=[-1, 1])
+        self.acc1 = fluid.layers.accuracy(pred, label, k=1)
+
+    def get_feeds(self):
+        return self.feed_vars
+
+    def get_outputs(self):
+        return {"loss": self.loss, "accuracy": self.acc1}
+
+    def get_pyreader(self):
+        return self.pyreader
+
+
+class PointNet2SemSegSSG(PointNet2SemSeg):
+    def __init__(self, num_classes, use_xyz=True):
+        super(PointNet2SemSegSSG, self).__init__(num_classes, use_xyz)
+
+    def model_config(self):
+        self.SA_confs = [
+            {
+                "npoint": 1024,
+                "radiuss": [0.1],
+                "nsamples": [32],
+                "mlps": [[32, 32, 64]],
+            },
+            {
+                "npoint": 256,
+                "radiuss": [0.2],
+                "nsamples": [32],
+                "mlps": [[64, 64, 128]],
+            },
+            {
+                "npoint": 64,
+                "radiuss": [0.4],
+                "nsamples": [32],
+                "mlps": [[128, 128, 256]],
+            },
+            {
+                "npoint": 16,
+                "radiuss": [0.8],
+                "nsamples": [32],
+                "mlps": [[256, 256, 512]],
+            },
+        ]
+
+        self.FP_confs = [
+            {"mlp": [128, 128, 128]},
+            {"mlp": [256, 128]},
+            {"mlp": [256, 256]},
+            {"mlp": [256, 256]},
+        ]
+
+
+class PointNet2SemSegMSG(PointNet2SemSeg):
+    def __init__(self, num_classes, use_xyz=True):
+        super(PointNet2SemSegMSG, self).__init__(num_classes, use_xyz)
+
+    def model_config(self):
+        self.SA_confs = [
+            {
+                "npoint": 1024,
+                "radiuss": [0.05, 0.1],
+                "nsamples": [16, 32],
+                "mlps": [[16, 16, 32], [32, 32, 64]],
+            },
+            {
+                "npoint": 256,
+                "radiuss": [0.1, 0.2],
+                "nsamples": [16, 32],
+                "mlps": [[64, 64, 128], [64, 96, 128]],
+            },
+            {
+                "npoint": 64,
+                "radiuss": [0.2, 0.4],
+                "nsamples": [16, 32],
+                "mlps": [[128, 196, 256], [128, 196, 256]],
+            },
+            {
+                "npoint": 16,
+                "radiuss": [0.4, 0.8],
+                "nsamples": [16, 32],
+                "mlps": [[256, 256, 512], [256, 384, 512]],
+            },
+        ]
+
+        self.FP_confs = [
+            {"mlp": [128, 128]},
+            {"mlp": [256, 256]},
+            {"mlp": [512, 512]},
+            {"mlp": [512, 512]},
+        ]
+

PaddleCV/Paddle3D/PointNet++/scripts/eval_cls.sh

+export CUDA_VISIBLE_DEVICES=0
+
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+python eval_cls.py --model=MSG --weights=checkpoints/200

PaddleCV/Paddle3D/PointNet++/scripts/eval_seg.sh

+export CUDA_VISIBLE_DEVICES=0
+
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+python eval_seg.py --model=MSG --weights=checkpoints/200

PaddleCV/Paddle3D/PointNet++/scripts/train_cls.sh

+export CUDA_VISIBLE_DEVICES=0
+
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+python train_cls.py --model=MSG --batch_size=16 --num_points=4096 --epoch=200

PaddleCV/Paddle3D/PointNet++/scripts/train_seg.sh

+export CUDA_VISIBLE_DEVICES=0
+
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`python -c 'import paddle; print(paddle.sysconfig.get_lib())'`
+python train_seg.py --model=MSG --batch_size=32 --num_points=4096 --epoch=201

PaddleCV/Paddle3D/PointNet++/train_cls.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+
+import os
+import sys
+import time
+import shutil
+import argparse
+import ast
+import logging
+import numpy as np
+import paddle.fluid as fluid
+import paddle.fluid.framework as framework
+
+from models import *
+from data.modelnet40_reader import ModelNet40ClsReader
+from data.data_utils import *
+from utils import *
+
+logging.root.handlers = []
+FORMAT = '%(asctime)s-%(levelname)s: %(message)s'
+logging.basicConfig(level=logging.INFO, format=FORMAT, stream=sys.stdout)
+logger = logging.getLogger(__name__)
+
+
+def parse_args():
+    parser = argparse.ArgumentParser("PointNet++ classification train script")
+    parser.add_argument(
+        '--model',
+        type=str,
+        default='MSG',
+        help='SSG or MSG model to train, default MSG')
+    parser.add_argument(
+        '--use_gpu',
+        type=ast.literal_eval,
+        default=True,
+        help='default use gpu.')
+    parser.add_argument(
+        '--batch_size',
+        type=int,
+        default=16,
+        help='training batch size, default 16')
+    parser.add_argument(
+        '--num_points',
+        type=int,
+        default=4096,
+        help='number of points in a sample, default: 4096')
+    parser.add_argument(
+        '--num_classes',
+        type=int,
+        default=40,
+        help='number of classes in dataset, default: 40')
+    parser.add_argument(
+        '--lr',
+        type=float,
+        default=0.01,
+        help='initial learning rate, default 0.01')
+    parser.add_argument(
+        '--lr_decay',
+        type=float,
+        default=0.7,
+        help='learning rate decay gamma, default 0.5')
+    parser.add_argument(
+        '--bn_momentum',
+        type=float,
+        default=0.99,
+        help='initial batch norm momentum, default 0.99')
+    parser.add_argument(
+        '--decay_steps',
+        type=int,
+        default=12500,
+        help='learning rate and batch norm momentum decay steps, default 12500')
+    parser.add_argument(
+        '--weight_decay',
+        type=float,
+        default=1e-5,
+        help='L2 regularization weight decay coeff, default 1e-5.')
+    parser.add_argument(
+        '--epoch',
+        type=int,
+        default=201,
+        help='epoch number. default 201.')
+    parser.add_argument(
+        '--data_dir',
+        type=str,
+        default='dataset/ModelNet40/modelnet40_ply_hdf5_2048',
+        help='dataset directory')
+    parser.add_argument(
+        '--save_dir',
+        type=str,
+        default='checkpoints_cls',
+        help='directory name to save train snapshoot')
+    parser.add_argument(
+        '--resume',
+        type=str,
+        default=None,
+        help='path to resume training based on previous checkpoints. '
+        'None for not resuming any checkpoints.')
+    parser.add_argument(
+        '--log_interval',
+        type=int,
+        default=1,
+        help='mini-batch interval for logging.')
+    parser.add_argument(
+        '--enable_ce',
+        action='store_true',
+        help='The flag indicating whether to run the task '
+        'for continuous evaluation.')
+    args = parser.parse_args()
+    return args
+
+
+def train():
+    args = parse_args()
+    print_arguments(args)
+    # check whether the installed paddle is compiled with GPU
+    check_gpu(args.use_gpu)
+
+    if not os.path.isdir(args.save_dir):
+        os.makedirs(args.save_dir)
+
+    assert args.model in ['MSG', 'SSG'], \
+            "--model can only be 'MSG' or 'SSG'"
+
+    # build model
+    if args.enable_ce:
+        SEED = 102
+        fluid.default_main_program().random_seed = SEED
+        framework.default_startup_program().random_seed = SEED
+
+    startup = fluid.Program()
+    train_prog = fluid.Program()
+    with fluid.program_guard(train_prog, startup):
+        with fluid.unique_name.guard():
+            train_model = PointNet2ClsMSG(args.num_classes, args.num_points) \
+                            if args.model == "MSG" else \
+                          PointNet2ClsSSG(args.num_classes, args.num_points)
+            train_model.build_model(bn_momentum=args.bn_momentum)
+            train_feeds = train_model.get_feeds()
+            train_pyreader = train_model.get_pyreader()
+            train_outputs = train_model.get_outputs()
+            train_loss = train_outputs['loss']
+            lr = fluid.layers.exponential_decay(
+                    learning_rate=args.lr,
+                    decay_steps=args.decay_steps,
+                    decay_rate=args.lr_decay,
+                    staircase=True)
+            lr = fluid.layers.clip(lr, 1e-5, args.lr)
+            optimizer = fluid.optimizer.Adam(learning_rate=lr,
+                    regularization=fluid.regularizer.L2Decay(args.weight_decay))
+            optimizer.minimize(train_loss)
+    train_keys, train_values = parse_outputs(train_outputs)
+
+    test_prog = fluid.Program()
+    with fluid.program_guard(test_prog, startup):
+        with fluid.unique_name.guard():
+            test_model = PointNet2ClsMSG(args.num_classes, args.num_points) \
+                           if args.model == "MSG" else \
+                         PointNet2ClsSSG(args.num_classes, args.num_points)
+            test_model.build_model()
+            test_feeds = test_model.get_feeds()
+            test_outputs = test_model.get_outputs()
+            test_pyreader = test_model.get_pyreader()
+    test_prog = test_prog.clone(True)
+    test_keys, test_values = parse_outputs(test_outputs)
+
+    place = fluid.CUDAPlace(0) if args.use_gpu else fluid.CPUPlace()
+    exe = fluid.Executor(place)
+    exe.run(startup)
+
+    if args.resume:
+        assert os.path.exists(args.resume), \
+                "Given resume weight dir {} not exist.".format(args.resume)
+        def if_exist(var):
+            return os.path.exists(os.path.join(args.resume, var.name))
+        fluid.io.load_vars(
+            exe, args.resume, predicate=if_exist, main_program=train_prog)
+
+    build_strategy = fluid.BuildStrategy()
+    build_strategy.memory_optimize = False
+    build_strategy.enable_inplace = False
+    build_strategy.fuse_all_optimizer_ops = False
+    train_compile_prog = fluid.compiler.CompiledProgram(
+            train_prog).with_data_parallel(loss_name=train_loss.name,
+                    build_strategy=build_strategy)
+    test_compile_prog = fluid.compiler.CompiledProgram(test_prog)
+
+    def save_model(exe, prog, path):
+        if os.path.isdir(path):
+            shutil.rmtree(path)
+        logger.info("Save model to {}".format(path))
+        fluid.io.save_persistables(exe, path, prog)
+
+    # get reader
+    trans_list = [
+        PointcloudScale(),
+        PointcloudRotate(),
+        PointcloudRotatePerturbation(),
+        PointcloudTranslate(),
+        PointcloudJitter(),
+        PointcloudRandomInputDropout(),
+    ]
+    modelnet_reader = ModelNet40ClsReader(args.data_dir, mode='train', transforms=trans_list)
+    train_reader = modelnet_reader.get_reader(args.batch_size, args.num_points)
+    train_pyreader.decorate_sample_list_generator(train_reader, place)
+    modelnet_reader = ModelNet40ClsReader(args.data_dir, mode='test', transforms=None)
+    test_reader = modelnet_reader.get_reader(args.batch_size, args.num_points)
+    test_pyreader.decorate_sample_list_generator(test_reader, place)
+
+    train_stat = Stat()
+    test_stat = Stat()
+
+    ce_time = 0
+    ce_loss = []
+
+    for epoch_id in range(args.epoch):
+        try:
+            train_pyreader.start()
+            train_iter = 0
+            train_periods = []
+            while True:
+                cur_time = time.time()
+                train_outs = exe.run(train_compile_prog, fetch_list=train_values + [lr.name])
+                period = time.time() - cur_time
+                train_periods.append(period)
+                train_stat.update(train_keys, train_outs[:-1])
+                if train_iter % args.log_interval == 0:
+                    log_str = ""
+                    for name, values in zip(train_keys + ['learning_rate'], train_outs):
+                        log_str += "{}: {:.5f}, ".format(name, np.mean(values))
+                        if name == 'loss':
+                            ce_loss.append(np.mean(values))
+                    logger.info("[TRAIN] Epoch {}, batch {}: {}time: {:.2f}".format(epoch_id, train_iter, log_str, period))
+                train_iter += 1
+        except fluid.core.EOFException:
+            logger.info("[TRAIN] Epoch {} finished, {}average time: {:.2f}".format(epoch_id, train_stat.get_mean_log(), np.mean(train_periods[1:])))
+            ce_time = np.mean(train_periods[1:])
+            save_model(exe, train_prog, os.path.join(args.save_dir, str(epoch_id)))
+
+            # evaluation
+            if not args.enable_ce:
+                try:
+                    test_pyreader.start()
+                    test_iter = 0
+                    test_periods = []
+                    while True:
+                        cur_time = time.time()
+                        test_outs = exe.run(test_compile_prog, fetch_list=test_values)
+                        period = time.time() - cur_time
+                        test_periods.append(period)
+                        test_stat.update(test_keys, test_outs)
+                        if test_iter % args.log_interval == 0:
+                            log_str = ""
+                            for name, value in zip(test_keys, test_outs):
+                                log_str += "{}: {:.4f}, ".format(name, np.mean(value))
+                            logger.info("[TEST] Epoch {}, batch {}: {}time: {:.2f}".format(epoch_id, test_iter, log_str, period))
+                        test_iter += 1
+                except fluid.core.EOFException:
+                    logger.info("[TEST] Epoch {} finished, {}average time: {:.2f}".format(epoch_id, test_stat.get_mean_log(), np.mean(test_periods[1:])))
+                finally:
+                    test_pyreader.reset()
+                    test_stat.reset()
+                    test_periods = []
+
+        finally:
+            train_pyreader.reset()
+            train_stat.reset()
+            train_periods = []
+
+    # only for ce
+    if args.enable_ce:
+        card_num = get_cards()
+        _loss = 0
+        _time = 0
+        try:
+            _time = ce_time
+            _loss = np.mean(ce_loss[1:])
+        except:
+            print("ce info error")
+        print("kpis\ttrain_cls_%s_duration_card%s\t%s" % (args.model, card_num, _time))
+        print("kpis\ttrain_cls_%s_loss_card%s\t%f" % (args.model, card_num, _loss))
+
+def get_cards():
+    num = 0
+    cards = os.environ.get('CUDA_VISIBLE_DEVICES', '')
+    if cards != '':
+        num = len(cards.split(","))
+    return num
+
+if __name__ == "__main__":
+    train()

PaddleCV/Paddle3D/PointNet++/train_seg.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+
+import os
+import sys
+import time
+import shutil
+import argparse
+import ast
+import logging
+import numpy as np
+import paddle.fluid as fluid
+import paddle.fluid.framework as framework
+
+from models import *
+from data.indoor3d_reader import Indoor3DReader
+from utils import *
+
+logging.root.handlers = []
+FORMAT = '%(asctime)s-%(levelname)s: %(message)s'
+logging.basicConfig(level=logging.INFO, format=FORMAT, stream=sys.stdout)
+logger = logging.getLogger(__name__)
+
+
+def parse_args():
+    parser = argparse.ArgumentParser("PointNet++ semantic segmentation train script")
+    parser.add_argument(
+        '--model',
+        type=str,
+        default='MSG',
+        help='SSG or MSG model to train, default MSG')
+    parser.add_argument(
+        '--use_gpu',
+        type=ast.literal_eval,
+        default=True,
+        help='default use gpu.')
+    parser.add_argument(
+        '--batch_size',
+        type=int,
+        default=32,
+        help='training batch size, default 32')
+    parser.add_argument(
+        '--num_points',
+        type=int,
+        default=4096,
+        help='number of points in a sample, default: 4096')
+    parser.add_argument(
+        '--num_classes',
+        type=int,
+        default=13,
+        help='number of classes in dataset, default: 13')
+    parser.add_argument(
+        '--lr',
+        type=float,
+        default=0.01,
+        help='initial learning rate, default 0.01')
+    parser.add_argument(
+        '--lr_decay',
+        type=float,
+        default=0.5,
+        help='learning rate decay gamma, default 0.5')
+    parser.add_argument(
+        '--bn_momentum',
+        type=float,
+        default=0.99,
+        help='initial batch norm momentum, default 0.99')
+    parser.add_argument(
+        '--decay_steps',
+        type=int,
+        default=6250,
+        help='learning rate and batch norm momentum decay steps, default 6250')
+    parser.add_argument(
+        '--weight_decay',
+        type=float,
+        default=0.,
+        help='L2 regularization weight decay coeff, default 0.')
+    parser.add_argument(
+        '--epoch',
+        type=int,
+        default=201,
+        help='epoch number. default 201.')
+    parser.add_argument(
+        '--data_dir',
+        type=str,
+        default='dataset/Indoor3DSemSeg/indoor3d_sem_seg_hdf5_data',
+        help='dataset directory')
+    parser.add_argument(
+        '--save_dir',
+        type=str,
+        default='checkpoints_seg',
+        help='directory name to save train snapshoot')
+    parser.add_argument(
+        '--resume',
+        type=str,
+        default=None,
+        help='path to resume training based on previous checkpoints. '
+        'None for not resuming any checkpoints.')
+    parser.add_argument(
+        '--log_interval',
+        type=int,
+        default=1,
+        help='mini-batch interval for logging.')
+    parser.add_argument(
+        '--enable_ce',
+        action='store_true',
+        help='The flag indicating whether to run the task '
+        'for continuous evaluation.')
+    args = parser.parse_args()
+    return args
+
+
+def train():
+    args = parse_args()
+    print_arguments(args)
+    # check whether the installed paddle is compiled with GPU
+    check_gpu(args.use_gpu)
+
+    if not os.path.isdir(args.save_dir):
+        os.makedirs(args.save_dir)
+
+    assert args.model in ['MSG', 'SSG'], \
+            "--model can only be 'MSG' or 'SSG'"
+
+    # build model
+    if args.enable_ce:
+        SEED = 102
+        fluid.default_main_program().random_seed = SEED
+        framework.default_startup_program().random_seed = SEED
+
+    startup = fluid.Program()
+    train_prog = fluid.Program()
+    with fluid.program_guard(train_prog, startup):
+        with fluid.unique_name.guard():
+            train_model = PointNet2SemSegMSG(args.num_classes, args.num_points) \
+                            if args.model == "MSG" else \
+                          PointNet2SemSegSSG(args.num_classes, args.num_points)
+            train_model.build_model(bn_momentum=args.bn_momentum)
+            train_feeds = train_model.get_feeds()
+            train_pyreader = train_model.get_pyreader()
+            train_outputs = train_model.get_outputs()
+            train_loss = train_outputs['loss']
+            lr = fluid.layers.exponential_decay(
+                    learning_rate=args.lr,
+                    decay_steps=args.decay_steps,
+                    decay_rate=args.lr_decay,
+                    staircase=True)
+            lr = fluid.layers.clip(lr, 1e-5, args.lr)
+            optimizer = fluid.optimizer.Adam(learning_rate=lr,
+                    regularization=fluid.regularizer.L2Decay(args.weight_decay))
+            optimizer.minimize(train_loss)
+    train_keys, train_values = parse_outputs(train_outputs)
+
+    test_prog = fluid.Program()
+    with fluid.program_guard(test_prog, startup):
+        with fluid.unique_name.guard():
+            test_model = PointNet2SemSegMSG(args.num_classes, args.num_points) \
+                           if args.model == "MSG" else \
+                         PointNet2SemSegSSG(args.num_classes, args.num_points)
+            test_model.build_model()
+            test_feeds = test_model.get_feeds()
+            test_outputs = test_model.get_outputs()
+            test_pyreader = test_model.get_pyreader()
+    test_prog = test_prog.clone(True)
+    test_keys, test_values = parse_outputs(test_outputs)
+
+    place = fluid.CUDAPlace(0) if args.use_gpu else fluid.CPUPlace()
+    exe = fluid.Executor(place)
+    exe.run(startup)
+
+    if args.resume:
+        assert os.path.exists(args.resume), \
+                "Given resume weight dir {} not exist.".format(args.resume)
+        def if_exist(var):
+            return os.path.exists(os.path.join(args.resume, var.name))
+        fluid.io.load_vars(
+            exe, args.resume, predicate=if_exist, main_program=train_prog)
+
+    build_strategy = fluid.BuildStrategy()
+    build_strategy.memory_optimize = False
+    build_strategy.enable_inplace = False
+    build_strategy.fuse_all_optimizer_ops = False
+    train_compile_prog = fluid.compiler.CompiledProgram(
+            train_prog).with_data_parallel(loss_name=train_loss.name,
+                    build_strategy=build_strategy)
+    test_compile_prog = fluid.compiler.CompiledProgram(test_prog)
+
+    def save_model(exe, prog, path):
+        if os.path.isdir(path):
+            shutil.rmtree(path)
+        logger.info("Save model to {}".format(path))
+        fluid.io.save_persistables(exe, path, prog)
+
+    # get reader
+    indoor_reader = Indoor3DReader(args.data_dir)
+    train_reader = indoor_reader.get_reader(args.batch_size, args.num_points, mode='train')
+    test_reader = indoor_reader.get_reader(args.batch_size, args.num_points, mode='test')
+    train_pyreader.decorate_sample_list_generator(train_reader, place)
+    test_pyreader.decorate_sample_list_generator(test_reader, place)
+
+    train_stat = Stat()
+    test_stat = Stat()
+
+    ce_time = 0
+    ce_loss = []
+
+    for epoch_id in range(args.epoch):
+        try:
+            train_pyreader.start()
+            train_iter = 0
+            train_periods = []
+            while True:
+                cur_time = time.time()
+                train_outs = exe.run(train_compile_prog, fetch_list=train_values + [lr.name])
+                period = time.time() - cur_time
+                train_periods.append(period)
+                train_stat.update(train_keys, train_outs[:-1])
+                if train_iter % args.log_interval == 0:
+                    log_str = ""
+                    for name, values in zip(train_keys + ['learning_rate'], train_outs):
+                        log_str += "{}: {:.5f}, ".format(name, np.mean(values))
+                        if name == 'loss':
+                            ce_loss.append(np.mean(values))
+                    logger.info("[TRAIN] Epoch {}, batch {}: {}time: {:.2f}".format(epoch_id, train_iter, log_str, period))
+                train_iter += 1
+        except fluid.core.EOFException:
+            logger.info("[TRAIN] Epoch {} finished, {}average time: {:.2f}".format(epoch_id, train_stat.get_mean_log(), np.mean(train_periods[1:])))
+            ce_time = np.mean(train_periods[1:])
+            save_model(exe, train_prog, os.path.join(args.save_dir, str(epoch_id)))
+            
+            # evaluation
+            if not args.enable_ce:
+                try:
+                    test_pyreader.start()
+                    test_iter = 0
+                    test_periods = []
+                    while True:
+                        cur_time = time.time()
+                        test_outs = exe.run(test_compile_prog, fetch_list=test_values)
+                        period = time.time() - cur_time
+                        test_periods.append(period)
+                        test_stat.update(test_keys, test_outs)
+                        if test_iter % args.log_interval == 0:
+                            log_str = ""
+                            for name, value in zip(test_keys, test_outs):
+                                log_str += "{}: {:.4f}, ".format(name, np.mean(value))
+                            logger.info("[TEST] Epoch {}, batch {}: {}time: {:.2f}".format(epoch_id, test_iter, log_str, period))
+                        test_iter += 1
+                except fluid.core.EOFException:
+                    logger.info("[TEST] Epoch {} finished, {}average time: {:.2f}".format(epoch_id, test_stat.get_mean_log(), np.mean(test_periods[1:])))
+                finally:
+                    test_pyreader.reset()
+                    test_stat.reset()
+                    test_periods = []
+
+        finally:
+            train_pyreader.reset()
+            train_stat.reset()
+            train_periods = []
+
+    # only for ce
+    if args.enable_ce:
+        card_num = get_cards()
+        _loss = 0
+        _time = 0
+        try:
+            _time = ce_time
+            _loss = np.mean(ce_loss[1:])
+        except:
+            print("ce info error")
+        print("kpis\ttrain_seg_%s_duration_card%s\t%s" % (args.model, card_num, _time))
+        print("kpis\ttrain_seg_%s_loss_card%s\t%f" % (args.model, card_num, _loss))
+
+def get_cards():
+    num = 0
+    cards = os.environ.get('CUDA_VISIBLE_DEVICES', '')
+    if cards != '':
+        num = len(cards.split(","))
+    return num
+
+if __name__ == "__main__":
+    train()

PaddleCV/Paddle3D/PointNet++/utils.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.
+"""
+Contains common utility functions.
+"""
+
+from __future__ import absolute_import
+from __future__ import division
+from __future__ import print_function
+
+import sys
+import six
+import logging
+import numpy as np
+import paddle.fluid as fluid
+
+__all__ = ["check_gpu", "print_arguments", "parse_outputs", "Stat"]
+
+logger = logging.getLogger(__name__)
+
+
+def check_gpu(use_gpu):
+    """
+    Log error and exit when set use_gpu=True in paddlepaddle
+    cpu version.
+    """
+    err = "Config use_gpu cannot be set as True while you are " \
+          "using paddlepaddle cpu version ! \nPlease try: \n" \
+          "\t1. Install paddlepaddle-gpu to run model on GPU \n" \
+          "\t2. Set --use_gpu=False to run model on CPU"
+
+    try:
+        if use_gpu and not fluid.is_compiled_with_cuda():
+            logger.error(err)
+            sys.exit(1)
+    except Exception as e:
+        pass
+
+
+def print_arguments(args):
+    """Print argparse's arguments.
+
+    Usage:
+
+    .. code-block:: python
+
+        parser = argparse.ArgumentParser()
+        parser.add_argument("name", default="Jonh", type=str, help="User name.")
+        args = parser.parse_args()
+        print_arguments(args)
+
+    :param args: Input argparse.Namespace for printing.
+    :type args: argparse.Namespace
+    """
+    logger.info("-----------  Configuration Arguments -----------")
+    for arg, value in sorted(six.iteritems(vars(args))):
+        logger.info("%s: %s" % (arg, value))
+    logger.info("------------------------------------------------")
+
+
+def parse_outputs(outputs):
+    keys, values = [], []
+    for k, v in outputs.items():
+        keys.append(k)
+        v.persistable = True
+        values.append(v.name)
+    return keys, values
+
+
+class Stat(object):
+    def __init__(self):
+        self.stats = {}
+
+    def update(self, keys, values):
+        for k, v in zip(keys, values):
+            if k not in self.stats:
+                self.stats[k] = []
+            self.stats[k].append(v)
+
+    def reset(self):
+        self.stats = {}
+
+    def get_mean_log(self):
+        log = ""
+        for k, v in self.stats.items():
+            log += "avg_{}: {:.4f}, ".format(k, np.mean(v))
+        return log

PaddleCV/Paddle3D/PointRCNN/.gitignore

+*log*
+checkpoints*
+build
+output
+result_dir
+pp_pointrcnn*
+data/gt_database
+utils/pts_utils/dist
+utils/pts_utils/build
+utils/pts_utils/pts_utils.egg-info
+utils/cyops/*.c
+utils/cyops/*.so
+ext_op/src/*.o
+ext_op/src/*.so

PaddleCV/Paddle3D/PointRCNN/build_and_install.sh

+# compile cyops
+python utils/cyops/setup.py develop
+
+# compile and install pts_utils
+cd utils/pts_utils
+python setup.py install
+cd ../..

PaddleCV/Paddle3D/PointRCNN/cfgs/default.yml

+# This config is based on https://github.com/sshaoshuai/PointRCNN/blob/master/tools/cfgs/default.yaml
+CLASSES: Car
+
+INCLUDE_SIMILAR_TYPE: True
+
+# config of augmentation
+AUG_DATA: True
+AUG_METHOD_LIST: ['rotation', 'scaling', 'flip']
+AUG_METHOD_PROB: [1.0, 1.0, 0.5]
+AUG_ROT_RANGE: 18
+
+GT_AUG_ENABLED: True
+GT_EXTRA_NUM: 15
+GT_AUG_RAND_NUM: True
+GT_AUG_APPLY_PROB: 1.0
+GT_AUG_HARD_RATIO: 0.6
+
+PC_REDUCE_BY_RANGE: True
+PC_AREA_SCOPE: [[-40, 40], [-1,   3], [0, 70.4]]  # x, y, z scope in rect camera coords
+CLS_MEAN_SIZE: [[1.52563191462, 1.62856739989, 3.88311640418]]
+
+
+# 1. config of rpn network
+RPN:
+    ENABLED: True
+    FIXED: False
+
+    # config of input
+    USE_INTENSITY: False
+
+    # config of bin-based loss
+    LOC_XZ_FINE: True
+    LOC_SCOPE: 3.0
+    LOC_BIN_SIZE: 0.5
+    NUM_HEAD_BIN: 12
+
+    # config of network structure
+    BACKBONE: pointnet2_msg
+    USE_BN: True
+    NUM_POINTS: 16384
+
+    SA_CONFIG:
+        NPOINTS: [4096, 1024, 256, 64]
+        RADIUS: [[0.1, 0.5], [0.5, 1.0], [1.0, 2.0], [2.0, 4.0]]
+        NSAMPLE: [[16, 32], [16, 32], [16, 32], [16, 32]]
+        MLPS: [[[16, 16, 32], [32, 32, 64]],
+              [[64, 64, 128], [64, 96, 128]],
+              [[128, 196, 256], [128, 196, 256]],
+              [[256, 256, 512], [256, 384, 512]]]
+    FP_MLPS: [[128, 128], [256, 256], [512, 512], [512, 512]]
+    CLS_FC: [128]
+    REG_FC: [128]
+    DP_RATIO: 0.5
+
+    # config of training
+    LOSS_CLS: SigmoidFocalLoss
+    FG_WEIGHT: 15
+    FOCAL_ALPHA: [0.25, 0.75]
+    FOCAL_GAMMA: 2.0
+    REG_LOSS_WEIGHT: [1.0, 1.0, 1.0, 1.0]
+    LOSS_WEIGHT: [1.0, 1.0]
+    NMS_TYPE: normal
+
+    # config of testing
+    SCORE_THRESH: 0.3
+
+# 2. config of rcnn network
+RCNN:
+    ENABLED: True
+
+    # config of input
+    ROI_SAMPLE_JIT: False 
+    REG_AUG_METHOD: multiple  # multiple, single, normal
+    ROI_FG_AUG_TIMES: 10
+
+    USE_RPN_FEATURES: True
+    USE_MASK: True
+    MASK_TYPE: seg
+    USE_INTENSITY: False
+    USE_DEPTH: True
+    USE_SEG_SCORE: False
+
+    POOL_EXTRA_WIDTH: 1.0
+
+    # config of bin-based loss
+    LOC_SCOPE: 1.5
+    LOC_BIN_SIZE: 0.5
+    NUM_HEAD_BIN: 9
+    LOC_Y_BY_BIN: False
+    LOC_Y_SCOPE: 0.5
+    LOC_Y_BIN_SIZE: 0.25
+    SIZE_RES_ON_ROI: False
+
+    # config of network structure
+    USE_BN: False
+    DP_RATIO: 0.0
+
+    BACKBONE: pointnet  # pointnet
+    XYZ_UP_LAYER: [128, 128]
+
+    NUM_POINTS: 512
+    SA_CONFIG:
+        NPOINTS: [128, 32, -1]
+        RADIUS: [0.2, 0.4, 100]
+        NSAMPLE: [64, 64, 64]
+        MLPS: [[128, 128, 128],
+               [128, 128, 256],
+               [256, 256, 512]]
+    CLS_FC: [256, 256]
+    REG_FC: [256, 256]
+
+    # config of training
+    LOSS_CLS: BinaryCrossEntropy
+    FOCAL_ALPHA: [0.25, 0.75]
+    FOCAL_GAMMA: 2.0
+    CLS_WEIGHT: [1.0, 1.0, 1.0]
+    CLS_FG_THRESH: 0.6
+    CLS_BG_THRESH: 0.45
+    CLS_BG_THRESH_LO: 0.05
+    REG_FG_THRESH: 0.55
+    FG_RATIO: 0.5
+    ROI_PER_IMAGE: 64
+    HARD_BG_RATIO: 0.8
+
+    # config of testing
+    SCORE_THRESH: 0.3
+    NMS_THRESH: 0.1
+
+# general training config
+TRAIN:
+    SPLIT: train
+    VAL_SPLIT: smallval
+
+    LR: 0.002
+    LR_CLIP: 0.00001
+    LR_DECAY: 0.5
+    DECAY_STEP_LIST: [100, 150, 180, 200]
+    LR_WARMUP: True
+    WARMUP_MIN: 0.0002
+    WARMUP_EPOCH: 1
+
+    BN_MOMENTUM: 0.1
+    BN_DECAY: 0.5
+    BNM_CLIP: 0.01
+    BN_DECAY_STEP_LIST: [1000]
+
+    OPTIMIZER: adam # adam, adam_onecycle
+    WEIGHT_DECAY: 0.001  # L2 regularization
+    MOMENTUM: 0.9
+
+    MOMS: [0.95, 0.85]
+    DIV_FACTOR: 10.0
+    PCT_START: 0.4
+
+    GRAD_NORM_CLIP: 1.0
+
+    RPN_PRE_NMS_TOP_N: 9000
+    RPN_POST_NMS_TOP_N: 512
+    RPN_NMS_THRESH: 0.85
+    RPN_DISTANCE_BASED_PROPOSE: True
+
+TEST:
+    SPLIT: val
+    RPN_PRE_NMS_TOP_N: 9000
+    RPN_POST_NMS_TOP_N: 100
+    RPN_NMS_THRESH: 0.8
+    RPN_DISTANCE_BASED_PROPOSE: True

PaddleCV/Paddle3D/PointRCNN/data/KITTI/object/download.sh

+DIR="$( cd "$(dirname "$0")" ; pwd -P  )"
+cd "$DIR"
+
+echo "Downloading https://s3.eu-central-1.amazonaws.com/avg-kitti/data_object_velodyne.zip"
+wget https://s3.eu-central-1.amazonaws.com/avg-kitti/data_object_velodyne.zip
+echo "https://s3.eu-central-1.amazonaws.com/avg-kitti/data_object_image_2.zip"
+wget https://s3.eu-central-1.amazonaws.com/avg-kitti/data_object_image_2.zip
+echo "https://s3.eu-central-1.amazonaws.com/avg-kitti/data_object_calib.zip"
+wget https://s3.eu-central-1.amazonaws.com/avg-kitti/data_object_calib.zip
+echo "https://s3.eu-central-1.amazonaws.com/avg-kitti/data_object_label_2.zip"
+wget https://s3.eu-central-1.amazonaws.com/avg-kitti/data_object_label_2.zip
+
+echo "Decompressing data_object_velodyne.zip"
+unzip data_object_velodyne.zip
+echo "Decompressing data_object_image_2.zip"
+unzip "data_object_image_2.zip"
+echo "Decompressing data_object_calib.zip"
+unzip data_object_calib.zip
+echo "Decompressing data_object_label_2.zip"
+unzip data_object_label_2.zip
+
+echo "Download KITTI ImageSets"
+wget https://paddlemodels.bj.bcebos.com/Paddle3D/pointrcnn_kitti_imagesets.tar
+tar xf pointrcnn_kitti_imagesets.tar
+mv ImageSets ..

PaddleCV/Paddle3D/PointRCNN/data/__init__.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.

PaddleCV/Paddle3D/PointRCNN/ext_op

+../PointNet++/ext_op
\ No newline at end of file

PaddleCV/Paddle3D/PointRCNN/models/__init__.py

+#  Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserve.
+#
+#Licensed under the Apache License, Version 2.0 (the "License");
+#you may not use this file except in compliance with the License.
+#You may obtain a copy of the License at
+#
+#    http://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.

PaddleCV/Paddle3D/PointRCNN/requirement.txt

+Cython
+opencv-python
+shapely
+scikit-image
+Numba
+fire

PaddleCV/Paddle3D/PointRCNN/utils/pts_utils/CMakeLists.txt

+
+cmake_minimum_required(VERSION 2.8.12)
+project(pts_utils)
+
+add_subdirectory(pybind11)
+pybind11_add_module(pts_utils pts_utils.cpp)