Cuda Thrust - max vec3

Question

When i want to perform a reduction on an array of float i usually do the following :

    float res = *thrust::max_element(thrust::device, 
            thrust::device_ptr<float>(dDensities),
            thrust::device_ptr<float>(dDensities+numParticles)
            );

However what i would like to do now is pretty much the same thing on a vec3 (the glm library type) array :

    float res = *thrust::max_element(thrust::device, 
            thrust::device_ptr<glm::vec3>(dDensities),
            thrust::device_ptr<glm::vec3>(dDensities+numParticles)
            );

As you can see, this has no sense because the '<' operator is not defined on. But i would like to get the maximum vec3 based on his length :

len = sqrtf(v.x*v.x + v.y*v.y + v.z*v.z);

Is that possible ?

Answer 1

Yes, its possible. You may want to read the thrust quickstart guide if you're not already familiar with it.

If you review the thrust extrema documentation, you'll note that thrust::max_element comes in several different varieties (as do most thrust algorithms). One of these accepts a binary comparison functor. We can define a comparison functor which will do what you want.

Here's a trivial worked example:

$ cat t134.cu
#include <thrust/extrema.h>
#include <thrust/device_ptr.h>
#include <glm/glm.hpp>
#include <iostream>

struct comp
{
template <typename T>
__host__ __device__
bool operator()(T &t1, T &t2){
  return ((t1.x*t1.x+t1.y*t1.y+t1.z*t1.z) < (t2.x*t2.x+t2.y*t2.y+t2.z*t2.z));
  }
};

int main(){

  int numParticles = 3;
  glm::vec3 d[numParticles];
  d[0].x = 0; d[0].y = 0; d[0].z = 0;
  d[1].x = 2; d[1].y = 2; d[1].z = 2;
  d[2].x = 1; d[2].y = 1; d[2].z = 1;

  glm::vec3 *dDensities;
  cudaMalloc(&dDensities, numParticles*sizeof(glm::vec3));
  cudaMemcpy(dDensities, d, numParticles*sizeof(glm::vec3), cudaMemcpyHostToDevice);
  glm::vec3 res = *thrust::max_element(thrust::device,
            thrust::device_ptr<glm::vec3>(dDensities),
            thrust::device_ptr<glm::vec3>(dDensities+numParticles),
            comp()
            );
  std::cout << "max element x: " << res.x << " y: " << res.y << " z: " << res.z << std::endl;
}
$ nvcc -arch=sm_61 -o t134 t134.cu
$ ./t134
max element x: 2 y: 2 z: 2
$