Thrust::transform with lambda not working

[rafael.o.scatena]

Hello,

I am trying to use thrust::transform to select a struct variable from a thrust::vectror and copy it to another thurst vector. But it is not copying.

Thanks,
Rafael Scatena

Those are mine results:

First 5 elements of h_triangles:
h_p1x[0] = 16.6072
h_p1x[1] = 16.6072
h_p1x[2] = 16.6072
h_p1x[3] = 16.6951
h_p1x[4] = 16.6951
h_p1x[5] = 16.783
h_p1x[6] = 16.6072
h_p1x[7] = 16.5633
h_p1x[8] = 16.6951
h_p1x[9] = 16.6951
h_p1x[10] = 16.783
h_p1x[11] = 16.783
h_p1x[12] = 16.8269
h_p1x[13] = 16.783
h_p1x[14] = 16.6072
h_p1x[15] = 16.5633
h_p1x[16] = 16.6951
h_p1x[17] = 16.6951
h_p1x[18] = 16.783
h_p1x[19] = 16.783

First 5 elements of d_p1x:
h_p1x[0] = 0
h_p1x[1] = 0
h_p1x[2] = 0
h_p1x[3] = 0
h_p1x[4] = 0
h_p1x[5] = 0
h_p1x[6] = 0
h_p1x[7] = 0
h_p1x[8] = 0
h_p1x[9] = 0
h_p1x[10] = 0
h_p1x[11] = 0
h_p1x[12] = 0
h_p1x[13] = 0
h_p1x[14] = 0
h_p1x[15] = 0
h_p1x[16] = 0
h_p1x[17] = 0
h_p1x[18] = 0
h_p1x[19] = 0

The struct:

  struct Triangle{ double p1x; double p1y;  
  double p1z;double p2x; double p2y;  double p2z;
                 double p3x; double p3y;  double p3z;
                 
                 int Material;	int Body;
	             int64_t  triangle_id;
	             
	            __host__ __device__
                Triangle() {}

                __host__ __device__
                Triangle(double p1x, double p1y, double p1z,double p2x, double p2y, double p2z,
                         double p3x, double p3y, double p3z,int material, int body, int64_t id)
        : p1x(p1x), p1y(p1y), p1z(p1z),p2x(p2x), p2y(p2y), p2z(p2z),
          p3x(p3x), p3y(p3y), p3z(p3z),Material(material), Body(body), triangle_id(id) {}
};

The lambda

struct ExtractP1x {__host__ __device__  double operator()(const Triangle& t)  const {return t.p1x;}};

The code:

        thrust::device_vector<Triangle> d_triangles(triangles.begin(), triangles.end());


        // Create thrust device vectors to hold the extracted data
        thrust::device_vector<double>  d_p1x(h_numTriangles), d_p1y(h_numTriangles), d_p1z(h_numTriangles);
        thrust::device_vector<double>  d_p2x(h_numTriangles), d_p2y(h_numTriangles), d_p2z(h_numTriangles);
        thrust::device_vector<double>  d_p3x(h_numTriangles), d_p3y(h_numTriangles), d_p3z(h_numTriangles);
        thrust::device_vector<int>     d_Material(h_numTriangles), d_Body(h_numTriangles);
        thrust::device_vector<int64_t> d_triangle_id(h_numTriangles);
  
  
         thrust::host_vector<Triangle> h_triangles = d_triangles;  // Copy data from device to host

        // Print some elements from the host vector (say the first 5)
        std::cout << "First 5 elements of h_triangles:" << std::endl;
        for (int i = 0; i < 20 && i < h_numTriangles; i++) {std::cout << "h_p1x[" << i << "] = " << h_triangles[i].p1x << std::endl;}

  
 
        // Copy data from host to device using thrust::transform
        thrust::transform(d_triangles.begin(), d_triangles.end(), d_p1x.begin(), ExtractP1x());

// Print some elements from the host vector (say the first 5)
      std::cout << "First 5 elements of d_p1x:" << std::endl;
      for (int i = 0; i < 20 && i < h_numTriangles; i++) {
          std::cout << "h_p1x[" << i << "] = " << h_p1x[i] << std::endl;
  }

[Robert_Crovella]

your comments in the code don’t make sense to me.

Putting that aside, the thrust selection operation appears to be here:

@rafael.o.scatena

        thrust::transform(d_triangles.begin(), d_triangles.end(), d_p1x.begin(), ExtractP1x());

After that, the things you are printing out:

@rafael.o.scatena

          std::cout << "h_p1x[" << i << "] = " << h_p1x[i] << std::endl;

are from a vector that doesn’t have any contents from the previous select operation. Perhaps you meant to copy the d_p1x vector to the h_p1x vector?

[rafael.o.scatena]

I didn`t post the full code. Here it is:

    thrust::device_vector<Triangle> d_triangles(triangles.begin(), triangles.end());


    // Create thrust device vectors to hold the extracted data
    thrust::device_vector<double>  d_p1x(h_numTriangles);thrust::device_vector<double>  d_p1y(h_numTriangles), d_p1z(h_numTriangles);
    thrust::device_vector<double>  d_p2x(h_numTriangles), d_p2y(h_numTriangles), d_p2z(h_numTriangles);
    thrust::device_vector<double>  d_p3x(h_numTriangles), d_p3y(h_numTriangles), d_p3z(h_numTriangles);
    thrust::device_vector<int>     d_Material(h_numTriangles), d_Body(h_numTriangles);
    thrust::device_vector<int64_t> d_triangle_id(h_numTriangles);


     thrust::host_vector<Triangle> h_triangles = d_triangles;  // Copy data from device to host

      // Print some elements from the host vector (say the first 5)
      std::cout << "First 5 elements of h_triangles:" << std::endl;
      for (int i = 0; i < 20 && i < h_numTriangles; i++) {std::cout << "h_p1x[" << i << "] = " << h_triangles[i].p2x << std::endl;}


      std::cout <<  "PRISMAMESH: Thrust Init" << std::endl;

    

    // Copy data from host to device using thrust::transform
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p1x.begin(), ExtractP1x());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p1y.begin(), ExtractP1y());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p1z.begin(), ExtractP1z());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p2x.begin(), ExtractP2x());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p2y.begin(), ExtractP2y());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p2z.begin(), ExtractP2z());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p3x.begin(), ExtractP3x());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p3y.begin(), ExtractP3y());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p3z.begin(), ExtractP3z());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_Material.begin(), ExtractMaterial());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_Body.begin(), ExtractBody());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_triangle_id.begin(), ExtractTriangleID());

      std::cout <<  "PRISMAMESH: Thrust Extract" << std::endl;
      
   //     size_t freeMem, totalMem;
   //     cudaMemGetInfo(&freeMem, &totalMem);
 //       printf("Free Memory: %zu, Total Memory: %zu\n", freeMem, totalMem);


    //  thrust::host_vector<double> h_p1x = d_p1x;  // Copy data from device to host
      

    // Print some elements from the host vector (say the first 5)
    std::cout << "First 5 elements of d_p1x:" << std::endl;
    for (int i = 0; i < 20 && i < h_numTriangles; i++) {std::cout << "d_p1x[" << i << "] = " << d_p1x[i] << std::endl;}

But I still get the same result:

First 5 elements of h_triangles:
h_p1x[0] = 16.6072
h_p1x[1] = 16.6951
h_p1x[2] = 16.6072
h_p1x[3] = 16.783
h_p1x[4] = 16.6951
h_p1x[5] = 16.783
h_p1x[6] = 16.6072
h_p1x[7] = 16.6072
h_p1x[8] = 16.6072
h_p1x[9] = 16.6072
h_p1x[10] = 16.6951
h_p1x[11] = 16.6951
h_p1x[12] = 16.8269
h_p1x[13] = 16.8269
h_p1x[14] = 16.6072
h_p1x[15] = 16.6072
h_p1x[16] = 16.6072
h_p1x[17] = 16.6072
h_p1x[18] = 16.6951
h_p1x[19] = 16.6951
PRISMAMESH: Thrust Init
PRISMAMESH: Thrust Extract
First 5 elements of d_p1x:
d_p1x[0] = 0
d_p1x[1] = 0
d_p1x[2] = 0
d_p1x[3] = 0
d_p1x[4] = 0
d_p1x[5] = 0
d_p1x[6] = 0
d_p1x[7] = 0
d_p1x[8] = 0
d_p1x[9] = 0
d_p1x[10] = 0
d_p1x[11] = 0
d_p1x[12] = 0
d_p1x[13] = 0
d_p1x[14] = 0
d_p1x[15] = 0
d_p1x[16] = 0
d_p1x[17] = 0
d_p1x[18] = 0
d_p1x[19] = 0

[rafael.o.scatena]

The structs are these:

struct ExtractP1x {__host__ __device__  double operator()(const Triangle& t)  const {return t.p1x;}};
struct ExtractP1y {__host__ __device__  double operator()(const Triangle& t) const {return t.p1y;}};
struct ExtractP1z {__host__ __device__  double operator()(const  Triangle& t)  const  {return t.p1z;}};
struct ExtractP2x {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p2x;}};
struct ExtractP2y {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p2y;}};
struct ExtractP2z {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p2z;}};
struct ExtractP3x {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p3x;}};
struct ExtractP3y {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p3y;}};
struct ExtractP3z {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p3z;}};
struct ExtractMaterial {__host__ __device__ int operator()(const  Triangle& t)const  {return t.Material;}};
struct ExtractBody {__host__ __device__ int operator()(const  Triangle& t)const  {return t.Body;}};
struct ExtractTriangleID {__host__ __device__   int64_t operator()(const Triangle& t) const  {return t.triangle_id;}};

[Robert_Crovella]

When I build a complete code around what you have shown, I don’t seem to have any trouble with the “selection” operations:

# cat t285.cu
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/execution_policy.h>
#include <cstdint>
#include <iostream>


  struct Triangle{ double p1x; double p1y;
  double p1z;double p2x; double p2y;  double p2z;
                 double p3x; double p3y;  double p3z;

                 int Material;  int Body;
                     int64_t  triangle_id;

                    __host__ __device__
                Triangle() {}

                __host__ __device__
                Triangle(double p1x, double p1y, double p1z,double p2x, double p2y, double p2z,
                         double p3x, double p3y, double p3z,int material, int body, int64_t id)
        : p1x(p1x), p1y(p1y), p1z(p1z),p2x(p2x), p2y(p2y), p2z(p2z),
          p3x(p3x), p3y(p3y), p3z(p3z),Material(material), Body(body), triangle_id(id) {}
};

struct ExtractP1x {__host__ __device__  double operator()(const Triangle& t)  const {return t.p1x;}};
struct ExtractP1y {__host__ __device__  double operator()(const Triangle& t) const {return t.p1y;}};
struct ExtractP1z {__host__ __device__  double operator()(const  Triangle& t)  const  {return t.p1z;}};
struct ExtractP2x {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p2x;}};
struct ExtractP2y {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p2y;}};
struct ExtractP2z {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p2z;}};
struct ExtractP3x {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p3x;}};
struct ExtractP3y {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p3y;}};
struct ExtractP3z {__host__ __device__  double operator()(const  Triangle& t) const  {return t.p3z;}};
struct ExtractMaterial {__host__ __device__ int operator()(const  Triangle& t)const  {return t.Material;}};
struct ExtractBody {__host__ __device__ int operator()(const  Triangle& t)const  {return t.Body;}};
struct ExtractTriangleID {__host__ __device__   int64_t operator()(const Triangle& t) const  {return t.triangle_id;}};


int main(){
    int h_numTriangles = 256;
    struct Triangle my_init;
    my_init.p1x = 16.7;
    my_init.p2x = 16.5;
    thrust::host_vector<Triangle> triangles(h_numTriangles, my_init);
    thrust::device_vector<Triangle> d_triangles(triangles.begin(), triangles.end());


    // Create thrust device vectors to hold the extracted data
    thrust::device_vector<double>  d_p1x(h_numTriangles);thrust::device_vector<double>  d_p1y(h_numTriangles), d_p1z(h_numTriangles);
    thrust::device_vector<double>  d_p2x(h_numTriangles), d_p2y(h_numTriangles), d_p2z(h_numTriangles);
    thrust::device_vector<double>  d_p3x(h_numTriangles), d_p3y(h_numTriangles), d_p3z(h_numTriangles);
    thrust::device_vector<int>     d_Material(h_numTriangles), d_Body(h_numTriangles);
    thrust::device_vector<int64_t> d_triangle_id(h_numTriangles);


     thrust::host_vector<Triangle> h_triangles = d_triangles;  // Copy data from device to host

      // Print some elements from the host vector (say the first 5)
      std::cout << "First 5 elements of h_triangles:" << std::endl;
      for (int i = 0; i < 20 && i < h_numTriangles; i++) {std::cout << "h_p1x[" << i << "] = " << h_triangles[i].p2x << std::endl;}


      std::cout <<  "PRISMAMESH: Thrust Init" << std::endl;



    // Copy data from host to device using thrust::transform
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p1x.begin(), ExtractP1x());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p1y.begin(), ExtractP1y());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p1z.begin(), ExtractP1z());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p2x.begin(), ExtractP2x());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p2y.begin(), ExtractP2y());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p2z.begin(), ExtractP2z());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p3x.begin(), ExtractP3x());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p3y.begin(), ExtractP3y());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_p3z.begin(), ExtractP3z());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_Material.begin(), ExtractMaterial());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_Body.begin(), ExtractBody());
    thrust::transform(thrust::device,d_triangles.begin(), d_triangles.end(), d_triangle_id.begin(), ExtractTriangleID());

      std::cout <<  "PRISMAMESH: Thrust Extract" << std::endl;

   //     size_t freeMem, totalMem;
   //     cudaMemGetInfo(&freeMem, &totalMem);
 //       printf("Free Memory: %zu, Total Memory: %zu\n", freeMem, totalMem);


    //  thrust::host_vector<double> h_p1x = d_p1x;  // Copy data from device to host


    // Print some elements from the host vector (say the first 5)
    std::cout << "First 5 elements of d_p1x:" << std::endl;
    for (int i = 0; i < 20 && i < h_numTriangles; i++) {std::cout << "d_p1x[" << i << "] = " << d_p1x[i] << std::endl;}
}
# nvcc -o t285 t285.cu -std=c++17
# ./t285
First 5 elements of h_triangles:
h_p1x[0] = 16.5
h_p1x[1] = 16.5
h_p1x[2] = 16.5
h_p1x[3] = 16.5
h_p1x[4] = 16.5
h_p1x[5] = 16.5
h_p1x[6] = 16.5
h_p1x[7] = 16.5
h_p1x[8] = 16.5
h_p1x[9] = 16.5
h_p1x[10] = 16.5
h_p1x[11] = 16.5
h_p1x[12] = 16.5
h_p1x[13] = 16.5
h_p1x[14] = 16.5
h_p1x[15] = 16.5
h_p1x[16] = 16.5
h_p1x[17] = 16.5
h_p1x[18] = 16.5
h_p1x[19] = 16.5
PRISMAMESH: Thrust Init
PRISMAMESH: Thrust Extract
First 5 elements of d_p1x:
d_p1x[0] = 16.7
d_p1x[1] = 16.7
d_p1x[2] = 16.7
d_p1x[3] = 16.7
d_p1x[4] = 16.7
d_p1x[5] = 16.7
d_p1x[6] = 16.7
d_p1x[7] = 16.7
d_p1x[8] = 16.7
d_p1x[9] = 16.7
d_p1x[10] = 16.7
d_p1x[11] = 16.7
d_p1x[12] = 16.7
d_p1x[13] = 16.7
d_p1x[14] = 16.7
d_p1x[15] = 16.7
d_p1x[16] = 16.7
d_p1x[17] = 16.7
d_p1x[18] = 16.7
d_p1x[19] = 16.7
#

CUDA 12.2, L4 GPU

Your code is printing out p2x but calling it p1x in the printout, I haven’t tried to fix any of that.

I probably won’t be able to help further unless you provide a complete code, just as I have done. If you continue to post snippets, I’m unlikely to be able to comment further.