I'm trying to transform a float3 array by using containers into a container of a specific structure. Code below:
#include <thrust/device_ptr.h>
#include <thrust/extrema.h>
#include <thrust/reduce.h>
#include <thrust/execution_policy.h>
#include <thrust/functional.h>
// Ouput structure for thrust::transform
struct cloud_point_index_idx {
unsigned int idx;
float3 cloud_point;
cloud_point_index_idx(unsigned int idx_, float3 cloud_point_) :
idx(idx_), cloud_point(cloud_point_) {}
bool operator < (const cloud_point_index_idx &p) const {
return (idx < p.idx);
}
};
// Functor for thrust::transform
struct point2voxelcentroid {
float3 leaf_size;
int min_x, min_y, min_z;
point2voxelcentroid(float3 leaf_size,int min_x, int min_y, int min_z) {
this->leaf_size = leaf_size;
this->min_x = min_x; this->min_y = min_y; this->min_z = min_z;
}
__host__ __device__
cloud_point_index_idx operator()(const float3 point) const {
int ijk0 = static_cast<int>(floor(point.x / leaf_size.x) -
static_cast<float>(min_x));
int ijk1 = static_cast<int>(floor(point.y / leaf_size.y) -
static_cast<float>(min_y));
int ijk2 = static_cast<int>(floor(point.z / leaf_size.z) -
static_cast<float>(min_z));
int voxel_idx = ijk0 + ijk1 + ijk2;
return cloud_point_index_idx(static_cast<unsigned int>(voxel_idx), point);
}
};
int main() { // Example
int num_points = 5;
float3 data[5] = {{1, 0, 2}, {2, 1, 3}, {1, 1, -5}, {-1, 3, -2}, {-5, -2, 0}}; // Set the data
int min_b_[3] = {-5, -2, -5};
float3 leaf_size = {0.5, 0.5, 0.5};
thrust::device_vector<float3> d_ptr(data, data + num_points); // Wrap it into a device_vector
thrust::device_vector<cloud_point_index_idx> voxel_idx_vector; // Output
voxel_idx_vector.reserve(num_points);
thrust::transform(
thrust::device,
d_ptr.begin(), d_ptr.end(),
voxel_idx_vector.begin(),
point2voxelcentroid(leaf_size, min_b_[0], min_b_[1], min_b_[2]));
thrust::host_vector<cloud_point_index_idx> indices; // Host vector to verify
indices.reserve(num_points);
thrust::copy(voxel_idx_vector.begin(), voxel_idx_vector.end(), indices.begin()); // Copy device to host
// Print out values
std::cout << "\n---\nAfter assignment\n";
for (int i = 0; i < num_points; i++) {
std::cout << "Voxel idx: " << indices[i].idx << ". Point: [" << indices[i].cloud_point.x << ", "
<< indices[i].cloud_point.y << ", " << indices[i].cloud_point.z << std::endl;
}
}
I inspected my functor values and they seem to correctly parse the data but when I print my host_vector I get really weird values, nothing related to what my output must be. I suspect I'm not initializing my output host/device vectors correctly. I tried other methods to initialize them but they all give me errors when compiling. I'm not sure what I'm doing wrong.
There are a couple of problems here, but the most severe is the use of reserve which doesn't actually allocate memory for a thrust container.
What you need to do is define a default constructor and explicitly allocate a size at instantiation. Something like this:
struct cloud_point_index_idx {
int idx;
float3 cloud_point;
cloud_point_index_idx()=default;
__host__ __device__
cloud_point_index_idx(unsigned int idx_, float3 cloud_point_) :
idx(idx_), cloud_point(cloud_point_) {}
__host__ __device__
bool operator < (const cloud_point_index_idx &p) const {
return (idx < p.idx);
}
};
(requires -std=c++11) will define a default constructor on both device and host which the container must call during initialization of each class instance.
This modification of your code works for me:
$ cat bodgey.cu
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/extrema.h>
#include <thrust/reduce.h>
#include <thrust/execution_policy.h>
#include <thrust/functional.h>
#include <iostream>
// Ouput structure for thrust::transform
struct cloud_point_index_idx {
int idx;
float3 cloud_point;
cloud_point_index_idx()=default;
__host__ __device__
cloud_point_index_idx(unsigned int idx_, float3 cloud_point_) :
idx(idx_), cloud_point(cloud_point_) {}
__host__ __device__
bool operator < (const cloud_point_index_idx &p) const {
return (idx < p.idx);
}
};
// Functor for thrust::transform
struct point2voxelcentroid {
float3 leaf_size;
int min_x, min_y, min_z;
point2voxelcentroid(float3 leaf_size,int min_x, int min_y, int min_z) {
this->leaf_size = leaf_size;
this->min_x = min_x; this->min_y = min_y; this->min_z = min_z;
}
__host__ __device__
cloud_point_index_idx operator()(const float3 point) const {
int ijk0 = static_cast<int>(floor(point.x / leaf_size.x) -
static_cast<float>(min_x));
int ijk1 = static_cast<int>(floor(point.y / leaf_size.y) -
static_cast<float>(min_y));
int ijk2 = static_cast<int>(floor(point.z / leaf_size.z) -
static_cast<float>(min_z));
int voxel_idx = ijk0 + ijk1 + ijk2;
return cloud_point_index_idx(voxel_idx, point);
}
};
int main() { // Example
int num_points = 5;
float3 data[5] = {{1, 0, 2}, {2, 1, 3}, {1, 1, -5}, {-1, 3, -2}, {-5, -2, 0}}; // Set the data
int min_b_[3] = {-5, -2, -5};
float3 leaf_size = {0.5, 0.5, 0.5};
thrust::device_vector<float3> d_ptr(data, data + num_points); // Wrap it into a device_vector
thrust::device_vector<cloud_point_index_idx> voxel_idx_vector(num_points); // Output
thrust::transform(
thrust::device,
d_ptr.begin(), d_ptr.end(),
voxel_idx_vector.begin(),
point2voxelcentroid(leaf_size, min_b_[0], min_b_[1], min_b_[2]));
thrust::host_vector<cloud_point_index_idx> indices(num_points); // Host vector to verify
thrust::copy(voxel_idx_vector.begin(), voxel_idx_vector.end(), indices.begin()); // Copy device to host
// Print out values
std::cout << "\n---\nAfter assignment\n";
for (int i = 0; i < num_points; i++) {
std::cout << "Voxel idx: " << indices[i].idx << ". Point: [" << indices[i].cloud_point.x << ", "
<< indices[i].cloud_point.y << ", " << indices[i].cloud_point.z << std::endl;
}
}
$ nvcc -std=c++11 -arch=sm_52 -o bodgey bodgey.cu
$ ./bodgey
---
After assignment
Voxel idx: 18. Point: [1, 0, 2
Voxel idx: 24. Point: [2, 1, 3
Voxel idx: 6. Point: [1, 1, -5
Voxel idx: 12. Point: [-1, 3, -2
Voxel idx: -2. Point: [-5, -2, 0