CUDA — Отдельные мин / макс (x, y, z) компоненты массива float3?

Question

CUDA — Отдельные мин / макс (x, y, z) компоненты массива float3?

У меня есть массив точек float3 на устройстве CUDA. Я хотел бы быстро найти минимальное и максимальное значения (отдельно) компонентов x, y и z.

Я знаю, что уже существует огромное количество реализаций, чтобы найти минимальные / максимальные значения для массивов с использованием алгоритмов редукции, но я не могу найти ни одного, который выполняет ту же операцию над многокомпонентными переменными или предоставляет возможность предоставить шаг.

Кто-нибудь знает какие-либо существующие реализации, которые будут выполнять эту операцию?

-1

c++cuda max min minmax

Решение

Другие решения

Других решений пока нет …

Источник

Accepted Answer

Я выяснил метод с использованием тяги и пользовательских компараторов,

Для тех, кто хочет сделать что-то подобное, вы можете сделать следующее. На моей машине производительность в 10 раз выше для наборов данных, превышающих 10 миллионов элементов:

// Comparators
struct comp_float3_x{
__host__ __device__
bool operator()(const float3& lhs, const float3& rhs){
return lhs.x < rhs.x;
}
};

struct comp_float3_y {
__host__ __device__
bool operator()(const float3& lhs, const float3& rhs) {
return lhs.y < rhs.y;
}
};

struct comp_float3_z {
__host__ __device__
bool operator()(const float3& lhs, const float3& rhs) {
return lhs.z < rhs.z;
}
};

void getMinMaxDeviceFloat3(Vec3i& min, Vec3i& max, const DeviceArray<float3>& points)
{
// Thrust does not deal with raw pointers well, wrapping is necessary
thrust::device_ptr<float3> ptr = thrust::device_pointer_cast(points.ptr());
thrust::pair<thrust::device_ptr<float3>, thrust::device_ptr<float3>> minmax_x = thrust::minmax_element(thrust::device, ptr, ptr + occupied_voxels, comp_float3_x());
thrust::pair<thrust::device_ptr<float3>, thrust::device_ptr<float3>> minmax_y = thrust::minmax_element(thrust::device, ptr, ptr + occupied_voxels, comp_float3_y());
thrust::pair<thrust::device_ptr<float3>, thrust::device_ptr<float3>> minmax_z = thrust::minmax_element(thrust::device, ptr, ptr + occupied_voxels, comp_float3_z());

// Host buffers
float3  min_x_host, min_y_host, min_z_host, max_x_host, max_y_host, max_z_host;

// Copy data to host
cudaMemcpy(&min_x_host, minmax_x.first.get(), sizeof(float3), cudaMemcpyDeviceToHost);
cudaMemcpy(&min_y_host, minmax_y.first.get(), sizeof(float3), cudaMemcpyDeviceToHost);
cudaMemcpy(&min_z_host, minmax_z.first.get(), sizeof(float3), cudaMemcpyDeviceToHost);
cudaMemcpy(&max_x_host, minmax_x.second.get(), sizeof(float3), cudaMemcpyDeviceToHost);
cudaMemcpy(&max_y_host, minmax_y.second.get(), sizeof(float3), cudaMemcpyDeviceToHost);
cudaMemcpy(&max_z_host, minmax_z.second.get(), sizeof(float3), cudaMemcpyDeviceToHost);

// Assign output
min[0] = (int)min_x_host.x;
min[1] = (int)min_y_host.y;
min[2] = (int)min_z_host.z;
max[0] = (int)max_x_host.x;
max[1] = (int)max_y_host.y;
max[2] = (int)max_z_host.z;
}

0