This repository has been archived on 2023-08-20. You can view files and clone it, but cannot push or open issues or pull requests.
yap-6.3/packages/cuda/union2.cu
2013-11-03 14:13:08 +00:00

453 lines
7.6 KiB
Plaintext

#include <thrust/device_vector.h>
#include <thrust/unique.h>
#include <thrust/distance.h>
#include <iostream>
typedef struct n2
{
int v[2];
}s2;
typedef struct n3
{
int v[3];
}s3;
typedef struct n4
{
int v[4];
}s4;
typedef struct n5
{
int v[5];
}s5;
typedef struct n6
{
int v[6];
}s6;
typedef struct n7
{
int v[7];
}s7;
struct p2
{
__host__ __device__
bool operator()(const s2 &r1, const s2 &r2)
{
int x;
for(x = 0; x < 2; x++)
{
if(r1.v[x] != r2.v[x])
return false;
}
return true;
}
};
struct o2
{
__host__ __device__
bool operator()(const s2 &r1, const s2 &r2)
{
int x;
for(x = 0; x < 2; x++)
{
if(r1.v[x] > r2.v[x])
return true;
if(r1.v[x] < r2.v[x])
return false;
}
return false;
}
};
struct p3
{
__host__ __device__
bool operator()(const s3 &r1, const s3 &r2)
{
int x;
for(x = 0; x < 3; x++)
{
if(r1.v[x] != r2.v[x])
return false;
}
return true;
}
};
struct o3
{
__host__ __device__
bool operator()(const s3 &r1, const s3 &r2)
{
int x;
for(x = 0; x < 3; x++)
{
if(r1.v[x] > r2.v[x])
return true;
if(r1.v[x] < r2.v[x])
return false;
}
return false;
}
};
struct p4
{
__host__ __device__
bool operator()(const s4 &r1, const s4 &r2)
{
int x;
for(x = 0; x < 4; x++)
{
if(r1.v[x] != r2.v[x])
return false;
}
return true;
}
};
struct o4
{
__host__ __device__
bool operator()(const s4 &r1, const s4 &r2)
{
int x;
for(x = 0; x < 4; x++)
{
if(r1.v[x] > r2.v[x])
return true;
if(r1.v[x] < r2.v[x])
return false;
}
return false;
}
};
struct p5
{
__host__ __device__
bool operator()(const s5 &r1, const s5 &r2)
{
int x;
for(x = 0; x < 5; x++)
{
if(r1.v[x] != r2.v[x])
return false;
}
return true;
}
};
struct o5
{
__host__ __device__
bool operator()(const s5 &r1, const s5 &r2)
{
int x;
for(x = 0; x < 5; x++)
{
if(r1.v[x] > r2.v[x])
return true;
if(r1.v[x] < r2.v[x])
return false;
}
return false;
}
};
struct p6
{
__host__ __device__
bool operator()(const s6 &r1, const s6 &r2)
{
int x;
for(x = 0; x < 6; x++)
{
if(r1.v[x] != r2.v[x])
return false;
}
return true;
}
};
struct o6
{
__host__ __device__
bool operator()(const s6 &r1, const s6 &r2)
{
int x;
for(x = 0; x < 6; x++)
{
if(r1.v[x] > r2.v[x])
return true;
if(r1.v[x] < r2.v[x])
return false;
}
return false;
}
};
struct p7
{
__host__ __device__
bool operator()(const s7 &r1, const s7 &r2)
{
int x;
for(x = 0; x < 7; x++)
{
if(r1.v[x] != r2.v[x])
return false;
}
return true;
}
};
struct o7
{
__host__ __device__
bool operator()(const s7 &r1, const s7 &r2)
{
int x;
for(x = 0; x < 7; x++)
{
if(r1.v[x] > r2.v[x])
return true;
if(r1.v[x] < r2.v[x])
return false;
}
return false;
}
};
int unir(int *res, int rows, int tipo)
{
int flag, nrows;
#if TIMER
cuda_stats.unions++;
#endif
switch(tipo)
{
case 1:
{
thrust::device_ptr<int> pt, re;
pt = thrust::device_pointer_cast(res);
flag = 0;
while(flag != 1)
{
try
{
thrust::sort(pt, pt + rows);
re = thrust::unique(pt, pt + rows);
flag = 1;
}
catch(std::bad_alloc &e)
{
limpiar("sort/unique in unir", 0);
}
}
nrows = thrust::distance(pt, re);
thrust::device_vector<int> iVec(pt, pt + rows);
iVec.resize(nrows);
iVec.shrink_to_fit();
return nrows;
}
case 2:
{
thrust::device_ptr<s2> pt2, re2;
s2 *t2;
t2 = (s2*)res;
/*int *a, x, y;
a = (int *)malloc(rows * 2 * sizeof(int));
cudaMemcpy(a, res, rows * 2 * sizeof(int), cudaMemcpyDeviceToHost);
cout << "INI" << endl;
for(x = 0; x < rows; x++)
{
for(y = 0; y < 2; y++)
cout << a[x * 2 + y] << " ";
cout << endl;
}
cout << "INI fin" << endl;
free(a);*/
pt2 = thrust::device_pointer_cast(t2);
flag = 0;
while(flag != 1)
{
try
{
thrust::sort(pt2, pt2 + rows, o2());
re2 = thrust::unique(pt2, pt2 + rows, p2());
flag = 1;
}
catch(std::bad_alloc &e)
{
limpiar("sort/unique in unir", 0);
}
}
nrows = thrust::distance(pt2, re2);
thrust::device_vector<s2> iVec(pt2, pt2 + rows);
iVec.resize(nrows);
iVec.shrink_to_fit();
/*tam = (int)(re2.get() - pt2.get());
a = (int *)malloc(tam * 2 * sizeof(int));
cudaMemcpy(a, res, tam * 2 * sizeof(int), cudaMemcpyDeviceToHost);
cout << "FIN" << endl;
for(x = 0; x < tam; x++)
{
for(y = 0; y < 2; y++)
cout << a[x * 2 + y] << " ";
cout << endl;
}
cout << "FIN fin" << endl;
free(a);
cout << "antes = " << rows << " despues = " << thrust::distance(pt2, re2) << endl;*/
return nrows;
}
case 3:
{
thrust::device_ptr<s3> pt3, re3;
s3 *t3;
t3 = (s3*)res;
pt3 = thrust::device_pointer_cast(t3);
flag = 0;
while(flag != 1)
{
try
{
thrust::sort(pt3, pt3 + rows, o3());
re3 = thrust::unique(pt3, pt3 + rows, p3());
flag = 1;
}
catch(std::bad_alloc &e)
{
limpiar("sort/unique in unir", 0);
}
}
nrows = thrust::distance(pt3, re3);
thrust::device_vector<s3> iVec(pt3, pt3 + rows);
iVec.resize(nrows);
iVec.shrink_to_fit();
return nrows;
}
case 4:
{
thrust::device_ptr<s4> pt4, re4;
s4 *t4;
t4 = (s4*)res;
pt4 = thrust::device_pointer_cast(t4);
flag = 0;
while(flag != 1)
{
try
{
thrust::sort(pt4, pt4 + rows, o4());
re4 = thrust::unique(pt4, pt4 + rows, p4());
flag = 1;
}
catch(std::bad_alloc &e)
{
limpiar("sort/unique in unir", 0);
}
}
nrows = thrust::distance(pt4, re4);
thrust::device_vector<s4> iVec(pt4, pt4 + rows);
iVec.resize(nrows);
iVec.shrink_to_fit();
return nrows;
}
case 5:
{
thrust::device_ptr<s5> pt5, re5;
s5 *t5;
t5 = (s5*)res;
pt5 = thrust::device_pointer_cast(t5);
flag = 0;
while(flag != 1)
{
try
{
thrust::sort(pt5, pt5 + rows, o5());
re5 = thrust::unique(pt5, pt5 + rows, p5());
flag = 1;
}
catch(std::bad_alloc &e)
{
limpiar("sort/unique in unir", 0);
}
}
nrows = thrust::distance(pt5, re5);
thrust::device_vector<s5> iVec(pt5, pt5 + rows);
iVec.resize(nrows);
iVec.shrink_to_fit();
return nrows;
}
case 6:
{
thrust::device_ptr<s6> pt6, re6;
s6 *t6;
t6 = (s6*)res;
pt6 = thrust::device_pointer_cast(t6);
flag = 0;
while(flag != 1)
{
try
{
thrust::sort(pt6, pt6 + rows, o6());
re6 = thrust::unique(pt6, pt6 + rows, p6());
flag = 1;
}
catch(std::bad_alloc &e)
{
limpiar("sort/unique in unir", 0);
}
}
nrows = thrust::distance(pt6, re6);
thrust::device_vector<s6> iVec(pt6, pt6 + rows);
iVec.resize(nrows);
iVec.shrink_to_fit();
return nrows;
}
case 7:
{
thrust::device_ptr<s7> pt7, re7;
s7 *t7;
t7 = (s7*)res;
pt7 = thrust::device_pointer_cast(t7);
flag = 0;
while(flag != 1)
{
try
{
thrust::sort(pt7, pt7 + rows, o7());
re7 = thrust::unique(pt7, pt7 + rows, p7());
flag = 1;
}
catch(std::bad_alloc &e)
{
limpiar("sort/unique in unir", 0);
}
}
nrows = thrust::distance(pt7, re7);
thrust::device_vector<s7> iVec(pt7, pt7 + rows);
iVec.resize(nrows);
iVec.shrink_to_fit();
return nrows;
}
default:
cerr << "Union: " << tipo << " columns are too many." << endl;
exit(1);
}
return 0;
}