2 Commits

Author SHA1 Message Date
Jeremy Reizenstein
c9a23bb832 [DONOTMERGE] one-off builds for pytorch 1.13.1
do not merge
2023-01-03 14:38:32 +00:00
Jeremy Reizenstein
3388d3f0aa Windows fix for marching cubes #1398
Summary: See https://github.com/facebookresearch/pytorch3d/issues/1398 .

Reviewed By: davidsonic

Differential Revision: D42139493

fbshipit-source-id: 972fc33b9c3017554ce704f2f10190eba406b7c8
2022-12-20 04:07:04 -08:00
2 changed files with 34 additions and 421 deletions

View File

@@ -178,428 +178,45 @@ workflows:
version: 2
build_and_test:
jobs:
# - main:
# context: DOCKERHUB_TOKEN
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py38_cu102_pyt190
python_version: '3.8'
pytorch_version: 1.9.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py38_cu111_pyt190
python_version: '3.8'
pytorch_version: 1.9.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py38_cu102_pyt191
python_version: '3.8'
pytorch_version: 1.9.1
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py38_cu111_pyt191
python_version: '3.8'
pytorch_version: 1.9.1
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py38_cu102_pyt1100
python_version: '3.8'
pytorch_version: 1.10.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py38_cu111_pyt1100
python_version: '3.8'
pytorch_version: 1.10.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py38_cu113_pyt1100
python_version: '3.8'
pytorch_version: 1.10.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py38_cu102_pyt1101
python_version: '3.8'
pytorch_version: 1.10.1
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py38_cu111_pyt1101
python_version: '3.8'
pytorch_version: 1.10.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py38_cu113_pyt1101
python_version: '3.8'
pytorch_version: 1.10.1
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py38_cu102_pyt1102
python_version: '3.8'
pytorch_version: 1.10.2
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py38_cu111_pyt1102
python_version: '3.8'
pytorch_version: 1.10.2
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py38_cu113_pyt1102
python_version: '3.8'
pytorch_version: 1.10.2
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py38_cu102_pyt1110
python_version: '3.8'
pytorch_version: 1.11.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py38_cu111_pyt1110
python_version: '3.8'
pytorch_version: 1.11.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py38_cu113_pyt1110
python_version: '3.8'
pytorch_version: 1.11.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda115
context: DOCKERHUB_TOKEN
cu_version: cu115
name: linux_conda_py38_cu115_pyt1110
python_version: '3.8'
pytorch_version: 1.11.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py38_cu102_pyt1120
python_version: '3.8'
pytorch_version: 1.12.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py38_cu113_pyt1120
python_version: '3.8'
pytorch_version: 1.12.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda116
context: DOCKERHUB_TOKEN
cu_version: cu116
name: linux_conda_py38_cu116_pyt1120
name: linux_conda_py38_cu116_pyt1131
python_version: '3.8'
pytorch_version: 1.12.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py38_cu102_pyt1121
python_version: '3.8'
pytorch_version: 1.12.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py38_cu113_pyt1121
python_version: '3.8'
pytorch_version: 1.12.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda116
context: DOCKERHUB_TOKEN
cu_version: cu116
name: linux_conda_py38_cu116_pyt1121
python_version: '3.8'
pytorch_version: 1.12.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda116
context: DOCKERHUB_TOKEN
cu_version: cu116
name: linux_conda_py38_cu116_pyt1130
python_version: '3.8'
pytorch_version: 1.13.0
pytorch_version: 1.13.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda117
context: DOCKERHUB_TOKEN
cu_version: cu117
name: linux_conda_py38_cu117_pyt1130
name: linux_conda_py38_cu117_pyt1131
python_version: '3.8'
pytorch_version: 1.13.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py39_cu102_pyt190
python_version: '3.9'
pytorch_version: 1.9.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py39_cu111_pyt190
python_version: '3.9'
pytorch_version: 1.9.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py39_cu102_pyt191
python_version: '3.9'
pytorch_version: 1.9.1
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py39_cu111_pyt191
python_version: '3.9'
pytorch_version: 1.9.1
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py39_cu102_pyt1100
python_version: '3.9'
pytorch_version: 1.10.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py39_cu111_pyt1100
python_version: '3.9'
pytorch_version: 1.10.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py39_cu113_pyt1100
python_version: '3.9'
pytorch_version: 1.10.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py39_cu102_pyt1101
python_version: '3.9'
pytorch_version: 1.10.1
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py39_cu111_pyt1101
python_version: '3.9'
pytorch_version: 1.10.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py39_cu113_pyt1101
python_version: '3.9'
pytorch_version: 1.10.1
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py39_cu102_pyt1102
python_version: '3.9'
pytorch_version: 1.10.2
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py39_cu111_pyt1102
python_version: '3.9'
pytorch_version: 1.10.2
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py39_cu113_pyt1102
python_version: '3.9'
pytorch_version: 1.10.2
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py39_cu102_pyt1110
python_version: '3.9'
pytorch_version: 1.11.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py39_cu111_pyt1110
python_version: '3.9'
pytorch_version: 1.11.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py39_cu113_pyt1110
python_version: '3.9'
pytorch_version: 1.11.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda115
context: DOCKERHUB_TOKEN
cu_version: cu115
name: linux_conda_py39_cu115_pyt1110
python_version: '3.9'
pytorch_version: 1.11.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py39_cu102_pyt1120
python_version: '3.9'
pytorch_version: 1.12.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py39_cu113_pyt1120
python_version: '3.9'
pytorch_version: 1.12.0
pytorch_version: 1.13.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda116
context: DOCKERHUB_TOKEN
cu_version: cu116
name: linux_conda_py39_cu116_pyt1120
name: linux_conda_py39_cu116_pyt1131
python_version: '3.9'
pytorch_version: 1.12.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py39_cu102_pyt1121
python_version: '3.9'
pytorch_version: 1.12.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py39_cu113_pyt1121
python_version: '3.9'
pytorch_version: 1.12.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda116
context: DOCKERHUB_TOKEN
cu_version: cu116
name: linux_conda_py39_cu116_pyt1121
python_version: '3.9'
pytorch_version: 1.12.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda116
context: DOCKERHUB_TOKEN
cu_version: cu116
name: linux_conda_py39_cu116_pyt1130
python_version: '3.9'
pytorch_version: 1.13.0
pytorch_version: 1.13.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda117
context: DOCKERHUB_TOKEN
cu_version: cu117
name: linux_conda_py39_cu117_pyt1130
name: linux_conda_py39_cu117_pyt1131
python_version: '3.9'
pytorch_version: 1.13.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py310_cu102_pyt1110
python_version: '3.10'
pytorch_version: 1.11.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu111
name: linux_conda_py310_cu111_pyt1110
python_version: '3.10'
pytorch_version: 1.11.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py310_cu113_pyt1110
python_version: '3.10'
pytorch_version: 1.11.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda115
context: DOCKERHUB_TOKEN
cu_version: cu115
name: linux_conda_py310_cu115_pyt1110
python_version: '3.10'
pytorch_version: 1.11.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py310_cu102_pyt1120
python_version: '3.10'
pytorch_version: 1.12.0
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py310_cu113_pyt1120
python_version: '3.10'
pytorch_version: 1.12.0
pytorch_version: 1.13.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda116
context: DOCKERHUB_TOKEN
cu_version: cu116
name: linux_conda_py310_cu116_pyt1120
name: linux_conda_py310_cu116_pyt1131
python_version: '3.10'
pytorch_version: 1.12.0
- binary_linux_conda:
context: DOCKERHUB_TOKEN
cu_version: cu102
name: linux_conda_py310_cu102_pyt1121
python_version: '3.10'
pytorch_version: 1.12.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda113
context: DOCKERHUB_TOKEN
cu_version: cu113
name: linux_conda_py310_cu113_pyt1121
python_version: '3.10'
pytorch_version: 1.12.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda116
context: DOCKERHUB_TOKEN
cu_version: cu116
name: linux_conda_py310_cu116_pyt1121
python_version: '3.10'
pytorch_version: 1.12.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda116
context: DOCKERHUB_TOKEN
cu_version: cu116
name: linux_conda_py310_cu116_pyt1130
python_version: '3.10'
pytorch_version: 1.13.0
pytorch_version: 1.13.1
- binary_linux_conda:
conda_docker_image: pytorch/conda-builder:cuda117
context: DOCKERHUB_TOKEN
cu_version: cu117
name: linux_conda_py310_cu117_pyt1130
name: linux_conda_py310_cu117_pyt1131
python_version: '3.10'
pytorch_version: 1.13.0
- binary_linux_conda_cuda:
name: testrun_conda_cuda_py38_cu102_pyt190
context: DOCKERHUB_TOKEN
python_version: "3.8"
pytorch_version: '1.9.0'
cu_version: "cu102"
- binary_macos_wheel:
cu_version: cpu
name: macos_wheel_py3.8_cpu
python_version: '3.8'
pytorch_version: '1.13.0'
- binary_macos_wheel:
cu_version: cpu
name: macos_wheel_py3.9_cpu
python_version: '3.9'
pytorch_version: '1.13.0'
- binary_macos_wheel:
cu_version: cpu
name: macos_wheel_py3.10_cpu
python_version: '3.10'
pytorch_version: '1.13.0'
pytorch_version: 1.13.1

View File

@@ -13,7 +13,6 @@
#include <thrust/scan.h>
#include <cstdio>
#include "marching_cubes/tables.h"
#include "utils/pytorch3d_cutils.h"
/*
Parallelized marching cubes for pytorch extension
@@ -267,13 +266,12 @@ __global__ void CompactVoxelsKernel(
// isolevel: threshold to determine isosurface intersection
//
__global__ void GenerateFacesKernel(
torch::PackedTensorAccessor32<float, 2, torch::RestrictPtrTraits> verts,
torch::PackedTensorAccessor<int64_t, 2, torch::RestrictPtrTraits> faces,
torch::PackedTensorAccessor<int64_t, 1, torch::RestrictPtrTraits> ids,
torch::PackedTensorAccessor32<int, 1, torch::RestrictPtrTraits>
at::PackedTensorAccessor32<float, 2, at::RestrictPtrTraits> verts,
at::PackedTensorAccessor<int64_t, 2, at::RestrictPtrTraits> faces,
at::PackedTensorAccessor<int64_t, 1, at::RestrictPtrTraits> ids,
at::PackedTensorAccessor32<int, 1, at::RestrictPtrTraits>
compactedVoxelArray,
torch::PackedTensorAccessor32<int, 1, torch::RestrictPtrTraits>
numVertsScanned,
at::PackedTensorAccessor32<int, 1, at::RestrictPtrTraits> numVertsScanned,
const uint activeVoxels,
const at::PackedTensorAccessor32<float, 3, at::RestrictPtrTraits> vol,
const at::PackedTensorAccessor32<int, 2, at::RestrictPtrTraits> faceTable,
@@ -436,15 +434,15 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> MarchingCubesCuda(
cudaStream_t stream = at::cuda::getCurrentCUDAStream();
// transfer _FACE_TABLE data to device
torch::Tensor face_table_tensor = torch::zeros(
{256, 16}, torch::TensorOptions().dtype(at::kInt).device(at::kCPU));
at::Tensor face_table_tensor = at::zeros(
{256, 16}, at::TensorOptions().dtype(at::kInt).device(at::kCPU));
auto face_table_a = face_table_tensor.accessor<int, 2>();
for (int i = 0; i < 256; i++) {
for (int j = 0; j < 16; j++) {
face_table_a[i][j] = _FACE_TABLE[i][j];
}
}
torch::Tensor faceTable = face_table_tensor.to(vol.device());
at::Tensor faceTable = face_table_tensor.to(vol.device());
// get numVoxels
int threads = 128;
@@ -458,10 +456,10 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> MarchingCubesCuda(
}
auto d_voxelVerts =
torch::zeros({numVoxels}, torch::TensorOptions().dtype(at::kInt))
at::zeros({numVoxels}, at::TensorOptions().dtype(at::kInt))
.to(vol.device());
auto d_voxelOccupied =
torch::zeros({numVoxels}, torch::TensorOptions().dtype(at::kInt))
at::zeros({numVoxels}, at::TensorOptions().dtype(at::kInt))
.to(vol.device());
// Execute "ClassifyVoxelKernel" kernel to precompute
@@ -480,7 +478,7 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> MarchingCubesCuda(
// If the number of active voxels is 0, return zero tensor for verts and
// faces.
auto d_voxelOccupiedScan =
torch::zeros({numVoxels}, torch::TensorOptions().dtype(at::kInt))
at::zeros({numVoxels}, at::TensorOptions().dtype(at::kInt))
.to(vol.device());
ThrustScanWrapper(
d_voxelOccupiedScan.data_ptr<int>(),
@@ -493,23 +491,21 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> MarchingCubesCuda(
int activeVoxels = lastElement + lastScan;
const int device_id = vol.device().index();
auto opt =
torch::TensorOptions().dtype(torch::kInt).device(torch::kCUDA, device_id);
auto opt_long = torch::TensorOptions()
.dtype(torch::kInt64)
.device(torch::kCUDA, device_id);
auto opt = at::TensorOptions().dtype(at::kInt).device(at::kCUDA, device_id);
auto opt_long =
at::TensorOptions().dtype(at::kLong).device(at::kCUDA, device_id);
if (activeVoxels == 0) {
int ntris = 0;
torch::Tensor verts = torch::zeros({ntris * 3, 3}, vol.options());
torch::Tensor faces = torch::zeros({ntris, 3}, opt_long);
torch::Tensor ids = torch::zeros({ntris}, opt_long);
at::Tensor verts = at::zeros({ntris * 3, 3}, vol.options());
at::Tensor faces = at::zeros({ntris, 3}, opt_long);
at::Tensor ids = at::zeros({ntris}, opt_long);
return std::make_tuple(verts, faces, ids);
}
// Execute "CompactVoxelsKernel" kernel to compress voxels for accleration.
// This allows us to run triangle generation on only the occupied voxels.
auto d_compVoxelArray = torch::zeros({activeVoxels}, opt);
auto d_compVoxelArray = at::zeros({activeVoxels}, opt);
CompactVoxelsKernel<<<grid, threads, 0, stream>>>(
d_compVoxelArray.packed_accessor32<int, 1, at::RestrictPtrTraits>(),
d_voxelOccupied.packed_accessor32<int, 1, at::RestrictPtrTraits>(),
@@ -519,7 +515,7 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> MarchingCubesCuda(
cudaDeviceSynchronize();
// Scan d_voxelVerts array to generate offsets of vertices for each voxel
auto d_voxelVertsScan = torch::zeros({numVoxels}, opt);
auto d_voxelVertsScan = at::zeros({numVoxels}, opt);
ThrustScanWrapper(
d_voxelVertsScan.data_ptr<int>(),
d_voxelVerts.data_ptr<int>(),
@@ -533,10 +529,10 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> MarchingCubesCuda(
// Execute "GenerateFacesKernel" kernel
// This runs only on the occupied voxels.
// It looks up the field values and generates the triangle data.
torch::Tensor verts = torch::zeros({totalVerts, 3}, vol.options());
torch::Tensor faces = torch::zeros({totalVerts / 3, 3}, opt_long);
at::Tensor verts = at::zeros({totalVerts, 3}, vol.options());
at::Tensor faces = at::zeros({totalVerts / 3, 3}, opt_long);
torch::Tensor ids = torch::zeros({totalVerts}, opt_long);
at::Tensor ids = at::zeros({totalVerts}, opt_long);
dim3 grid2((activeVoxels + threads - 1) / threads, 1, 1);
if (grid2.x > 65535) {