1 Commits

Author SHA1 Message Date
bottler
9c586b1351 Run tests in github action not circleci (#1896)
Summary: Pull Request resolved: https://github.com/facebookresearch/pytorch3d/pull/1896

Differential Revision: D65272512

Pulled By: bottler
2024-10-31 08:41:20 -07:00
307 changed files with 2460 additions and 2850 deletions

View File

@@ -88,6 +88,7 @@ def workflow_pair(
upload=False,
filter_branch,
):
w = []
py = python_version.replace(".", "")
pyt = pytorch_version.replace(".", "")
@@ -126,6 +127,7 @@ def generate_base_workflow(
btype,
filter_branch=None,
):
d = {
"name": base_workflow_name,
"python_version": python_version,

View File

@@ -3,9 +3,6 @@ on:
pull_request:
branches:
- main
push:
branches:
- main
jobs:
binary_linux_conda_cuda:
runs-on: 4-core-ubuntu-gpu-t4
@@ -21,45 +18,3 @@ jobs:
run: |-
conda create --name env --yes --quiet conda-build
conda run --no-capture-output --name env python3 ./packaging/build_conda.py --use-conda-cuda
# Build-only verification for the ROCm/HIP code paths. Runs in an AMD ROCm
# dev container on a CPU-only GitHub runner; we don't need an AMD GPU just
# to compile, and not running tests keeps the CI cost low. Catches build
# regressions in the ROCm code paths (USE_ROCM guards, hipify-touched sources,
# the pulsar HIP intrinsic replacements, etc.).
linux_rocm_build:
runs-on: ubuntu-latest
container:
# `-complete` tag bundles the full ROCm math stack (rocThrust, hipCUB,
# rocPRIM, ...). The plain `7.2.3` tag is HIP-runtime-only and fails to
# find <thrust/complex.h> when including PyTorch headers.
image: rocm/dev-ubuntu-22.04:7.2.3-complete
env:
PYTORCH_VERSION: "2.11.0"
ROCM_INDEX: "rocm7.2"
steps:
- uses: actions/checkout@v4
- name: Install Python and torch+rocm
run: |-
apt-get update
apt-get install -y --no-install-recommends python3 python3-dev python3-pip git
python3 -m pip install --upgrade pip
python3 -m pip install --index-url https://download.pytorch.org/whl/${ROCM_INDEX} torch==${PYTORCH_VERSION}
- name: Verify torch is ROCm-built
run: |-
python3 -c "import torch; assert torch.version.hip is not None, 'torch is not HIP-built'; print('torch.version.hip:', torch.version.hip)"
- name: Build pytorch3d _C extension (build only, no tests)
env:
# CPU-only runner: torch.cuda.is_available() is False, so force the
# CUDAExtension path. ROCM_HOME is auto-detected from /opt/rocm in
# the rocm/dev-ubuntu container.
FORCE_CUDA: "1"
run: |-
python3 -m pip install --no-build-isolation -v .
- name: Smoke import
# cd out of the checkout root so the source-tree pytorch3d/ directory
# (which has no _C.so since the build doesn't install in-place) doesn't
# shadow the site-packages install via sys.path[0] for `python -c`.
run: |-
cd /tmp
python3 -c "import torch; from pytorch3d import _C; print('PulsarRenderer:', hasattr(_C, 'PulsarRenderer')); print('n_symbols:', len([s for s in dir(_C) if not s.startswith('_')]))"

11
.gitignore vendored
View File

@@ -6,17 +6,6 @@ dist/
**/.ipynb_checkpoints
**/.ipynb_checkpoints/**
# Build artifacts produced in-place by torch.utils.cpp_extension auto-hipify
# when pytorch3d is built against a ROCm PyTorch.
*.so
pytorch3d/csrc/**/*.hip
pytorch3d/csrc/**/*_hip.cpp
pytorch3d/csrc/**/*_hip.h
pytorch3d/csrc/**/*_hip.cuh
# Debug PNG dumps written by pulsar tests when FB_TEST is unset.
tests/pulsar/test_out/
# Docusaurus site
website/yarn.lock

View File

@@ -10,7 +10,7 @@
DIR="$( cd "$( dirname "${BASH_SOURCE[0]}" )" >/dev/null 2>&1 && pwd )"
DIR=$(dirname "${DIR}")
if [[ -f "${DIR}/BUCK" ]]
if [[ -f "${DIR}/TARGETS" ]]
then
pyfmt "${DIR}"
else
@@ -36,5 +36,5 @@ then
echo "Running pyre..."
echo "To restart/kill pyre server, run 'pyre restart' or 'pyre kill' in fbcode/"
( cd ~/fbsource/fbcode; arc pyre check //vision/fair/pytorch3d/... )
( cd ~/fbsource/fbcode; pyre -l vision/fair/pytorch3d/ )
fi

View File

@@ -19,6 +19,7 @@
#
import os
import sys
import unittest.mock as mock
from recommonmark.parser import CommonMarkParser

View File

@@ -10,7 +10,6 @@ This example demonstrates the most trivial, direct interface of the pulsar
sphere renderer. It renders and saves an image with 10 random spheres.
Output: basic.png.
"""
import logging
import math
from os import path

View File

@@ -11,7 +11,6 @@ interface for sphere renderering. It renders and saves an image with
10 random spheres.
Output: basic-pt3d.png.
"""
import logging
from os import path

View File

@@ -14,7 +14,6 @@ distorted. Gradient-based optimization is used to converge towards the
original camera parameters.
Output: cam.gif.
"""
import logging
import math
from os import path

View File

@@ -14,7 +14,6 @@ distorted. Gradient-based optimization is used to converge towards the
original camera parameters.
Output: cam-pt3d.gif
"""
import logging
from os import path

View File

@@ -18,7 +18,6 @@ This example is not available yet through the 'unified' interface,
because opacity support has not landed in PyTorch3D for general data
structures yet.
"""
import logging
import math
from os import path

View File

@@ -13,7 +13,6 @@ The scene is initialized with random spheres. Gradient-based
optimization is used to converge towards a faithful
scene representation.
"""
import logging
import math

View File

@@ -13,7 +13,6 @@ The scene is initialized with random spheres. Gradient-based
optimization is used to converge towards a faithful
scene representation.
"""
import logging
import math

View File

@@ -3,6 +3,11 @@ pytorch3d.implicitron.dataset specific datasets
specific datasets
.. automodule:: pytorch3d.implicitron.dataset.blender_dataset_map_provider
:members:
:undoc-members:
:show-inheritance:
.. automodule:: pytorch3d.implicitron.dataset.json_index_dataset_map_provider
:members:
:undoc-members:
@@ -13,6 +18,11 @@ specific datasets
:undoc-members:
:show-inheritance:
.. automodule:: pytorch3d.implicitron.dataset.llff_dataset_map_provider
:members:
:undoc-members:
:show-inheritance:
.. automodule:: pytorch3d.implicitron.dataset.rendered_mesh_dataset_map_provider
:members:
:undoc-members:

View File

@@ -32,6 +32,7 @@ requirements:
build:
string: py{{py}}_{{ environ['CU_VERSION'] }}_pyt{{ environ['PYTORCH_VERSION_NODOT']}}
# script: LD_LIBRARY_PATH=$PREFIX/lib:$BUILD_PREFIX/lib:$LD_LIBRARY_PATH python setup.py install --single-version-externally-managed --record=record.txt # [not win]
script: python setup.py install --single-version-externally-managed --record=record.txt # [not win]
script_env:
- CUDA_HOME
@@ -56,6 +57,7 @@ test:
- pandas
- sqlalchemy
commands:
#pytest .
python -m unittest discover -v -s tests -t .

View File

@@ -0,0 +1,56 @@
defaults:
- overfit_singleseq_base
- _self_
exp_dir: "./data/overfit_nerf_blender_repro/${oc.env:BLENDER_SINGLESEQ_CLASS}"
data_source_ImplicitronDataSource_args:
data_loader_map_provider_SequenceDataLoaderMapProvider_args:
dataset_length_train: 100
dataset_map_provider_class_type: BlenderDatasetMapProvider
dataset_map_provider_BlenderDatasetMapProvider_args:
base_dir: ${oc.env:BLENDER_DATASET_ROOT}/${oc.env:BLENDER_SINGLESEQ_CLASS}
n_known_frames_for_test: null
object_name: ${oc.env:BLENDER_SINGLESEQ_CLASS}
path_manager_factory_class_type: PathManagerFactory
path_manager_factory_PathManagerFactory_args:
silence_logs: true
model_factory_ImplicitronModelFactory_args:
model_class_type: "OverfitModel"
model_OverfitModel_args:
mask_images: false
raysampler_class_type: AdaptiveRaySampler
raysampler_AdaptiveRaySampler_args:
n_pts_per_ray_training: 64
n_pts_per_ray_evaluation: 64
n_rays_per_image_sampled_from_mask: 4096
stratified_point_sampling_training: true
stratified_point_sampling_evaluation: false
scene_extent: 2.0
scene_center:
- 0.0
- 0.0
- 0.0
renderer_MultiPassEmissionAbsorptionRenderer_args:
density_noise_std_train: 0.0
n_pts_per_ray_fine_training: 128
n_pts_per_ray_fine_evaluation: 128
raymarcher_EmissionAbsorptionRaymarcher_args:
blend_output: false
loss_weights:
loss_rgb_mse: 1.0
loss_prev_stage_rgb_mse: 1.0
loss_mask_bce: 0.0
loss_prev_stage_mask_bce: 0.0
loss_autodecoder_norm: 0.00
optimizer_factory_ImplicitronOptimizerFactory_args:
exponential_lr_step_size: 3001
lr_policy: LinearExponential
linear_exponential_lr_milestone: 200
training_loop_ImplicitronTrainingLoop_args:
max_epochs: 6000
metric_print_interval: 10
store_checkpoints_purge: 3
test_when_finished: true
validation_interval: 100

View File

@@ -0,0 +1,55 @@
defaults:
- repro_singleseq_base
- _self_
exp_dir: "./data/nerf_blender_repro/${oc.env:BLENDER_SINGLESEQ_CLASS}"
data_source_ImplicitronDataSource_args:
data_loader_map_provider_SequenceDataLoaderMapProvider_args:
dataset_length_train: 100
dataset_map_provider_class_type: BlenderDatasetMapProvider
dataset_map_provider_BlenderDatasetMapProvider_args:
base_dir: ${oc.env:BLENDER_DATASET_ROOT}/${oc.env:BLENDER_SINGLESEQ_CLASS}
n_known_frames_for_test: null
object_name: ${oc.env:BLENDER_SINGLESEQ_CLASS}
path_manager_factory_class_type: PathManagerFactory
path_manager_factory_PathManagerFactory_args:
silence_logs: true
model_factory_ImplicitronModelFactory_args:
model_GenericModel_args:
mask_images: false
raysampler_class_type: AdaptiveRaySampler
raysampler_AdaptiveRaySampler_args:
n_pts_per_ray_training: 64
n_pts_per_ray_evaluation: 64
n_rays_per_image_sampled_from_mask: 4096
stratified_point_sampling_training: true
stratified_point_sampling_evaluation: false
scene_extent: 2.0
scene_center:
- 0.0
- 0.0
- 0.0
renderer_MultiPassEmissionAbsorptionRenderer_args:
density_noise_std_train: 0.0
n_pts_per_ray_fine_training: 128
n_pts_per_ray_fine_evaluation: 128
raymarcher_EmissionAbsorptionRaymarcher_args:
blend_output: false
loss_weights:
loss_rgb_mse: 1.0
loss_prev_stage_rgb_mse: 1.0
loss_mask_bce: 0.0
loss_prev_stage_mask_bce: 0.0
loss_autodecoder_norm: 0.00
optimizer_factory_ImplicitronOptimizerFactory_args:
exponential_lr_step_size: 3001
lr_policy: LinearExponential
linear_exponential_lr_milestone: 200
training_loop_ImplicitronTrainingLoop_args:
max_epochs: 6000
metric_print_interval: 10
store_checkpoints_purge: 3
test_when_finished: true
validation_interval: 100

View File

@@ -7,7 +7,7 @@
# pyre-unsafe
""" "
""""
This file is the entry point for launching experiments with Implicitron.
Launch Training
@@ -44,22 +44,25 @@ The outputs of the experiment are saved and logged in multiple ways:
config file.
"""
import logging
import os
import warnings
from dataclasses import field
import hydra
import torch
from accelerate import Accelerator
from omegaconf import DictConfig, OmegaConf
from packaging import version
from pytorch3d.implicitron.dataset.data_source import (
DataSourceBase,
ImplicitronDataSource,
)
from pytorch3d.implicitron.models.base_model import ImplicitronModelBase
from pytorch3d.implicitron.models.renderer.multipass_ea import (
MultiPassEmissionAbsorptionRenderer,
)
@@ -117,12 +120,16 @@ class Experiment(Configurable):
will be saved here.
"""
# pyre-fixme[13]: Attribute `data_source` is never initialized.
data_source: DataSourceBase
data_source_class_type: str = "ImplicitronDataSource"
# pyre-fixme[13]: Attribute `model_factory` is never initialized.
model_factory: ModelFactoryBase
model_factory_class_type: str = "ImplicitronModelFactory"
# pyre-fixme[13]: Attribute `optimizer_factory` is never initialized.
optimizer_factory: OptimizerFactoryBase
optimizer_factory_class_type: str = "ImplicitronOptimizerFactory"
# pyre-fixme[13]: Attribute `training_loop` is never initialized.
training_loop: TrainingLoopBase
training_loop_class_type: str = "ImplicitronTrainingLoop"
@@ -189,7 +196,6 @@ class Experiment(Configurable):
last_epoch=start_epoch,
model=model,
resume=self.model_factory.resume,
# pyrefly: ignore [missing-attribute]
resume_epoch=self.model_factory.resume_epoch,
)
@@ -209,7 +215,6 @@ class Experiment(Configurable):
# Enter the main training loop.
self.training_loop.run(
# pyrefly: ignore [bad-argument-type]
train_loader=train_loader,
val_loader=val_loader,
test_loader=test_loader,

View File

@@ -11,6 +11,7 @@ import os
from typing import Optional
import torch.optim
from accelerate import Accelerator
from pytorch3d.implicitron.models.base_model import ImplicitronModelBase
from pytorch3d.implicitron.tools import model_io
@@ -25,6 +26,7 @@ logger = logging.getLogger(__name__)
class ModelFactoryBase(ReplaceableBase):
resume: bool = True # resume from the last checkpoint
def __call__(self, **kwargs) -> ImplicitronModelBase:
@@ -59,6 +61,7 @@ class ImplicitronModelFactory(ModelFactoryBase):
"""
# pyre-fixme[13]: Attribute `model` is never initialized.
model: ImplicitronModelBase
model_class_type: str = "GenericModel"
resume: bool = True
@@ -113,9 +116,7 @@ class ImplicitronModelFactory(ModelFactoryBase):
"cuda:%d" % 0: "cuda:%d" % accelerator.local_process_index
}
model_state_dict = torch.load(
model_io.get_model_path(model_path),
map_location=map_location,
weights_only=True,
model_io.get_model_path(model_path), map_location=map_location
)
try:

View File

@@ -14,7 +14,9 @@ from dataclasses import field
from typing import Any, Dict, List, Optional, Tuple
import torch.optim
from accelerate import Accelerator
from pytorch3d.implicitron.models.base_model import ImplicitronModelBase
from pytorch3d.implicitron.tools import model_io
from pytorch3d.implicitron.tools.config import (
@@ -121,7 +123,6 @@ class ImplicitronOptimizerFactory(OptimizerFactoryBase):
"""
# Get the parameters to optimize
if hasattr(model, "_get_param_groups"): # use the model function
# pyre-fixme[29]: `Union[Tensor, Module]` is not a function.
p_groups = model._get_param_groups(self.lr, wd=self.weight_decay)
else:
p_groups = [
@@ -172,7 +173,6 @@ class ImplicitronOptimizerFactory(OptimizerFactoryBase):
scheduler = torch.optim.lr_scheduler.LambdaLR(
optimizer,
lambda epoch: self.gamma ** (epoch / self.exponential_lr_step_size),
# pyrefly: ignore [unexpected-keyword]
verbose=False,
)
elif self.lr_policy.casefold() == "LinearExponential".casefold():
@@ -190,10 +190,7 @@ class ImplicitronOptimizerFactory(OptimizerFactoryBase):
return gamma
scheduler = torch.optim.lr_scheduler.LambdaLR(
optimizer,
_get_lr,
# pyrefly: ignore [unexpected-keyword]
verbose=False,
optimizer, _get_lr, verbose=False
)
else:
raise ValueError("no such lr policy %s" % self.lr_policy)
@@ -244,7 +241,7 @@ class ImplicitronOptimizerFactory(OptimizerFactoryBase):
map_location = {
"cuda:%d" % 0: "cuda:%d" % accelerator.local_process_index
}
optimizer_state = torch.load(opt_path, map_location, weights_only=True)
optimizer_state = torch.load(opt_path, map_location)
else:
raise FileNotFoundError(f"Optimizer state {opt_path} does not exist.")
return optimizer_state

View File

@@ -36,6 +36,7 @@ class TrainingLoopBase(ReplaceableBase):
evaluator: An EvaluatorBase instance, used to evaluate training results.
"""
# pyre-fixme[13]: Attribute `evaluator` is never initialized.
evaluator: Optional[EvaluatorBase]
evaluator_class_type: Optional[str] = "ImplicitronEvaluator"
@@ -160,6 +161,7 @@ class ImplicitronTrainingLoop(TrainingLoopBase):
for epoch in range(start_epoch, self.max_epochs):
# automatic new_epoch and plotting of stats at every epoch start
with stats:
# Make sure to re-seed random generators to ensure reproducibility
# even after restart.
seed_all_random_engines(seed + epoch)
@@ -198,7 +200,6 @@ class ImplicitronTrainingLoop(TrainingLoopBase):
and self.test_interval > 0
and epoch % self.test_interval == 0
):
# pyrefly: ignore [missing-attribute]
self.evaluator.run(
device=device,
dataloader=test_loader,
@@ -215,7 +216,6 @@ class ImplicitronTrainingLoop(TrainingLoopBase):
if self.test_when_finished:
if test_loader is not None:
# pyrefly: ignore [missing-attribute]
self.evaluator.run(
device=device,
dump_to_json=True,
@@ -379,6 +379,7 @@ class ImplicitronTrainingLoop(TrainingLoopBase):
# update the stats logger
stats.update(preds, time_start=t_start, stat_set=trainmode)
# pyre-ignore [16]
assert stats.it[trainmode] == it, "inconsistent stat iteration number!"
# print textual status update
@@ -394,7 +395,6 @@ class ImplicitronTrainingLoop(TrainingLoopBase):
):
prefix = f"e{stats.epoch}_it{stats.it[trainmode]}"
if hasattr(model, "visualize"):
# pyre-fixme[29]: `Union[Tensor, Module]` is not a function.
model.visualize(
viz,
visdom_env_imgs,

View File

@@ -13,6 +13,13 @@ hydra:
data_source_ImplicitronDataSource_args:
dataset_map_provider_class_type: ???
data_loader_map_provider_class_type: SequenceDataLoaderMapProvider
dataset_map_provider_BlenderDatasetMapProvider_args:
base_dir: ???
object_name: ???
path_manager_factory_class_type: PathManagerFactory
n_known_frames_for_test: null
path_manager_factory_PathManagerFactory_args:
silence_logs: true
dataset_map_provider_JsonIndexDatasetMapProvider_args:
category: ???
task_str: singlesequence
@@ -84,6 +91,14 @@ data_source_ImplicitronDataSource_args:
sort_frames: false
path_manager_factory_PathManagerFactory_args:
silence_logs: true
dataset_map_provider_LlffDatasetMapProvider_args:
base_dir: ???
object_name: ???
path_manager_factory_class_type: PathManagerFactory
n_known_frames_for_test: null
path_manager_factory_PathManagerFactory_args:
silence_logs: true
downscale_factor: 4
dataset_map_provider_RenderedMeshDatasetMapProvider_args:
num_views: 40
data_file: null

View File

@@ -12,6 +12,7 @@ import unittest
from pathlib import Path
import torch
from hydra import compose, initialize_config_dir
from omegaconf import OmegaConf
from projects.implicitron_trainer.impl.optimizer_factory import (
@@ -52,8 +53,12 @@ class TestExperiment(unittest.TestCase):
cfg.data_source_ImplicitronDataSource_args.dataset_map_provider_class_type = (
"JsonIndexDatasetMapProvider"
)
dataset_args = cfg.data_source_ImplicitronDataSource_args.dataset_map_provider_JsonIndexDatasetMapProvider_args
dataloader_args = cfg.data_source_ImplicitronDataSource_args.data_loader_map_provider_SequenceDataLoaderMapProvider_args
dataset_args = (
cfg.data_source_ImplicitronDataSource_args.dataset_map_provider_JsonIndexDatasetMapProvider_args
)
dataloader_args = (
cfg.data_source_ImplicitronDataSource_args.data_loader_map_provider_SequenceDataLoaderMapProvider_args
)
dataset_args.category = "skateboard"
dataset_args.test_restrict_sequence_id = 0
dataset_args.dataset_root = "manifold://co3d/tree/extracted"
@@ -89,8 +94,12 @@ class TestExperiment(unittest.TestCase):
cfg.data_source_ImplicitronDataSource_args.dataset_map_provider_class_type = (
"JsonIndexDatasetMapProvider"
)
dataset_args = cfg.data_source_ImplicitronDataSource_args.dataset_map_provider_JsonIndexDatasetMapProvider_args
dataloader_args = cfg.data_source_ImplicitronDataSource_args.data_loader_map_provider_SequenceDataLoaderMapProvider_args
dataset_args = (
cfg.data_source_ImplicitronDataSource_args.dataset_map_provider_JsonIndexDatasetMapProvider_args
)
dataloader_args = (
cfg.data_source_ImplicitronDataSource_args.data_loader_map_provider_SequenceDataLoaderMapProvider_args
)
dataset_args.category = "skateboard"
dataset_args.test_restrict_sequence_id = 0
dataset_args.dataset_root = "manifold://co3d/tree/extracted"
@@ -102,7 +111,9 @@ class TestExperiment(unittest.TestCase):
cfg.training_loop_ImplicitronTrainingLoop_args.max_epochs = 2
cfg.training_loop_ImplicitronTrainingLoop_args.store_checkpoints = False
cfg.optimizer_factory_ImplicitronOptimizerFactory_args.lr_policy = "Exponential"
cfg.optimizer_factory_ImplicitronOptimizerFactory_args.exponential_lr_step_size = 2
cfg.optimizer_factory_ImplicitronOptimizerFactory_args.exponential_lr_step_size = (
2
)
if DEBUG:
experiment.dump_cfg(cfg)

View File

@@ -81,9 +81,8 @@ class TestOptimizerFactory(unittest.TestCase):
def test_param_overrides_self_param_group_assignment(self):
pa, pb, pc = [torch.nn.Parameter(data=torch.tensor(i * 1.0)) for i in range(3)]
na, nb = (
Node(params=[pa]),
Node(params=[pb], param_groups={"self": "pb_self", "p1": "pb_param"}),
na, nb = Node(params=[pa]), Node(
params=[pb], param_groups={"self": "pb_self", "p1": "pb_param"}
)
root = Node(children=[na, nb], params=[pc], param_groups={"m1": "pb_member"})
param_groups = self._get_param_groups(root)

View File

@@ -135,7 +135,6 @@ def visualize_reconstruction(
"device": device,
**render_flyaround_kwargs,
}
# pyrefly: ignore [bad-argument-type]
render_flyaround(**render_kwargs)

View File

@@ -84,9 +84,9 @@ def get_nerf_datasets(
if autodownload and any(not os.path.isfile(p) for p in (cameras_path, image_path)):
# Automatically download the data files if missing.
download_data([dataset_name], data_root=data_root)
download_data((dataset_name,), data_root=data_root)
train_data = torch.load(cameras_path, weights_only=True)
train_data = torch.load(cameras_path)
n_cameras = train_data["cameras"]["R"].shape[0]
_image_max_image_pixels = Image.MAX_IMAGE_PIXELS

View File

@@ -194,6 +194,7 @@ class Stats:
it = self.it[stat_set]
for stat in self.log_vars:
if stat not in self.stats[stat_set]:
self.stats[stat_set][stat] = AverageMeter()

View File

@@ -24,6 +24,7 @@ CONFIG_DIR = os.path.join(os.path.dirname(os.path.realpath(__file__)), "configs"
@hydra.main(config_path=CONFIG_DIR, config_name="lego")
def main(cfg: DictConfig):
# Device on which to run.
if torch.cuda.is_available():
device = "cuda"
@@ -62,7 +63,7 @@ def main(cfg: DictConfig):
raise ValueError(f"Model checkpoint {checkpoint_path} does not exist!")
print(f"Loading checkpoint {checkpoint_path}.")
loaded_data = torch.load(checkpoint_path, weights_only=True)
loaded_data = torch.load(checkpoint_path)
# Do not load the cached xy grid.
# - this allows setting an arbitrary evaluation image size.
state_dict = {

View File

@@ -42,6 +42,7 @@ class TestRaysampler(unittest.TestCase):
cameras, rays = [], []
for _ in range(batch_size):
R = random_rotations(1)
T = torch.randn(1, 3)
focal_length = torch.rand(1, 2) + 0.5

View File

@@ -25,6 +25,7 @@ CONFIG_DIR = os.path.join(os.path.dirname(os.path.realpath(__file__)), "configs"
@hydra.main(config_path=CONFIG_DIR, config_name="lego")
def main(cfg: DictConfig):
# Set the relevant seeds for reproducibility.
np.random.seed(cfg.seed)
torch.manual_seed(cfg.seed)
@@ -76,7 +77,7 @@ def main(cfg: DictConfig):
# Resume training if requested.
if cfg.resume and os.path.isfile(checkpoint_path):
print(f"Resuming from checkpoint {checkpoint_path}.")
loaded_data = torch.load(checkpoint_path, weights_only=True)
loaded_data = torch.load(checkpoint_path)
model.load_state_dict(loaded_data["model"])
stats = pickle.loads(loaded_data["stats"])
print(f" => resuming from epoch {stats.epoch}.")
@@ -218,6 +219,7 @@ def main(cfg: DictConfig):
# Validation
if epoch % cfg.validation_epoch_interval == 0 and epoch > 0:
# Sample a validation camera/image.
val_batch = next(val_dataloader.__iter__())
val_image, val_camera, camera_idx = val_batch[0].values()

View File

@@ -6,4 +6,4 @@
# pyre-unsafe
__version__ = "0.7.9"
__version__ = "0.7.8"

View File

@@ -17,12 +17,13 @@ Some functions which depend on PyTorch or Python versions.
def meshgrid_ij(
*A: Union[torch.Tensor, Sequence[torch.Tensor]],
*A: Union[torch.Tensor, Sequence[torch.Tensor]]
) -> Tuple[torch.Tensor, ...]: # pragma: no cover
"""
Like torch.meshgrid was before PyTorch 1.10.0, i.e. with indexing set to ij
"""
if (
# pyre-fixme[16]: Callable `meshgrid` has no attribute `__kwdefaults__`.
torch.meshgrid.__kwdefaults__ is not None
and "indexing" in torch.meshgrid.__kwdefaults__
):

View File

@@ -77,34 +77,17 @@ class _SymEig3x3(nn.Module):
if inputs.shape[-2:] != (3, 3):
raise ValueError("Only inputs of shape (..., 3, 3) are supported.")
# The thresholds and regularizers used below are absolute, so they are
# only meaningful for inputs of order 1: rescale each matrix to unit
# maximum magnitude and undo it on the eigenvalues. The clamp only
# matters for the all-zero matrix. The scale is detached because the
# decomposition is homogeneous of degree 1, so its Jacobian is
# homogeneous of degree 0 and holding the scale constant is exact.
scale = (
inputs.detach()
.abs()
.amax(dim=(-2, -1), keepdim=True)
.clamp(min=torch.finfo(inputs.dtype).tiny)
)
inputs = inputs / scale
inputs_diag = inputs.diagonal(dim1=-2, dim2=-1)
inputs_trace = inputs_diag.sum(-1)
q = inputs_trace / 3.0
# Calculate squared sum of elements outside the main diagonal / 2
p1 = (
torch.square(inputs).sum(dim=(-1, -2)) - torch.square(inputs_diag).sum(-1)
) / 2
p2 = torch.square(inputs_diag - q[..., None]).sum(dim=-1) + 2.0 * p1.clamp(
self._eps
)
# pyre-fixme[58]: `**` is not supported for operand types `Tensor` and `int`.
p1 = ((inputs**2).sum(dim=(-1, -2)) - (inputs_diag**2).sum(-1)) / 2
# pyre-fixme[58]: `**` is not supported for operand types `Tensor` and `int`.
p2 = ((inputs_diag - q[..., None]) ** 2).sum(dim=-1) + 2.0 * p1.clamp(self._eps)
p = torch.sqrt(p2 / 6.0)
# pyrefly: ignore [unsupported-operation]
B = (inputs - q[..., None, None] * self._identity) / p[..., None, None]
r = torch.det(B) / 2.0
@@ -121,9 +104,7 @@ class _SymEig3x3(nn.Module):
# Soft dispatch between the degenerate case (diagonal A) and general.
# diag_soft_cond -> 1.0 when p1 < 6 * eps and diag_soft_cond -> 0.0 otherwise.
# We use 6 * eps to take into account the error accumulated during the p1 summation
diag_soft_cond = torch.exp(-torch.square(p1 / (6 * self._eps))).detach()[
..., None
]
diag_soft_cond = torch.exp(-((p1 / (6 * self._eps)) ** 2)).detach()[..., None]
# Eigenvalues are the ordered elements of main diagonal in the degenerate case
diag_eigenvals, _ = torch.sort(inputs_diag, dim=-1)
@@ -134,7 +115,7 @@ class _SymEig3x3(nn.Module):
else:
eigenvecs = None
return eigenvals * scale.squeeze(-1), eigenvecs
return eigenvals, eigenvecs
def _construct_eigenvecs_set(
self, inputs: torch.Tensor, eigenvals: torch.Tensor
@@ -189,10 +170,8 @@ class _SymEig3x3(nn.Module):
"""
# Find the eigenvector corresponding to alpha0, its eigenvalue is distinct
# pyrefly: ignore [unsupported-operation]
ev0 = self._get_ev0(inputs - alpha0[..., None, None] * self._identity)
u, v = self._get_uv(ev0)
# pyrefly: ignore [unsupported-operation]
ev1 = self._get_ev1(inputs - alpha1[..., None, None] * self._identity, u, v)
# Third eigenvector is computed as the cross-product of the other two
ev2 = torch.cross(ev0, ev1, dim=-1)
@@ -220,7 +199,8 @@ class _SymEig3x3(nn.Module):
cross_products[..., :1, :]
)
norms_sq = torch.square(cross_products).sum(dim=-1)
# pyre-fixme[58]: `**` is not supported for operand types `Tensor` and `int`.
norms_sq = (cross_products**2).sum(dim=-1)
max_norms_index = norms_sq.argmax(dim=-1)
# Pick only the cross-product with highest squared norm for each input
@@ -267,7 +247,6 @@ class _SymEig3x3(nn.Module):
"""
min_idx = w.abs().argmin(dim=-1)
# pyrefly: ignore [bad-index]
rotation_2d = self._rotations_3d[min_idx].to(w)
u = F.normalize((rotation_2d @ w[..., None])[..., 0], dim=-1)

View File

@@ -32,9 +32,7 @@ __global__ void BallQueryKernel(
at::PackedTensorAccessor64<int64_t, 3, at::RestrictPtrTraits> idxs,
at::PackedTensorAccessor64<scalar_t, 3, at::RestrictPtrTraits> dists,
const int64_t K,
const float radius,
const float radius2,
const bool skip_points_outside_cube) {
const float radius2) {
const int64_t N = p1.size(0);
const int64_t chunks_per_cloud = (1 + (p1.size(1) - 1) / blockDim.x);
const int64_t chunks_to_do = N * chunks_per_cloud;
@@ -53,19 +51,7 @@ __global__ void BallQueryKernel(
// Iterate over points in p2 until desired count is reached or
// all points have been considered
for (int64_t j = 0, count = 0; j < lengths2[n] && count < K; ++j) {
if (skip_points_outside_cube) {
bool is_within_radius = true;
// Filter when any one coordinate is already outside the radius
for (int d = 0; is_within_radius && d < D; ++d) {
scalar_t abs_diff = fabs(p1[n][i][d] - p2[n][j][d]);
is_within_radius = (abs_diff <= radius);
}
if (!is_within_radius) {
continue;
}
}
// Else, calculate the distance between the points and compare
// Calculate the distance between the points
scalar_t dist2 = 0.0;
for (int d = 0; d < D; ++d) {
scalar_t diff = p1[n][i][d] - p2[n][j][d];
@@ -91,8 +77,7 @@ std::tuple<at::Tensor, at::Tensor> BallQueryCuda(
const at::Tensor& lengths1, // (N,)
const at::Tensor& lengths2, // (N,)
int K,
float radius,
bool skip_points_outside_cube) {
float radius) {
// Check inputs are on the same device
at::TensorArg p1_t{p1, "p1", 1}, p2_t{p2, "p2", 2},
lengths1_t{lengths1, "lengths1", 3}, lengths2_t{lengths2, "lengths2", 4};
@@ -135,9 +120,7 @@ std::tuple<at::Tensor, at::Tensor> BallQueryCuda(
idxs.packed_accessor64<int64_t, 3, at::RestrictPtrTraits>(),
dists.packed_accessor64<float, 3, at::RestrictPtrTraits>(),
K_64,
radius,
radius2,
skip_points_outside_cube);
radius2);
}));
AT_CUDA_CHECK(cudaGetLastError());

View File

@@ -25,9 +25,6 @@
// within the radius
// radius: the radius around each point within which the neighbors need to be
// located
// skip_points_outside_cube: If true, reduce multiplications of float values
// by not explicitly calculating distances to points that fall outside the
// D-cube with side length (2*radius) centered at each point in p1.
//
// Returns:
// p1_neighbor_idx: LongTensor of shape (N, P1, K), where
@@ -49,8 +46,7 @@ std::tuple<at::Tensor, at::Tensor> BallQueryCpu(
const at::Tensor& lengths1,
const at::Tensor& lengths2,
const int K,
const float radius,
const bool skip_points_outside_cube);
const float radius);
// CUDA implementation
std::tuple<at::Tensor, at::Tensor> BallQueryCuda(
@@ -59,8 +55,7 @@ std::tuple<at::Tensor, at::Tensor> BallQueryCuda(
const at::Tensor& lengths1,
const at::Tensor& lengths2,
const int K,
const float radius,
const bool skip_points_outside_cube);
const float radius);
// Implementation which is exposed
// Note: the backward pass reuses the KNearestNeighborBackward kernel
@@ -70,8 +65,7 @@ inline std::tuple<at::Tensor, at::Tensor> BallQuery(
const at::Tensor& lengths1,
const at::Tensor& lengths2,
int K,
float radius,
bool skip_points_outside_cube) {
float radius) {
if (p1.is_cuda() || p2.is_cuda()) {
#ifdef WITH_CUDA
CHECK_CUDA(p1);
@@ -82,20 +76,16 @@ inline std::tuple<at::Tensor, at::Tensor> BallQuery(
lengths1.contiguous(),
lengths2.contiguous(),
K,
radius,
skip_points_outside_cube);
radius);
#else
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(p1);
CHECK_CPU(p2);
return BallQueryCpu(
p1.contiguous(),
p2.contiguous(),
lengths1.contiguous(),
lengths2.contiguous(),
K,
radius,
skip_points_outside_cube);
radius);
}

View File

@@ -6,8 +6,8 @@
* LICENSE file in the root directory of this source tree.
*/
#include <math.h>
#include <torch/extension.h>
#include <queue>
#include <tuple>
std::tuple<at::Tensor, at::Tensor> BallQueryCpu(
@@ -16,8 +16,7 @@ std::tuple<at::Tensor, at::Tensor> BallQueryCpu(
const at::Tensor& lengths1,
const at::Tensor& lengths2,
int K,
float radius,
bool skip_points_outside_cube) {
float radius) {
const int N = p1.size(0);
const int P1 = p1.size(1);
const int D = p1.size(2);
@@ -39,16 +38,6 @@ std::tuple<at::Tensor, at::Tensor> BallQueryCpu(
const int64_t length2 = lengths2_a[n];
for (int64_t i = 0; i < length1; ++i) {
for (int64_t j = 0, count = 0; j < length2 && count < K; ++j) {
if (skip_points_outside_cube) {
bool is_within_radius = true;
for (int d = 0; is_within_radius && d < D; ++d) {
float abs_diff = fabs(p1_a[n][i][d] - p2_a[n][j][d]);
is_within_radius = (abs_diff <= radius);
}
if (!is_within_radius) {
continue;
}
}
float dist2 = 0;
for (int d = 0; d < D; ++d) {
float diff = p1_a[n][i][d] - p2_a[n][j][d];

View File

@@ -98,11 +98,6 @@ at::Tensor SigmoidAlphaBlendBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(distances);
CHECK_CPU(pix_to_face);
CHECK_CPU(alphas);
CHECK_CPU(grad_alphas);
return SigmoidAlphaBlendBackwardCpu(
grad_alphas, alphas, distances, pix_to_face, sigma);
}

View File

@@ -28,16 +28,17 @@ __global__ void alphaCompositeCudaForwardKernel(
const at::PackedTensorAccessor64<float, 4, at::RestrictPtrTraits> alphas,
const at::PackedTensorAccessor64<int64_t, 4, at::RestrictPtrTraits> points_idx) {
// clang-format on
const int64_t batch_size = result.size(0);
const int64_t C = features.size(0);
const int64_t H = points_idx.size(2);
const int64_t W = points_idx.size(3);
// Get the batch and index
const auto batch = blockIdx.x;
const int batch = blockIdx.x;
const int num_pixels = C * H * W;
const auto num_threads = gridDim.y * blockDim.x;
const auto tid = blockIdx.y * blockDim.x + threadIdx.x;
const int num_threads = gridDim.y * blockDim.x;
const int tid = blockIdx.y * blockDim.x + threadIdx.x;
// Iterate over each feature in each pixel
for (int pid = tid; pid < num_pixels; pid += num_threads) {
@@ -78,16 +79,17 @@ __global__ void alphaCompositeCudaBackwardKernel(
const at::PackedTensorAccessor64<float, 4, at::RestrictPtrTraits> alphas,
const at::PackedTensorAccessor64<int64_t, 4, at::RestrictPtrTraits> points_idx) {
// clang-format on
const int64_t batch_size = points_idx.size(0);
const int64_t C = features.size(0);
const int64_t H = points_idx.size(2);
const int64_t W = points_idx.size(3);
// Get the batch and index
const auto batch = blockIdx.x;
const int batch = blockIdx.x;
const int num_pixels = C * H * W;
const auto num_threads = gridDim.y * blockDim.x;
const auto tid = blockIdx.y * blockDim.x + threadIdx.x;
const int num_threads = gridDim.y * blockDim.x;
const int tid = blockIdx.y * blockDim.x + threadIdx.x;
// Parallelize over each feature in each pixel in images of size H * W,
// for each image in the batch of size batch_size

View File

@@ -74,9 +74,6 @@ torch::Tensor alphaCompositeForward(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(features);
CHECK_CPU(alphas);
CHECK_CPU(points_idx);
return alphaCompositeCpuForward(features, alphas, points_idx);
}
}
@@ -104,11 +101,6 @@ std::tuple<torch::Tensor, torch::Tensor> alphaCompositeBackward(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(grad_outputs);
CHECK_CPU(features);
CHECK_CPU(alphas);
CHECK_CPU(points_idx);
return alphaCompositeCpuBackward(
grad_outputs, features, alphas, points_idx);
}

View File

@@ -28,16 +28,17 @@ __global__ void weightedSumNormCudaForwardKernel(
const at::PackedTensorAccessor64<float, 4, at::RestrictPtrTraits> alphas,
const at::PackedTensorAccessor64<int64_t, 4, at::RestrictPtrTraits> points_idx) {
// clang-format on
const int64_t batch_size = result.size(0);
const int64_t C = features.size(0);
const int64_t H = points_idx.size(2);
const int64_t W = points_idx.size(3);
// Get the batch and index
const auto batch = blockIdx.x;
const int batch = blockIdx.x;
const int num_pixels = C * H * W;
const auto num_threads = gridDim.y * blockDim.x;
const auto tid = blockIdx.y * blockDim.x + threadIdx.x;
const int num_threads = gridDim.y * blockDim.x;
const int tid = blockIdx.y * blockDim.x + threadIdx.x;
// Parallelize over each feature in each pixel in images of size H * W,
// for each image in the batch of size batch_size
@@ -91,16 +92,17 @@ __global__ void weightedSumNormCudaBackwardKernel(
const at::PackedTensorAccessor64<float, 4, at::RestrictPtrTraits> alphas,
const at::PackedTensorAccessor64<int64_t, 4, at::RestrictPtrTraits> points_idx) {
// clang-format on
const int64_t batch_size = points_idx.size(0);
const int64_t C = features.size(0);
const int64_t H = points_idx.size(2);
const int64_t W = points_idx.size(3);
// Get the batch and index
const auto batch = blockIdx.x;
const int batch = blockIdx.x;
const int num_pixels = C * W * H;
const auto num_threads = gridDim.y * blockDim.x;
const auto tid = blockIdx.y * blockDim.x + threadIdx.x;
const int num_threads = gridDim.y * blockDim.x;
const int tid = blockIdx.y * blockDim.x + threadIdx.x;
// Parallelize over each feature in each pixel in images of size H * W,
// for each image in the batch of size batch_size

View File

@@ -73,10 +73,6 @@ torch::Tensor weightedSumNormForward(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(features);
CHECK_CPU(alphas);
CHECK_CPU(points_idx);
return weightedSumNormCpuForward(features, alphas, points_idx);
}
}
@@ -104,11 +100,6 @@ std::tuple<torch::Tensor, torch::Tensor> weightedSumNormBackward(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(grad_outputs);
CHECK_CPU(features);
CHECK_CPU(alphas);
CHECK_CPU(points_idx);
return weightedSumNormCpuBackward(
grad_outputs, features, alphas, points_idx);
}

View File

@@ -26,16 +26,17 @@ __global__ void weightedSumCudaForwardKernel(
const at::PackedTensorAccessor64<float, 4, at::RestrictPtrTraits> alphas,
const at::PackedTensorAccessor64<int64_t, 4, at::RestrictPtrTraits> points_idx) {
// clang-format on
const int64_t batch_size = result.size(0);
const int64_t C = features.size(0);
const int64_t H = points_idx.size(2);
const int64_t W = points_idx.size(3);
// Get the batch and index
const auto batch = blockIdx.x;
const int batch = blockIdx.x;
const int num_pixels = C * H * W;
const auto num_threads = gridDim.y * blockDim.x;
const auto tid = blockIdx.y * blockDim.x + threadIdx.x;
const int num_threads = gridDim.y * blockDim.x;
const int tid = blockIdx.y * blockDim.x + threadIdx.x;
// Parallelize over each feature in each pixel in images of size H * W,
// for each image in the batch of size batch_size
@@ -73,16 +74,17 @@ __global__ void weightedSumCudaBackwardKernel(
const at::PackedTensorAccessor64<float, 4, at::RestrictPtrTraits> alphas,
const at::PackedTensorAccessor64<int64_t, 4, at::RestrictPtrTraits> points_idx) {
// clang-format on
const int64_t batch_size = points_idx.size(0);
const int64_t C = features.size(0);
const int64_t H = points_idx.size(2);
const int64_t W = points_idx.size(3);
// Get the batch and index
const auto batch = blockIdx.x;
const int batch = blockIdx.x;
const int num_pixels = C * H * W;
const auto num_threads = gridDim.y * blockDim.x;
const auto tid = blockIdx.y * blockDim.x + threadIdx.x;
const int num_threads = gridDim.y * blockDim.x;
const int tid = blockIdx.y * blockDim.x + threadIdx.x;
// Iterate over each pixel to compute the contribution to the
// gradient for the features and weights

View File

@@ -72,9 +72,6 @@ torch::Tensor weightedSumForward(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(features);
CHECK_CPU(alphas);
CHECK_CPU(points_idx);
return weightedSumCpuForward(features, alphas, points_idx);
}
}
@@ -101,11 +98,6 @@ std::tuple<torch::Tensor, torch::Tensor> weightedSumBackward(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(grad_outputs);
CHECK_CPU(features);
CHECK_CPU(alphas);
CHECK_CPU(points_idx);
return weightedSumCpuBackward(grad_outputs, features, alphas, points_idx);
}
}

View File

@@ -8,6 +8,7 @@
// clang-format off
#include "./pulsar/global.h" // Include before <torch/extension.h>.
#include <torch/extension.h>
// clang-format on
#include "./pulsar/pytorch/renderer.h"
#include "./pulsar/pytorch/tensor_util.h"
@@ -98,22 +99,22 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
m.def("marching_cubes", &MarchingCubes);
// Pulsar.
// Pulsar not enabled on AMD.
#ifdef PULSAR_LOGGING_ENABLED
c10::ShowLogInfoToStderr();
#endif
py::class_<
pulsar::pytorch::Renderer,
std::shared_ptr<pulsar::pytorch::Renderer>>(m, "PulsarRenderer")
.def(
py::init<
const uint&,
const uint&,
const uint&,
const bool&,
const bool&,
const float&,
const uint&,
const uint&>())
.def(py::init<
const uint&,
const uint&,
const uint&,
const bool&,
const bool&,
const float&,
const uint&,
const uint&>())
.def(
"__eq__",
[](const pulsar::pytorch::Renderer& a,
@@ -148,10 +149,10 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
py::arg("gamma"),
py::arg("max_depth"),
py::arg("min_depth") /* = 0.f*/,
py::arg("bg_col") /* = std::nullopt not exposed properly in
pytorch 1.1. */
py::arg(
"bg_col") /* = at::nullopt not exposed properly in pytorch 1.1. */
,
py::arg("opacity") /* = std::nullopt ... */,
py::arg("opacity") /* = at::nullopt ... */,
py::arg("percent_allowed_difference") = 0.01f,
py::arg("max_n_hits") = MAX_UINT,
py::arg("mode") = 0)

View File

@@ -60,8 +60,6 @@ std::tuple<at::Tensor, at::Tensor> FaceAreasNormalsForward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(verts);
CHECK_CPU(faces);
return FaceAreasNormalsForwardCpu(verts, faces);
}
@@ -82,9 +80,5 @@ at::Tensor FaceAreasNormalsBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(grad_areas);
CHECK_CPU(grad_normals);
CHECK_CPU(verts);
CHECK_CPU(faces);
return FaceAreasNormalsBackwardCpu(grad_areas, grad_normals, verts, faces);
}

View File

@@ -20,14 +20,14 @@ __global__ void GatherScatterCudaKernel(
const size_t V,
const size_t D,
const size_t E) {
const auto tid = threadIdx.x;
const int tid = threadIdx.x;
// Reverse the vertex order if backward.
const int v0_idx = backward ? 1 : 0;
const int v1_idx = backward ? 0 : 1;
// Edges are split evenly across the blocks.
for (auto e = blockIdx.x; e < E; e += gridDim.x) {
for (int e = blockIdx.x; e < E; e += gridDim.x) {
// Get indices of vertices which form the edge.
const int64_t v0 = edges[2 * e + v0_idx];
const int64_t v1 = edges[2 * e + v1_idx];
@@ -35,7 +35,7 @@ __global__ void GatherScatterCudaKernel(
// Split vertex features evenly across threads.
// This implementation will be quite wasteful when D<128 since there will be
// a lot of threads doing nothing.
for (auto d = tid; d < D; d += blockDim.x) {
for (int d = tid; d < D; d += blockDim.x) {
const float val = input[v1 * D + d];
float* address = output + v0 * D + d;
atomicAdd(address, val);

View File

@@ -53,7 +53,5 @@ at::Tensor GatherScatter(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(input);
CHECK_CPU(edges);
return GatherScatterCpu(input, edges, directed, backward);
}

View File

@@ -20,8 +20,8 @@ __global__ void InterpFaceAttrsForwardKernel(
const size_t P,
const size_t F,
const size_t D) {
const auto tid = threadIdx.x + blockIdx.x * blockDim.x;
const auto num_threads = blockDim.x * gridDim.x;
const int tid = threadIdx.x + blockIdx.x * blockDim.x;
const int num_threads = blockDim.x * gridDim.x;
for (int pd = tid; pd < P * D; pd += num_threads) {
const int p = pd / D;
const int d = pd % D;
@@ -93,8 +93,8 @@ __global__ void InterpFaceAttrsBackwardKernel(
const size_t P,
const size_t F,
const size_t D) {
const auto tid = threadIdx.x + blockIdx.x * blockDim.x;
const auto num_threads = blockDim.x * gridDim.x;
const int tid = threadIdx.x + blockIdx.x * blockDim.x;
const int num_threads = blockDim.x * gridDim.x;
for (int pd = tid; pd < P * D; pd += num_threads) {
const int p = pd / D;
const int d = pd % D;

View File

@@ -57,8 +57,6 @@ at::Tensor InterpFaceAttrsForward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(face_attrs);
CHECK_CPU(barycentric_coords);
return InterpFaceAttrsForwardCpu(pix_to_face, barycentric_coords, face_attrs);
}
@@ -108,9 +106,6 @@ std::tuple<at::Tensor, at::Tensor> InterpFaceAttrsBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(face_attrs);
CHECK_CPU(barycentric_coords);
CHECK_CPU(grad_pix_attrs);
return InterpFaceAttrsBackwardCpu(
pix_to_face, barycentric_coords, face_attrs, grad_pix_attrs);
}

View File

@@ -44,7 +44,5 @@ inline std::tuple<at::Tensor, at::Tensor> IoUBox3D(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(boxes1);
CHECK_CPU(boxes2);
return IoUBox3DCpu(boxes1.contiguous(), boxes2.contiguous());
}

View File

@@ -7,6 +7,10 @@
*/
#include <torch/extension.h>
#include <torch/torch.h>
#include <list>
#include <numeric>
#include <queue>
#include <tuple>
#include "iou_box3d/iou_utils.h"

View File

@@ -461,8 +461,10 @@ __device__ inline std::tuple<float3, float3> ArgMaxVerts(
__device__ inline bool IsCoplanarTriTri(
const FaceVerts& tri1,
const FaceVerts& tri2) {
const float3 tri1_ctr = FaceCenter({tri1.v0, tri1.v1, tri1.v2});
const float3 tri1_n = FaceNormal({tri1.v0, tri1.v1, tri1.v2});
const float3 tri2_ctr = FaceCenter({tri2.v0, tri2.v1, tri2.v2});
const float3 tri2_n = FaceNormal({tri2.v0, tri2.v1, tri2.v2});
// Check if parallel
@@ -498,6 +500,7 @@ __device__ inline bool IsCoplanarTriPlane(
const FaceVerts& tri,
const FaceVerts& plane,
const float3& normal) {
const float3 tri_ctr = FaceCenter({tri.v0, tri.v1, tri.v2});
const float3 nt = FaceNormal({tri.v0, tri.v1, tri.v2});
// check if parallel
@@ -725,7 +728,7 @@ __device__ inline int BoxIntersections(
}
}
// Update the face_verts_out tris
num_tris = min(MAX_TRIS, offset);
num_tris = offset;
for (int j = 0; j < num_tris; ++j) {
face_verts_out[j] = tri_verts_updated[j];
}

View File

@@ -74,8 +74,6 @@ std::tuple<at::Tensor, at::Tensor> KNearestNeighborIdx(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(p1);
CHECK_CPU(p2);
return KNearestNeighborIdxCpu(p1, p2, lengths1, lengths2, norm, K);
}
@@ -142,8 +140,6 @@ std::tuple<at::Tensor, at::Tensor> KNearestNeighborBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(p1);
CHECK_CPU(p2);
return KNearestNeighborBackwardCpu(
p1, p2, lengths1, lengths2, idxs, norm, grad_dists);
}

View File

@@ -58,6 +58,5 @@ inline std::tuple<at::Tensor, at::Tensor, at::Tensor> MarchingCubes(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(vol);
return MarchingCubesCpu(vol.contiguous(), isolevel);
}

View File

@@ -34,18 +34,6 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> MarchingCubesCpu(
const int H = vol.size(1);
const int W = vol.size(2);
// The edge-id hash packs two grid-point ids as v1_id * (W*H*D) + v2_id, whose
// maximum value is (W*H*D)^2 - 1. Guard against signed int64 overflow so an
// oversized volume fails loudly instead of silently producing colliding edge
// ids and a corrupted mesh. floor(sqrt(INT64_MAX)) == 3037000499 (~1448^3).
const int64_t num_grid_points = (int64_t)W * H * D;
TORCH_CHECK(
num_grid_points <= 3037000499LL,
"Volume too large for CPU marching cubes: W*H*D (",
num_grid_points,
") exceeds 3037000499 (~1448^3), which would overflow the int64 edge-id "
"hash and corrupt the mesh.");
// Create tensor accessors
auto vol_a = vol.accessor<float, 3>();
// edge_id_to_v maps from an edge id to a vertex position
@@ -78,23 +66,19 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> MarchingCubesCpu(
tri.push_back(edge);
ps.push_back(interp_points[e]);
if (ps.size() == 3) {
// Check if the triangle face is degenerate. A triangle face
// is degenerate if any of the two vertices share the same 3D
// position
if (ps[0] != ps[1] && ps[1] != ps[2] && ps[2] != ps[0]) {
for (int k = 0; k < 3; k++) {
int64_t v = tri.at(k);
edge_id_to_v[v] = ps.at(k);
if (!uniq_edge_id.count(v)) {
uniq_edge_id[v] = verts.size();
verts.push_back(edge_id_to_v[v]);
}
faces.push_back(uniq_edge_id[v]);
// Check if the triangle face is degenerate. A triangle face
// is degenerate if any of the two verices share the same 3D position
if ((j + 1) % 3 == 0 && ps[0] != ps[1] && ps[1] != ps[2] &&
ps[2] != ps[0]) {
for (int k = 0; k < 3; k++) {
int64_t v = tri.at(k);
edge_id_to_v[v] = ps.at(k);
if (!uniq_edge_id.count(v)) {
uniq_edge_id[v] = verts.size();
verts.push_back(edge_id_to_v[v]);
}
faces.push_back(uniq_edge_id[v]);
}
// Clear unconditionally - a rejected degenerate triangle must not
// corrupt the buffer for the next triangle in this cube.
tri.clear();
ps.clear();
} // endif

View File

@@ -126,7 +126,8 @@ struct Cube {
// edge with an integer to address floating point precision issue.
//
// Args:
// edge: edge (ID) whose two endpoint vertices are hashed
// v1_id: global id of vertex 1
// v2_id: global id of vertex 2
// W: width of the 3d grid
// H: height of the 3d grid
// D: depth of the 3d grid
@@ -134,19 +135,11 @@ struct Cube {
// Returns:
// hashing for a pair of vertex ids
//
int64_t HashVpair(const int edge, int64_t W, int64_t H, int64_t D) const {
int64_t HashVpair(const int edge, int W, int H, int D) {
const int v1 = _EDGE_TO_VERTICES[edge][0];
const int v2 = _EDGE_TO_VERTICES[edge][1];
// p[v].x/y/z hold integral corner coordinates, so casting to int64_t is
// exact. We cast before multiplication to avoid float32 precision loss at
// grids > 256^3.
const int64_t v1_id =
(int64_t)p[v1].x + (int64_t)p[v1].y * W + (int64_t)p[v1].z * W * H;
const int64_t v2_id =
(int64_t)p[v2].x + (int64_t)p[v2].y * W + (int64_t)p[v2].z * W * H;
// v_id max is (W*H*D - 1). A stride of W*H*D perfectly separates v1 and v2.
// Note: This hash fits in int64_t safely up to grids of ~1448^3.
const int64_t stride = W * H * D;
return v1_id * stride + v2_id;
const int v1_id = p[v1].x + p[v1].y * W + p[v1].z * W * H;
const int v2_id = p[v2].x + p[v2].y * W + p[v2].z * W * H;
return (int64_t)v1_id * (W + W * H + W * H * D) + (int64_t)v2_id;
}
};

View File

@@ -88,8 +88,6 @@ at::Tensor PackedToPadded(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(inputs_packed);
CHECK_CPU(first_idxs);
return PackedToPaddedCpu(inputs_packed, first_idxs, max_size);
}
@@ -107,7 +105,5 @@ at::Tensor PaddedToPacked(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(inputs_padded);
CHECK_CPU(first_idxs);
return PaddedToPackedCpu(inputs_padded, first_idxs, num_inputs);
}

View File

@@ -174,8 +174,8 @@ std::tuple<at::Tensor, at::Tensor> HullHullDistanceForwardCpu(
at::Tensor idxs = at::zeros({A_N,}, as_first_idx.options());
// clang-format on
auto as_a = as.accessor<float, H1 == 1 ? 2 : 3>();
auto bs_a = bs.accessor<float, H2 == 1 ? 2 : 3>();
auto as_a = as.accessor < float, H1 == 1 ? 2 : 3 > ();
auto bs_a = bs.accessor < float, H2 == 1 ? 2 : 3 > ();
auto as_first_idx_a = as_first_idx.accessor<int64_t, 1>();
auto bs_first_idx_a = bs_first_idx.accessor<int64_t, 1>();
auto dists_a = dists.accessor<float, 1>();
@@ -230,10 +230,10 @@ std::tuple<at::Tensor, at::Tensor> HullHullDistanceBackwardCpu(
at::Tensor grad_as = at::zeros_like(as);
at::Tensor grad_bs = at::zeros_like(bs);
auto as_a = as.accessor<float, H1 == 1 ? 2 : 3>();
auto bs_a = bs.accessor<float, H2 == 1 ? 2 : 3>();
auto grad_as_a = grad_as.accessor<float, H1 == 1 ? 2 : 3>();
auto grad_bs_a = grad_bs.accessor<float, H2 == 1 ? 2 : 3>();
auto as_a = as.accessor < float, H1 == 1 ? 2 : 3 > ();
auto bs_a = bs.accessor < float, H2 == 1 ? 2 : 3 > ();
auto grad_as_a = grad_as.accessor < float, H1 == 1 ? 2 : 3 > ();
auto grad_bs_a = grad_bs.accessor < float, H2 == 1 ? 2 : 3 > ();
auto idx_bs_a = idx_bs.accessor<int64_t, 1>();
auto grad_dists_a = grad_dists.accessor<float, 1>();

View File

@@ -110,7 +110,7 @@ __global__ void DistanceForwardKernel(
__syncthreads();
// Perform reduction in shared memory.
for (auto s = blockDim.x / 2; s > 32; s >>= 1) {
for (int s = blockDim.x / 2; s > 32; s >>= 1) {
if (tid < s) {
if (min_dists[tid] > min_dists[tid + s]) {
min_dists[tid] = min_dists[tid + s];
@@ -502,8 +502,8 @@ __global__ void PointFaceArrayForwardKernel(
const float3* tris_f3 = (float3*)tris;
// Parallelize over P * S computations
const auto num_threads = gridDim.x * blockDim.x;
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
const int num_threads = gridDim.x * blockDim.x;
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
for (int t_i = tid; t_i < P * T; t_i += num_threads) {
const int t = t_i / P; // segment index.
@@ -576,8 +576,8 @@ __global__ void PointFaceArrayBackwardKernel(
const float3* tris_f3 = (float3*)tris;
// Parallelize over P * S computations
const auto num_threads = gridDim.x * blockDim.x;
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
const int num_threads = gridDim.x * blockDim.x;
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
for (int t_i = tid; t_i < P * T; t_i += num_threads) {
const int t = t_i / P; // triangle index.
@@ -683,8 +683,8 @@ __global__ void PointEdgeArrayForwardKernel(
float3* segms_f3 = (float3*)segms;
// Parallelize over P * S computations
const auto num_threads = gridDim.x * blockDim.x;
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
const int num_threads = gridDim.x * blockDim.x;
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
for (int t_i = tid; t_i < P * S; t_i += num_threads) {
const int s = t_i / P; // segment index.
@@ -752,8 +752,8 @@ __global__ void PointEdgeArrayBackwardKernel(
float3* segms_f3 = (float3*)segms;
// Parallelize over P * S computations
const auto num_threads = gridDim.x * blockDim.x;
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
const int num_threads = gridDim.x * blockDim.x;
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
for (int t_i = tid; t_i < P * S; t_i += num_threads) {
const int s = t_i / P; // segment index.

View File

@@ -88,10 +88,6 @@ std::tuple<torch::Tensor, torch::Tensor> PointFaceDistanceForward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(points_first_idx);
CHECK_CPU(tris);
CHECK_CPU(tris_first_idx);
return PointFaceDistanceForwardCpu(
points, points_first_idx, tris, tris_first_idx, min_triangle_area);
}
@@ -147,10 +143,6 @@ std::tuple<torch::Tensor, torch::Tensor> PointFaceDistanceBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(tris);
CHECK_CPU(idx_points);
CHECK_CPU(grad_dists);
return PointFaceDistanceBackwardCpu(
points, tris, idx_points, grad_dists, min_triangle_area);
}
@@ -229,10 +221,6 @@ std::tuple<torch::Tensor, torch::Tensor> FacePointDistanceForward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(points_first_idx);
CHECK_CPU(tris);
CHECK_CPU(tris_first_idx);
return FacePointDistanceForwardCpu(
points, points_first_idx, tris, tris_first_idx, min_triangle_area);
}
@@ -289,10 +277,6 @@ std::tuple<torch::Tensor, torch::Tensor> FacePointDistanceBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(tris);
CHECK_CPU(idx_tris);
CHECK_CPU(grad_dists);
return FacePointDistanceBackwardCpu(
points, tris, idx_tris, grad_dists, min_triangle_area);
}
@@ -362,10 +346,6 @@ std::tuple<torch::Tensor, torch::Tensor> PointEdgeDistanceForward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(points_first_idx);
CHECK_CPU(segms);
CHECK_CPU(segms_first_idx);
return PointEdgeDistanceForwardCpu(
points, points_first_idx, segms, segms_first_idx, max_points);
}
@@ -416,10 +396,6 @@ std::tuple<torch::Tensor, torch::Tensor> PointEdgeDistanceBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(segms);
CHECK_CPU(idx_points);
CHECK_CPU(grad_dists);
return PointEdgeDistanceBackwardCpu(points, segms, idx_points, grad_dists);
}
@@ -488,10 +464,6 @@ std::tuple<torch::Tensor, torch::Tensor> EdgePointDistanceForward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(points_first_idx);
CHECK_CPU(segms);
CHECK_CPU(segms_first_idx);
return EdgePointDistanceForwardCpu(
points, points_first_idx, segms, segms_first_idx, max_segms);
}
@@ -542,10 +514,6 @@ std::tuple<torch::Tensor, torch::Tensor> EdgePointDistanceBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(segms);
CHECK_CPU(idx_segms);
CHECK_CPU(grad_dists);
return EdgePointDistanceBackwardCpu(points, segms, idx_segms, grad_dists);
}
@@ -599,8 +567,6 @@ torch::Tensor PointFaceArrayDistanceForward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(tris);
return PointFaceArrayDistanceForwardCpu(points, tris, min_triangle_area);
}
@@ -647,9 +613,6 @@ std::tuple<torch::Tensor, torch::Tensor> PointFaceArrayDistanceBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(tris);
CHECK_CPU(grad_dists);
return PointFaceArrayDistanceBackwardCpu(
points, tris, grad_dists, min_triangle_area);
}
@@ -698,8 +661,6 @@ torch::Tensor PointEdgeArrayDistanceForward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(segms);
return PointEdgeArrayDistanceForwardCpu(points, segms);
}
@@ -742,8 +703,5 @@ std::tuple<torch::Tensor, torch::Tensor> PointEdgeArrayDistanceBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(segms);
CHECK_CPU(grad_dists);
return PointEdgeArrayDistanceBackwardCpu(points, segms, grad_dists);
}

View File

@@ -104,12 +104,6 @@ inline void PointsToVolumesForward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points_3d);
CHECK_CPU(points_features);
CHECK_CPU(volume_densities);
CHECK_CPU(volume_features);
CHECK_CPU(grid_sizes);
CHECK_CPU(mask);
PointsToVolumesForwardCpu(
points_3d,
points_features,
@@ -189,14 +183,6 @@ inline void PointsToVolumesBackward(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points_3d);
CHECK_CPU(points_features);
CHECK_CPU(grid_sizes);
CHECK_CPU(mask);
CHECK_CPU(grad_volume_densities);
CHECK_CPU(grad_volume_features);
CHECK_CPU(grad_points_3d);
CHECK_CPU(grad_points_features);
PointsToVolumesBackwardCpu(
points_3d,
points_features,

View File

@@ -8,7 +8,9 @@
#include <torch/csrc/autograd/VariableTypeUtils.h>
#include <torch/extension.h>
#include <algorithm>
#include <cmath>
#include <thread>
#include <vector>
// In the x direction, the location {0, ..., grid_size_x - 1} correspond to

View File

@@ -15,8 +15,8 @@
#endif
#if defined(_WIN64) || defined(_WIN32)
using uint = unsigned int;
using ushort = unsigned short;
#define uint unsigned int
#define ushort unsigned short
#endif
#include "./logging.h" // <- include before torch/extension.h

View File

@@ -125,43 +125,29 @@ INLINE DEVICE float3 WARP_SUM_FLOAT3(
// Floating point.
// #define FMUL(a, b) __fmul_rn((a), (b))
#define FMUL(a, b) ((a) * (b))
#define FDIV(a, b) __fdiv_rn((a), (b))
// #define FSUB(a, b) __fsub_rn((a), (b))
#define FSUB(a, b) ((a) - (b))
// HIP has no per-instruction rounding-mode override (round-to-nearest-even is
// the default at the hardware level), so the CUDA *_rn intrinsics have no
// runtime equivalent and we use plain operators. The HIP/clang compiler may
// fuse a+b*c into a single-rounding FMA where CUDA's _rn would have prevented
// it; if that becomes a numerical issue, add `-ffp-contract=off` to the pulsar
// nvcc_args in setup.py.
#if defined(USE_ROCM)
#define FDIV(a, b) ((a) / (b))
#define FADD(a, b) ((a) + (b))
#define FSQRT(a) sqrtf(a)
#define FPOW(a, b) powf((a), (b))
#define FSATURATE(x) fmaxf(0.0f, fminf(1.0f, (x)))
#define FMA(x, y, z) fmaf((x), (y), (z))
#define FRCP(x) (1.0f / (x))
#else
#define FDIV(a, b) __fdiv_rn((a), (b))
#define FADD(a, b) __fadd_rn((a), (b))
#define FSQRT(a) __fsqrt_rn(a)
#define FPOW(a, b) __powf((a), (b))
#define FSATURATE(x) __saturatef(x)
/** Calculates x*y+z. */
#define FMA(x, y, z) __fmaf_rn((x), (y), (z))
#define FRCP(x) __frcp_rn(x)
#endif
#define FEXP(a) fasterexp(a)
#define FLN(a) fasterlog(a)
#define FPOW(a, b) __powf((a), (b))
#define FMAX(a, b) fmax((a), (b))
#define FMIN(a, b) fmin((a), (b))
#define FCEIL(a) ceilf(a)
#define FFLOOR(a) floorf(a)
#define FROUND(x) nearbyintf(x)
#define FSATURATE(x) __saturatef(x)
#define FABS(a) abs(a)
#define IASF(a, loc) (loc) = __int_as_float(a)
#define FASI(a, loc) (loc) = __float_as_int(a)
#define FABSLEQAS(a, b, c) \
((a) <= (b) ? FSUB((b), (a)) <= (c) : FSUB((a), (b)) < (c))
/** Calculates x*y+z. */
#define FMA(x, y, z) __fmaf_rn((x), (y), (z))
#define I2F(a) __int2float_rn(a)
#define FRCP(x) __frcp_rn(x)
#if !defined(USE_ROCM)
__device__ static float atomicMax(float* address, float val) {
int* address_as_i = (int*)address;
@@ -215,16 +201,8 @@ __device__ static float atomicMin(float* address, float val) {
ATOMICADD(&((PTR)->x), VAL.x); \
ATOMICADD(&((PTR)->y), VAL.y); \
ATOMICADD(&((PTR)->z), VAL.z);
#if !defined(USE_ROCM) && (CUDART_VERSION >= 10000) && (__CUDA_ARCH__ >= 600)
#if (CUDART_VERSION >= 10000) && (__CUDA_ARCH__ >= 600)
#define ATOMICADD_B(PTR, VAL) atomicAdd_block((PTR), (VAL))
#elif defined(USE_ROCM)
// HIP has no atomicAdd_block, but the semantic equivalent is a relaxed
// fetch_add scoped to the workgroup (HIP's name for a CUDA thread block).
// This avoids device-wide L2-coherent atomics for what are block-local
// counters in pulsar's inner sphere-loading loop.
#define ATOMICADD_B(PTR, VAL) \
__hip_atomic_fetch_add( \
(PTR), (VAL), __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP)
#else
#define ATOMICADD_B(PTR, VAL) ATOMICADD(PTR, VAL)
#endif
@@ -439,7 +417,7 @@ __device__ static float atomicMin(float* address, float val) {
(OUT_PTR), \
(NUM_SELECTED_PTR), \
(NUM_ITEMS), \
(STREAM));
stream = (STREAM));
#define COPY_HOST_DEV(PTR_D, PTR_H, TYPE, SIZE) \
HANDLECUDA(cudaMemcpy( \

View File

@@ -114,6 +114,8 @@ INLINE float atomicMin(float* address, float val) {
#define FMA(x, y, z) ((x) * (y) + (z))
#define I2F(a) static_cast<float>(a)
#define FRCP(x) (1.f / (x))
#define IASF(x, loc) memcpy(&(loc), &(x), sizeof(x))
#define FASI(x, loc) memcpy(&(loc), &(x), sizeof(x))
#define DMAX(a, b) std::max((a), (b))
#define DMIN(a, b) std::min((a), (b))
#define DSATURATE(a) DMIN(1., DMAX(0., (a)))
@@ -355,11 +357,11 @@ void MAX_WS(
//
//
#define END_PARALLEL() \
end_parallel:; \
end_parallel :; \
}
#define END_PARALLEL_NORET() }
#define END_PARALLEL_2D() \
end_parallel:; \
end_parallel :; \
} \
}
#define END_PARALLEL_2D_NORET() \

View File

@@ -70,6 +70,11 @@ struct CamGradInfo {
float3 pixel_dir_y;
};
// TODO: remove once https://github.com/NVlabs/cub/issues/172 is resolved.
struct IntWrapper {
int val;
};
} // namespace pulsar
#endif

View File

@@ -9,11 +9,7 @@
#ifndef PULSAR_NATIVE_COMMANDS_ROUTING_H_
#define PULSAR_NATIVE_COMMANDS_ROUTING_H_
#include <cstring>
#include "../global.h"
#if __cplusplus >= 202002L
#include <bit>
#endif
// Commands available everywhere.
#define MALLOC_HOST(VAR, TYPE, SIZE) \
@@ -36,37 +32,4 @@
#include "../host/commands.h"
#endif
// Host + device int <-> float bit punning. Replaces IASF / FASI macros.
IHD float pulsar_int_as_float(int v) noexcept {
#ifdef __CUDA_ARCH__
return __int_as_float(v);
#elif defined(__HIP_DEVICE_COMPILE__)
return __int_as_float(v);
#else
#if __cpp_lib_bit_cast >= 201806L
return std::bit_cast<float>(v);
#else
float f;
std::memcpy(&f, &v, sizeof(f));
return f;
#endif
#endif
}
IHD int pulsar_float_as_int(float v) noexcept {
#ifdef __CUDA_ARCH__
return __float_as_int(v);
#elif defined(__HIP_DEVICE_COMPILE__)
return __float_as_int(v);
#else
#if __cpp_lib_bit_cast >= 201806L
return std::bit_cast<int>(v);
#else
int i;
std::memcpy(&i, &v, sizeof(i));
return i;
#endif
#endif
}
#endif

View File

@@ -149,6 +149,11 @@ IHD CamGradInfo operator*(const CamGradInfo& a, const float& b) {
return res;
}
IHD IntWrapper operator+(const IntWrapper& a, const IntWrapper& b) {
IntWrapper res;
res.val = a.val + b.val;
return res;
}
} // namespace pulsar
#endif

View File

@@ -155,8 +155,8 @@ void backward(
stream);
CHECKLAUNCH();
SUM_WS(
self->ids_sorted_d,
self->n_grad_contributions_d,
(IntWrapper*)(self->ids_sorted_d),
(IntWrapper*)(self->n_grad_contributions_d),
static_cast<int>(num_balls),
self->workspace_d,
self->workspace_size,

View File

@@ -121,8 +121,8 @@ GLOBAL void calc_gradients(
grad_im_l[2]);
// Start processing.
for (int grad_idx = 0; grad_idx < n_track; ++grad_idx) {
int sphere_idx =
pulsar_float_as_int(forw_info_d[fwi_loc + 3 + 2 * grad_idx]);
int sphere_idx;
FASI(forw_info_d[fwi_loc + 3 + 2 * grad_idx], sphere_idx);
PASSERT(
sphere_idx == -1 ||
sphere_idx >= 0 && static_cast<uint>(sphere_idx) < num_balls);

View File

@@ -52,7 +52,7 @@ HOST void construct(
self->cam.film_width = width;
self->cam.film_height = height;
self->max_num_balls = max_num_balls;
MALLOC(self->result_d, float, width * height * n_channels);
MALLOC(self->result_d, float, width* height* n_channels);
self->cam.orthogonal_projection = orthogonal_projection;
self->cam.right_handed = right_handed_system;
self->cam.background_normalization_depth = background_normalization_depth;
@@ -93,7 +93,7 @@ HOST void construct(
MALLOC(self->di_sorted_d, DrawInfo, max_num_balls);
MALLOC(self->region_flags_d, char, max_num_balls);
MALLOC(self->num_selected_d, size_t, 1);
MALLOC(self->forw_info_d, float, width * height * (3 + 2 * n_track));
MALLOC(self->forw_info_d, float, width* height * (3 + 2 * n_track));
MALLOC(self->min_max_pixels_d, IntersectInfo, 1);
MALLOC(self->grad_pos_d, float3, max_num_balls);
MALLOC(self->grad_col_d, float, max_num_balls* n_channels);

View File

@@ -18,89 +18,68 @@ namespace Renderer {
template <bool DEV>
HOST void destruct(Renderer* self) {
if (self->result_d != NULL) {
if (self->result_d != NULL)
FREE(self->result_d);
}
self->result_d = NULL;
if (self->min_depth_d != NULL) {
if (self->min_depth_d != NULL)
FREE(self->min_depth_d);
}
self->min_depth_d = NULL;
if (self->min_depth_sorted_d != NULL) {
if (self->min_depth_sorted_d != NULL)
FREE(self->min_depth_sorted_d);
}
self->min_depth_sorted_d = NULL;
if (self->ii_d != NULL) {
if (self->ii_d != NULL)
FREE(self->ii_d);
}
self->ii_d = NULL;
if (self->ii_sorted_d != NULL) {
if (self->ii_sorted_d != NULL)
FREE(self->ii_sorted_d);
}
self->ii_sorted_d = NULL;
if (self->ids_d != NULL) {
if (self->ids_d != NULL)
FREE(self->ids_d);
}
self->ids_d = NULL;
if (self->ids_sorted_d != NULL) {
if (self->ids_sorted_d != NULL)
FREE(self->ids_sorted_d);
}
self->ids_sorted_d = NULL;
if (self->workspace_d != NULL) {
if (self->workspace_d != NULL)
FREE(self->workspace_d);
}
self->workspace_d = NULL;
if (self->di_d != NULL) {
if (self->di_d != NULL)
FREE(self->di_d);
}
self->di_d = NULL;
if (self->di_sorted_d != NULL) {
if (self->di_sorted_d != NULL)
FREE(self->di_sorted_d);
}
self->di_sorted_d = NULL;
if (self->region_flags_d != NULL) {
if (self->region_flags_d != NULL)
FREE(self->region_flags_d);
}
self->region_flags_d = NULL;
if (self->num_selected_d != NULL) {
if (self->num_selected_d != NULL)
FREE(self->num_selected_d);
}
self->num_selected_d = NULL;
if (self->forw_info_d != NULL) {
if (self->forw_info_d != NULL)
FREE(self->forw_info_d);
}
self->forw_info_d = NULL;
if (self->min_max_pixels_d != NULL) {
if (self->min_max_pixels_d != NULL)
FREE(self->min_max_pixels_d);
}
self->min_max_pixels_d = NULL;
if (self->grad_pos_d != NULL) {
if (self->grad_pos_d != NULL)
FREE(self->grad_pos_d);
}
self->grad_pos_d = NULL;
if (self->grad_col_d != NULL) {
if (self->grad_col_d != NULL)
FREE(self->grad_col_d);
}
self->grad_col_d = NULL;
if (self->grad_rad_d != NULL) {
if (self->grad_rad_d != NULL)
FREE(self->grad_rad_d);
}
self->grad_rad_d = NULL;
if (self->grad_cam_d != NULL) {
if (self->grad_cam_d != NULL)
FREE(self->grad_cam_d);
}
self->grad_cam_d = NULL;
if (self->grad_cam_buf_d != NULL) {
if (self->grad_cam_buf_d != NULL)
FREE(self->grad_cam_buf_d);
}
self->grad_cam_buf_d = NULL;
if (self->grad_opy_d != NULL) {
if (self->grad_opy_d != NULL)
FREE(self->grad_opy_d);
}
self->grad_opy_d = NULL;
if (self->n_grad_contributions_d != NULL) {
if (self->n_grad_contributions_d != NULL)
FREE(self->n_grad_contributions_d);
}
self->n_grad_contributions_d = NULL;
}

View File

@@ -41,8 +41,7 @@ GLOBAL void fill_bg(
// sphere IDs and intersection depths.
for (int i = 0; i < renderer.n_track; ++i) {
int sphere_id = -1;
renderer.forw_info_d[write_loc + 3 + i * 2] =
pulsar_int_as_float(sphere_id);
IASF(sphere_id, renderer.forw_info_d[write_loc + 3 + i * 2]);
renderer.forw_info_d[write_loc + 3 + i * 2 + 1] = -1.f;
}
if (mode == 0) {

View File

@@ -255,7 +255,7 @@ GLOBAL void calc_signature(
* for every iteration through the loading loop every thread could add a
* 'hit' to the buffer.
*/
#define RENDER_BUFFER_SIZE RENDER_BLOCK_SIZE * RENDER_BLOCK_SIZE * 2
#define RENDER_BUFFER_SIZE RENDER_BLOCK_SIZE* RENDER_BLOCK_SIZE * 2
/**
* The threshold after which the spheres that are in the render buffer
* are rendered and the buffer is flushed.

View File

@@ -64,9 +64,8 @@ GLOBAL void norm_sphere_gradients(Renderer renderer, const int num_balls) {
// The sphere only contributes to the camera gradients if it is
// large enough in screen space.
if (renderer.ids_sorted_d[idx] > 0 && ii.max.x >= ii.min.x + 3 &&
ii.max.y >= ii.min.y + 3) {
ii.max.y >= ii.min.y + 3)
renderer.ids_sorted_d[idx] = 1;
}
END_PARALLEL_NORET();
};

View File

@@ -139,9 +139,8 @@ GLOBAL void render(
coord_y < cam_norm.film_border_top + cam_norm.film_height) {
// Initialize the result.
if (mode == 0u) {
for (uint c_id = 0; c_id < cam_norm.n_channels; ++c_id) {
for (uint c_id = 0; c_id < cam_norm.n_channels; ++c_id)
result[c_id] = bg_col[c_id];
}
} else {
result[0] = 0.f;
}
@@ -191,22 +190,20 @@ GLOBAL void render(
"render|found intersection with sphere %u.\n",
sphere_id_l[write_idx]);
}
if (ii.min.x == MAX_USHORT) {
if (ii.min.x == MAX_USHORT)
// This is an invalid sphere (out of image). These spheres have
// maximum depth. Since we ordered the spheres by earliest possible
// intersection depth we re certain that there will no other sphere
// that is relevant after this one.
loading_done = true;
}
}
// Reset n_pixels_done.
n_pixels_done = 0;
thread_block.sync(); // Make sure n_loaded is updated.
if (n_loaded > RENDER_BUFFER_LOAD_THRESH) {
// The load buffer is full enough. Draw.
if (thread_block.thread_rank() == 0) {
if (thread_block.thread_rank() == 0)
n_balls_loaded += n_loaded;
}
max_closest_possible_intersection = 0.f;
// This excludes threads outside of the image boundary. Also, it reduces
// block artifacts.
@@ -293,9 +290,8 @@ GLOBAL void render(
uint warp_done = thread_warp.ballot(done);
int warp_done_bit_cnt = POPC(warp_done);
#endif //__CUDACC__ && __HIP_PLATFORM_AMD__
if (thread_warp.thread_rank() == 0) {
if (thread_warp.thread_rank() == 0)
ATOMICADD_B(&n_pixels_done, warp_done_bit_cnt);
}
// This sync is necessary to keep n_loaded until all threads are done with
// painting.
thread_block.sync();
@@ -303,9 +299,8 @@ GLOBAL void render(
}
thread_block.sync();
}
if (thread_block.thread_rank() == 0) {
if (thread_block.thread_rank() == 0)
n_balls_loaded += n_loaded;
}
PULSAR_LOG_DEV_PIX(
PULSAR_LOG_RENDER_PIX,
"render|loaded %d balls in total.\n",
@@ -391,9 +386,8 @@ GLOBAL void render(
static_cast<float>(tracker.get_n_hits());
} else {
float sm_d_normfac = FRCP(FMAX(sm_d, FEPS));
for (uint c_id = 0; c_id < cam_norm.n_channels; ++c_id) {
for (uint c_id = 0; c_id < cam_norm.n_channels; ++c_id)
result[c_id] *= sm_d_normfac;
}
int write_loc = (coord_y - cam_norm.film_border_top) * cam_norm.film_width *
(3 + 2 * n_track) +
(coord_x - cam_norm.film_border_left) * (3 + 2 * n_track);
@@ -406,7 +400,7 @@ GLOBAL void render(
IMIN(n_track, tracker.get_n_hits()));
for (int i = 0; i < n_track; ++i) {
int sphere_id = tracker.get_closest_sphere_id(i);
forw_info_d[write_loc + 3 + i * 2] = pulsar_int_as_float(sphere_id);
IASF(sphere_id, forw_info_d[write_loc + 3 + i * 2]);
forw_info_d[write_loc + 3 + i * 2 + 1] =
tracker.get_closest_sphere_depth(i) == MAX_FLOAT
? -1.f

View File

@@ -7,11 +7,13 @@
*/
#include "./renderer.h"
#include "../include/commands.h"
#include "./camera.h"
#include "./util.h"
#include <ATen/ATen.h>
#ifdef WITH_CUDA
#include <c10/cuda/CUDAException.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#endif
@@ -107,10 +109,10 @@ void Renderer::ensure_on_device(torch::Device device, bool /*non_blocking*/) {
LOG_IF(INFO, PULSAR_LOG_INIT)
<< "Transferring render buffers between devices.";
int prev_active;
C10_CUDA_CHECK(cudaGetDevice(&prev_active));
cudaGetDevice(&prev_active);
if (this->device_type == c10::DeviceType::CUDA) {
LOG_IF(INFO, PULSAR_LOG_INIT) << " Destructing on CUDA.";
C10_CUDA_CHECK(cudaSetDevice(this->device_index));
cudaSetDevice(this->device_index);
for (auto& nrend : this->renderer_vec) {
PRE::destruct<true>(&nrend);
}
@@ -122,7 +124,7 @@ void Renderer::ensure_on_device(torch::Device device, bool /*non_blocking*/) {
}
if (device.type() == c10::DeviceType::CUDA) {
LOG_IF(INFO, PULSAR_LOG_INIT) << " Constructing on CUDA.";
C10_CUDA_CHECK(cudaSetDevice(device.index()));
cudaSetDevice(device.index());
for (auto& nrend : this->renderer_vec) {
PRE::construct<true>(
&nrend,
@@ -150,7 +152,7 @@ void Renderer::ensure_on_device(torch::Device device, bool /*non_blocking*/) {
this->n_track());
}
}
C10_CUDA_CHECK(cudaSetDevice(prev_active));
cudaSetDevice(prev_active);
this->device_type = device.type();
this->device_index = device.index();
#else
@@ -744,8 +746,8 @@ std::tuple<torch::Tensor, torch::Tensor> Renderer::forward(
// moved to a CUDA device).
#ifdef WITH_CUDA
int prev_active;
C10_CUDA_CHECK(cudaGetDevice(&prev_active));
C10_CUDA_CHECK(cudaSetDevice(this->device_index));
cudaGetDevice(&prev_active);
cudaSetDevice(this->device_index);
#ifdef PULSAR_TIMINGS_BATCHED_ENABLED
START_TIME_CU(batch_forward);
#endif
@@ -790,7 +792,7 @@ std::tuple<torch::Tensor, torch::Tensor> Renderer::forward(
std::cout << "Forward render batched time per example: "
<< time_ms / static_cast<float>(batch_size) << "ms" << std::endl;
#endif
C10_CUDA_CHECK(cudaSetDevice(prev_active));
cudaSetDevice(prev_active);
#endif
} else {
#ifdef PULSAR_TIMINGS_BATCHED_ENABLED
@@ -858,9 +860,8 @@ std::tuple<torch::Tensor, torch::Tensor> Renderer::forward(
? (cudaStream_t) nullptr
#endif
: (cudaStream_t) nullptr);
if (mode == 1) {
if (mode == 1)
results[batch_i] = results[batch_i].slice(2, 0, 1, 1);
}
forw_infos[batch_i] = from_blob(
this->renderer_vec[batch_i].forw_info_d,
{this->renderer_vec[0].cam.film_height,
@@ -1101,8 +1102,8 @@ Renderer::backward(
// the renderer to a CUDA device if not built with CUDA.
#ifdef WITH_CUDA
int prev_active;
C10_CUDA_CHECK(cudaGetDevice(&prev_active));
C10_CUDA_CHECK(cudaSetDevice(this->device_index));
cudaGetDevice(&prev_active);
cudaSetDevice(this->device_index);
#ifdef PULSAR_TIMINGS_BATCHED_ENABLED
START_TIME_CU(batch_backward);
#endif
@@ -1206,7 +1207,7 @@ Renderer::backward(
at::cuda::getCurrentCUDAStream());
}
}
C10_CUDA_CHECK(cudaSetDevice(prev_active));
cudaSetDevice(prev_active);
#ifdef PULSAR_TIMINGS_BATCHED_ENABLED
STOP_TIME_CU(batch_backward);
float time_ms;

View File

@@ -128,9 +128,8 @@ struct Renderer {
stream << "pulsar::Renderer[";
// Device info.
stream << self.device_type;
if (self.device_index != -1) {
if (self.device_index != -1)
stream << ", ID " << self.device_index;
}
stream << "]";
return stream;
}

View File

@@ -8,7 +8,6 @@
#ifdef WITH_CUDA
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime_api.h>
#endif
#include <torch/extension.h>
@@ -34,13 +33,13 @@ torch::Tensor sphere_ids_from_result_info_nograd(
.contiguous();
if (forw_info.device().type() == c10::DeviceType::CUDA) {
#ifdef WITH_CUDA
C10_CUDA_CHECK(cudaMemcpyAsync(
cudaMemcpyAsync(
result.data_ptr(),
tmp.data_ptr(),
sizeof(uint32_t) * tmp.size(0) * tmp.size(1) * tmp.size(2) *
tmp.size(3),
cudaMemcpyDeviceToDevice,
at::cuda::getCurrentCUDAStream()));
at::cuda::getCurrentCUDAStream());
#else
throw std::runtime_error(
"Copy on CUDA device initiated but built "

View File

@@ -7,7 +7,6 @@
*/
#ifdef WITH_CUDA
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime_api.h>
namespace pulsar {
@@ -18,8 +17,7 @@ void cudaDevToDev(
const void* src,
const int& size,
const cudaStream_t& stream) {
C10_CUDA_CHECK(
cudaMemcpyAsync(trg, src, size, cudaMemcpyDeviceToDevice, stream));
cudaMemcpyAsync(trg, src, size, cudaMemcpyDeviceToDevice, stream);
}
void cudaDevToHost(
@@ -27,8 +25,7 @@ void cudaDevToHost(
const void* src,
const int& size,
const cudaStream_t& stream) {
C10_CUDA_CHECK(
cudaMemcpyAsync(trg, src, size, cudaMemcpyDeviceToHost, stream));
cudaMemcpyAsync(trg, src, size, cudaMemcpyDeviceToHost, stream);
}
} // namespace pytorch

View File

@@ -6,6 +6,9 @@
* LICENSE file in the root directory of this source tree.
*/
#include "./global.h"
#include "./logging.h"
/**
* A compilation unit to provide warnings about the code and avoid
* repeated messages.

View File

@@ -25,7 +25,7 @@ class BitMask {
// Use all threads in the current block to clear all bits of this BitMask
__device__ void block_clear() {
for (auto i = threadIdx.x; i < H * W * D; i += blockDim.x) {
for (int i = threadIdx.x; i < H * W * D; i += blockDim.x) {
data[i] = 0;
}
__syncthreads();

View File

@@ -23,8 +23,8 @@ __global__ void TriangleBoundingBoxKernel(
const float blur_radius,
float* bboxes, // (4, F)
bool* skip_face) { // (F,)
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
const auto num_threads = blockDim.x * gridDim.x;
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
const int num_threads = blockDim.x * gridDim.x;
const float sqrt_radius = sqrt(blur_radius);
for (int f = tid; f < F; f += num_threads) {
const float v0x = face_verts[f * 9 + 0 * 3 + 0];
@@ -56,8 +56,8 @@ __global__ void PointBoundingBoxKernel(
const int P,
float* bboxes, // (4, P)
bool* skip_points) {
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
const auto num_threads = blockDim.x * gridDim.x;
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
const int num_threads = blockDim.x * gridDim.x;
for (int p = tid; p < P; p += num_threads) {
const float x = points[p * 3 + 0];
const float y = points[p * 3 + 1];
@@ -113,7 +113,7 @@ __global__ void RasterizeCoarseCudaKernel(
const int chunks_per_batch = 1 + (E - 1) / chunk_size;
const int num_chunks = N * chunks_per_batch;
for (auto chunk = blockIdx.x; chunk < num_chunks; chunk += gridDim.x) {
for (int chunk = blockIdx.x; chunk < num_chunks; chunk += gridDim.x) {
const int batch_idx = chunk / chunks_per_batch; // batch index
const int chunk_idx = chunk % chunks_per_batch;
const int elem_chunk_start_idx = chunk_idx * chunk_size;
@@ -123,7 +123,7 @@ __global__ void RasterizeCoarseCudaKernel(
const int64_t elem_stop_idx = elem_start_idx + elems_per_batch[batch_idx];
// Have each thread handle a different face within the chunk
for (auto e = threadIdx.x; e < chunk_size; e += blockDim.x) {
for (int e = threadIdx.x; e < chunk_size; e += blockDim.x) {
const int e_idx = elem_chunk_start_idx + e;
// Check that we are still within the same element of the batch
@@ -170,7 +170,7 @@ __global__ void RasterizeCoarseCudaKernel(
// Now we have processed every elem in the current chunk. We need to
// count the number of elems in each bin so we can write the indices
// out to global memory. We have each thread handle a different bin.
for (auto byx = threadIdx.x; byx < num_bins_y * num_bins_x;
for (int byx = threadIdx.x; byx < num_bins_y * num_bins_x;
byx += blockDim.x) {
const int by = byx / num_bins_x;
const int bx = byx % num_bins_x;

View File

@@ -260,8 +260,8 @@ __global__ void RasterizeMeshesNaiveCudaKernel(
float* pix_dists,
float* bary) {
// Simple version: One thread per output pixel
auto num_threads = gridDim.x * blockDim.x;
auto tid = blockDim.x * blockIdx.x + threadIdx.x;
int num_threads = gridDim.x * blockDim.x;
int tid = blockDim.x * blockIdx.x + threadIdx.x;
for (int i = tid; i < N * H * W; i += num_threads) {
// Convert linear index to 3D index
@@ -446,8 +446,8 @@ __global__ void RasterizeMeshesBackwardCudaKernel(
// Parallelize over each pixel in images of
// size H * W, for each image in the batch of size N.
const auto num_threads = gridDim.x * blockDim.x;
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
const int num_threads = gridDim.x * blockDim.x;
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
for (int t_i = tid; t_i < N * H * W; t_i += num_threads) {
// Convert linear index to 3D index
@@ -650,8 +650,8 @@ __global__ void RasterizeMeshesFineCudaKernel(
) {
// This can be more than H * W if H or W are not divisible by bin_size.
int num_pixels = N * BH * BW * bin_size * bin_size;
auto num_threads = gridDim.x * blockDim.x;
auto tid = blockIdx.x * blockDim.x + threadIdx.x;
int num_threads = gridDim.x * blockDim.x;
int tid = blockIdx.x * blockDim.x + threadIdx.x;
for (int pid = tid; pid < num_pixels; pid += num_threads) {
// Convert linear index into bin and pixel indices. We make the within

View File

@@ -138,9 +138,6 @@ RasterizeMeshesNaive(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(face_verts);
CHECK_CPU(mesh_to_face_first_idx);
CHECK_CPU(num_faces_per_mesh);
return RasterizeMeshesNaiveCpu(
face_verts,
mesh_to_face_first_idx,
@@ -235,11 +232,6 @@ torch::Tensor RasterizeMeshesBackward(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(face_verts);
CHECK_CPU(pix_to_face);
CHECK_CPU(grad_zbuf);
CHECK_CPU(grad_bary);
CHECK_CPU(grad_dists);
return RasterizeMeshesBackwardCpu(
face_verts,
pix_to_face,
@@ -314,9 +306,6 @@ torch::Tensor RasterizeMeshesCoarse(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(face_verts);
CHECK_CPU(mesh_to_face_first_idx);
CHECK_CPU(num_faces_per_mesh);
return RasterizeMeshesCoarseCpu(
face_verts,
mesh_to_face_first_idx,
@@ -434,8 +423,6 @@ RasterizeMeshesFine(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(face_verts);
CHECK_CPU(bin_faces);
AT_ERROR("NOT IMPLEMENTED");
}
}

View File

@@ -8,6 +8,8 @@
#include <torch/extension.h>
#include <algorithm>
#include <list>
#include <queue>
#include <thread>
#include <tuple>
#include "ATen/core/TensorAccessor.h"
@@ -105,8 +107,6 @@ auto ComputeFaceAreas(const torch::Tensor& face_verts) {
return face_areas;
}
namespace {
// Helper function to use with std::find_if to find the index of any
// values in the top k struct which match a given idx.
struct IsNeighbor {
@@ -119,6 +119,7 @@ struct IsNeighbor {
int neighbor_idx;
};
namespace {
void RasterizeMeshesNaiveCpu_worker(
const int start_yi,
const int end_yi,

View File

@@ -97,8 +97,8 @@ __global__ void RasterizePointsNaiveCudaKernel(
float* zbuf, // (N, H, W, K)
float* pix_dists) { // (N, H, W, K)
// Simple version: One thread per output pixel
const auto num_threads = gridDim.x * blockDim.x;
const auto tid = blockDim.x * blockIdx.x + threadIdx.x;
const int num_threads = gridDim.x * blockDim.x;
const int tid = blockDim.x * blockIdx.x + threadIdx.x;
for (int i = tid; i < N * H * W; i += num_threads) {
// Convert linear index to 3D index
const int n = i / (H * W); // Batch index
@@ -237,8 +237,8 @@ __global__ void RasterizePointsFineCudaKernel(
float* pix_dists) { // (N, H, W, K)
// This can be more than H * W if H or W are not divisible by bin_size.
const int num_pixels = N * BH * BW * bin_size * bin_size;
const auto num_threads = gridDim.x * blockDim.x;
const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
const int num_threads = gridDim.x * blockDim.x;
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
for (int pid = tid; pid < num_pixels; pid += num_threads) {
// Convert linear index into bin and pixel indices. We make the within
@@ -376,8 +376,8 @@ __global__ void RasterizePointsBackwardCudaKernel(
float* grad_points) { // (P, 3)
// Parallelized over each of K points per pixel, for each pixel in images of
// size H * W, for each image in the batch of size N.
auto num_threads = gridDim.x * blockDim.x;
auto tid = blockIdx.x * blockDim.x + threadIdx.x;
int num_threads = gridDim.x * blockDim.x;
int tid = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = tid; i < N * H * W * K; i += num_threads) {
// const int n = i / (H * W * K); // batch index (not needed).
const int yxk = i % (H * W * K);

View File

@@ -91,10 +91,6 @@ std::tuple<torch::Tensor, torch::Tensor, torch::Tensor> RasterizePointsNaive(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(points);
CHECK_CPU(cloud_to_packed_first_idx);
CHECK_CPU(num_points_per_cloud);
CHECK_CPU(radius);
return RasterizePointsNaiveCpu(
points,
cloud_to_packed_first_idx,
@@ -170,10 +166,6 @@ torch::Tensor RasterizePointsCoarse(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(points);
CHECK_CPU(cloud_to_packed_first_idx);
CHECK_CPU(num_points_per_cloud);
CHECK_CPU(radius);
return RasterizePointsCoarseCpu(
points,
cloud_to_packed_first_idx,
@@ -240,8 +232,6 @@ std::tuple<torch::Tensor, torch::Tensor, torch::Tensor> RasterizePointsFine(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(points);
CHECK_CPU(bin_points);
AT_ERROR("NOT IMPLEMENTED");
}
}
@@ -294,10 +284,6 @@ torch::Tensor RasterizePointsBackward(
AT_ERROR("Not compiled with GPU support");
#endif
} else {
CHECK_CPU(points);
CHECK_CPU(idxs);
CHECK_CPU(grad_zbuf);
CHECK_CPU(grad_dists);
return RasterizePointsBackwardCpu(points, idxs, grad_zbuf, grad_dists);
}
}

View File

@@ -35,6 +35,8 @@ __global__ void FarthestPointSamplingKernel(
__shared__ int64_t selected_store;
// Get constants
const int64_t N = points.size(0);
const int64_t P = points.size(1);
const int64_t D = points.size(2);
// Get batch index and thread index
@@ -107,8 +109,7 @@ at::Tensor FarthestPointSamplingCuda(
const at::Tensor& points, // (N, P, 3)
const at::Tensor& lengths, // (N,)
const at::Tensor& K, // (N,)
const at::Tensor& start_idxs,
const int64_t max_K_known = -1) {
const at::Tensor& start_idxs) {
// Check inputs are on the same device
at::TensorArg p_t{points, "points", 1}, lengths_t{lengths, "lengths", 2},
k_t{K, "K", 3}, start_idxs_t{start_idxs, "start_idxs", 4};
@@ -130,12 +131,7 @@ at::Tensor FarthestPointSamplingCuda(
const int64_t N = points.size(0);
const int64_t P = points.size(1);
int64_t max_K;
if (max_K_known > 0) {
max_K = max_K_known;
} else {
max_K = at::max(K).item<int64_t>();
}
const int64_t max_K = at::max(K).item<int64_t>();
// Initialize the output tensor with the sampled indices
auto idxs = at::full({N, max_K}, -1, lengths.options());

View File

@@ -43,8 +43,7 @@ at::Tensor FarthestPointSamplingCuda(
const at::Tensor& points,
const at::Tensor& lengths,
const at::Tensor& K,
const at::Tensor& start_idxs,
const int64_t max_K_known = -1);
const at::Tensor& start_idxs);
at::Tensor FarthestPointSamplingCpu(
const at::Tensor& points,
@@ -57,23 +56,17 @@ at::Tensor FarthestPointSampling(
const at::Tensor& points,
const at::Tensor& lengths,
const at::Tensor& K,
const at::Tensor& start_idxs,
const int64_t max_K_known = -1) {
const at::Tensor& start_idxs) {
if (points.is_cuda() || lengths.is_cuda() || K.is_cuda()) {
#ifdef WITH_CUDA
CHECK_CUDA(points);
CHECK_CUDA(lengths);
CHECK_CUDA(K);
CHECK_CUDA(start_idxs);
return FarthestPointSamplingCuda(
points, lengths, K, start_idxs, max_K_known);
return FarthestPointSamplingCuda(points, lengths, K, start_idxs);
#else
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(points);
CHECK_CPU(lengths);
CHECK_CPU(K);
CHECK_CPU(start_idxs);
return FarthestPointSamplingCpu(points, lengths, K, start_idxs);
}

View File

@@ -71,8 +71,6 @@ inline void SamplePdf(
AT_ERROR("Not compiled with GPU support.");
#endif
}
CHECK_CPU(weights);
CHECK_CPU(outputs);
CHECK_CONTIGUOUS(outputs);
SamplePdfCpu(bins, weights, outputs, eps);
}

View File

@@ -99,7 +99,8 @@ namespace {
// and increment it via template recursion until it is equal to the run-time
// argument N.
template <
template <typename, int64_t> class Kernel,
template <typename, int64_t>
class Kernel,
typename T,
int64_t minN,
int64_t maxN,
@@ -123,7 +124,8 @@ struct DispatchKernelHelper1D {
// 1D dispatch: Specialization when curN == maxN
// We need this base case to avoid infinite template recursion.
template <
template <typename, int64_t> class Kernel,
template <typename, int64_t>
class Kernel,
typename T,
int64_t minN,
int64_t maxN,
@@ -143,7 +145,8 @@ struct DispatchKernelHelper1D<Kernel, T, minN, maxN, maxN, Args...> {
// the run-time values of N and M, at which point we dispatch to the run
// method of the kernel.
template <
template <typename, int64_t, int64_t> class Kernel,
template <typename, int64_t, int64_t>
class Kernel,
typename T,
int64_t minN,
int64_t maxN,
@@ -200,7 +203,8 @@ struct DispatchKernelHelper2D {
// 2D dispatch, specialization for curN == maxN
template <
template <typename, int64_t, int64_t> class Kernel,
template <typename, int64_t, int64_t>
class Kernel,
typename T,
int64_t minN,
int64_t maxN,
@@ -239,7 +243,8 @@ struct DispatchKernelHelper2D<
// 2D dispatch, specialization for curM == maxM
template <
template <typename, int64_t, int64_t> class Kernel,
template <typename, int64_t, int64_t>
class Kernel,
typename T,
int64_t minN,
int64_t maxN,
@@ -278,7 +283,8 @@ struct DispatchKernelHelper2D<
// 2D dispatch, specialization for curN == maxN, curM == maxM
template <
template <typename, int64_t, int64_t> class Kernel,
template <typename, int64_t, int64_t>
class Kernel,
typename T,
int64_t minN,
int64_t maxN,
@@ -307,7 +313,8 @@ struct DispatchKernelHelper2D<
// This is the function we expect users to call to dispatch to 1D functions
template <
template <typename, int64_t> class Kernel,
template <typename, int64_t>
class Kernel,
typename T,
int64_t minN,
int64_t maxN,
@@ -323,7 +330,8 @@ void DispatchKernel1D(const int64_t N, Args... args) {
// This is the function we expect users to call to dispatch to 2D functions
template <
template <typename, int64_t, int64_t> class Kernel,
template <typename, int64_t, int64_t>
class Kernel,
typename T,
int64_t minN,
int64_t maxN,

View File

@@ -376,6 +376,8 @@ PointLineDistanceBackward(
float tt = t_top / t_bot;
tt = __saturatef(tt);
const float2 p_proj = (1.0f - tt) * v0 + tt * v1;
const float2 d = p - p_proj;
const float dist = sqrt(dot(d, d));
const float2 grad_p = -1.0f * grad_dist * 2.0f * (p_proj - p);
const float2 grad_v0 = grad_dist * (1.0f - tt) * 2.0f * (p_proj - p);

View File

@@ -15,7 +15,3 @@
#define CHECK_CONTIGUOUS_CUDA(x) \
CHECK_CUDA(x); \
CHECK_CONTIGUOUS(x)
#define CHECK_CPU(x) \
TORCH_CHECK( \
x.device().type() == torch::kCPU, \
"Cannot use CPU implementation: " #x " not on CPU.")

View File

@@ -19,7 +19,7 @@ template <
std::is_same<T, double>::value || std::is_same<T, float>::value>>
struct vec2 {
T x, y;
using scalar_t = T;
typedef T scalar_t;
vec2(T x, T y) : x(x), y(y) {}
};

View File

@@ -18,7 +18,7 @@ template <
std::is_same<T, double>::value || std::is_same<T, float>::value>>
struct vec3 {
T x, y, z;
using scalar_t = T;
typedef T scalar_t;
vec3(T x, T y, T z) : x(x), y(y), z(z) {}
};

View File

@@ -10,20 +10,10 @@
#include <math.h>
#include <cstdio>
// Helper functions WarpReduceMin and WarpReduceMax used in .cu files.
// Starting in Volta, instructions are no longer synchronous within a warp,
// so on CUDA __syncwarp() is required between dependent shared-memory
// accesses in the unrolled tail reduction.
//
// On AMD/HIP no __syncwarp() is needed here: all wavefront lanes execute
// in lockstep (AMD has no equivalent of NVIDIA's Independent Thread
// Scheduling), and the AMDGPU memory model guarantees that LDS operations
// issued by the same wavefront are observed in program order without an
// explicit s_waitcnt — see the LLVM AMDGPU backend memory-model rules
// (https://llvm.org/docs/AMDGPUUsage.html) and the HIP hardware-
// implementation docs. This holds for both wave32 (RDNA, gfx10xx/11xx/
// 12xx) and wave64 (CDNA, gfx9xx), so the USE_ROCM skip is
// architecture-independent.
// Helper functions WarpReduceMin and WarpReduceMax used in .cu files
// Starting in Volta, instructions are no longer synchronous within a warp.
// We need to call __syncwarp() to sync the 32 threads in the warp
// instead of all the threads in the block.
template <typename scalar_t>
__device__ void
@@ -33,6 +23,8 @@ WarpReduceMin(scalar_t* min_dists, int64_t* min_idxs, const size_t tid) {
min_idxs[tid] = min_idxs[tid + 32];
min_dists[tid] = min_dists[tid + 32];
}
// AMD does not use explicit syncwarp and instead automatically inserts memory
// fences during compilation.
#if !defined(USE_ROCM)
__syncwarp();
#endif

View File

@@ -200,7 +200,6 @@ class R2N2(ShapeNetBase): # pragma: no cover
) % (shapenet_dir, ", ".join(synset_not_present))
warnings.warn(msg)
# pyrefly: ignore [bad-override]
def __getitem__(self, model_idx, view_idxs: Optional[List[int]] = None) -> Dict:
"""
Read a model by the given index.
@@ -371,7 +370,6 @@ class R2N2(ShapeNetBase): # pragma: no cover
T = RT[3, :3]
return R, T
# pyrefly: ignore [bad-override]
def render(
self,
model_ids: Optional[List[str]] = None,

View File

@@ -62,10 +62,8 @@ def collate_batched_R2N2(batch: List[Dict]): # pragma: no cover
# all models have the same number of views V, stack the batches of
# views of each model into a new batch of shape (N, V, H, W, 3).
# Otherwise leave it as a list.
# pyrefly: ignore [not-iterable]
if "images" in collated_dict:
try:
# pyrefly: ignore [unsupported-operation]
collated_dict["images"] = torch.stack(collated_dict["images"])
except RuntimeError:
print(
@@ -77,14 +75,10 @@ def collate_batched_R2N2(batch: List[Dict]): # pragma: no cover
# matrices and that all models have the same number of views V, stack each
# type of matrices into a new batch of shape (N, V, ...).
# Otherwise leave them as lists.
# pyrefly: ignore [not-iterable]
if all(x in collated_dict for x in ["R", "T", "K"]):
try:
# pyrefly: ignore [unsupported-operation]
collated_dict["R"] = torch.stack(collated_dict["R"]) # (N, V, 3, 3)
# pyrefly: ignore [unsupported-operation]
collated_dict["T"] = torch.stack(collated_dict["T"]) # (N, V, 3)
# pyrefly: ignore [unsupported-operation]
collated_dict["K"] = torch.stack(collated_dict["K"]) # (N, V, 4, 4)
except RuntimeError:
print(
@@ -95,10 +89,8 @@ def collate_batched_R2N2(batch: List[Dict]): # pragma: no cover
# If collate_batched_meshes receives voxels and all models have the same
# number of views V, stack the batches of voxels into a new batch of shape
# (N, V, S, S, S), where S is the voxel size.
# pyrefly: ignore [not-iterable]
if "voxels" in collated_dict:
try:
# pyrefly: ignore [unsupported-operation]
collated_dict["voxels"] = torch.stack(collated_dict["voxels"])
except RuntimeError:
print(
@@ -466,8 +458,6 @@ class BlenderCamera(CamerasBase): # pragma: no cover
def get_projection_transform(self, **kwargs) -> Transform3d:
transform = Transform3d(device=self.device)
# pyre-fixme[29]: `Union[(self: TensorBase, dim0: int, dim1: int) -> Tensor,
# Tensor, Module]` is not a function.
transform._matrix = self.K.transpose(1, 2).contiguous()
return transform

View File

@@ -83,7 +83,7 @@ class ShapeNetCore(ShapeNetBase): # pragma: no cover
):
synset_set.add(synset)
elif (synset in self.synset_inv.keys()) and (
path.isdir(path.join(data_dir, self.synset_inv[synset]))
(path.isdir(path.join(data_dir, self.synset_inv[synset])))
):
synset_set.add(self.synset_inv[synset])
else:

View File

@@ -52,7 +52,6 @@ class ShapeNetBase(torch.utils.data.Dataset): # pragma: no cover
"""
return len(self.model_ids)
# pyrefly: ignore [bad-override-param-name]
def __getitem__(self, idx) -> Dict:
"""
Read a model by the given index. Need to be implemented for every child class
@@ -148,17 +147,12 @@ class ShapeNetBase(torch.utils.data.Dataset): # pragma: no cover
idxs = self._handle_render_inputs(model_ids, categories, sample_nums, idxs)
# Use the getitem method which loads mesh + texture
models = [self[idx] for idx in idxs]
# pyrefly: ignore [unsupported-operation]
meshes = collate_batched_meshes(models)["mesh"]
# pyrefly: ignore [missing-attribute]
if meshes.textures is None:
# pyrefly: ignore [missing-attribute]
meshes.textures = TexturesVertex(
# pyrefly: ignore [missing-attribute]
verts_features=torch.ones_like(meshes.verts_padded(), device=device)
)
# pyrefly: ignore [missing-attribute]
meshes = meshes.to(device)
cameras = kwargs.get("cameras", FoVPerspectiveCameras()).to(device)
if len(cameras) != 1 and len(cameras) % len(meshes) != 0:

Some files were not shown because too many files have changed in this diff Show More