cudaDB/device_string.h at release · ColdZoo/cudaDB · GitHub

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
//
// Created by adam on 6/30/19.
//

//Taken from http://ldn.linuxfoundation.org/article/c-gpu-and-thrust-strings-gpu
//Also, https://groups.google.com/group/thrust-users/msg/0eac80d2e41cbcfb?pli=1, https://groups.google.com/group/thrust-users/browse_thread/thread/f4b1b825cc927df9?pli=1,

#ifndef CUDADB_PROJECT_DEVICE_STRING_H
#define CUDADB_PROJECT_DEVICE_STRING_H


#define POOL_SZ (10*1024*102400)


class device_string
{
public:
    int cstr_len;
    char* raw;
    thrust::device_ptr<char> cstr;

    static char* pool_raw;
    static thrust::device_ptr<char> pool_cstr;
    static thrust::device_ptr<char> pool_top;

    // Sets the variables up the first time its used.
    __host__ static void init()
    {
        static bool v = true;
        if( v )
        {
            v = false;

            pool_cstr = thrust::device_malloc<char>(POOL_SZ);
            pool_raw  = (char*)raw_pointer_cast( pool_cstr );
            pool_top = pool_cstr;
        }
    }
    // Destructor for device variables used.
    __host__ static void fin()
    {
        init();
        thrust::device_free(pool_cstr);
    }

    // Parametrized constructor to copy one device_string to another.
    __host__ __device__ device_string( const device_string& s )
    {
        cstr_len = s.cstr_len;
        raw = s.raw;
        cstr = s.cstr;
    }

    // Parametrized constructor to copy a std::string to device_string type
    __host__ device_string( const std::string& s)
    {
        cstr_len = s.length();
        init();
        cstr = pool_top;
        pool_top += cstr_len+1;
        raw = (char *) raw_pointer_cast(cstr);
        cudaMemcpy( raw, s.c_str(), cstr_len+1, cudaMemcpyHostToDevice );
    }


    // Default constructor.
    __host__ __device__ device_string()
    {
        cstr_len = -1;
        raw = NULL;
    }


    // Conversion operator to copy device_string type to std::string
    // This is where the problem is

    __host__ operator std::string(void)
    {
        std::string ret;
        thrust::copy(cstr, cstr+cstr_len, back_inserter(ret));
        return ret;
    }
};


char* device_string::pool_raw;
thrust::device_ptr<char> device_string::pool_cstr;
thrust::device_ptr<char> device_string::pool_top;

// User-defined comparison operator
bool __device__ operator< (device_string lhs, device_string rhs)
{
char *l = lhs.raw;
char *r = rhs.raw;

for( ; *l && *r && *l==*r; ){
++l;
++r;
}
return *l < *r;
}


bool __device__ __host__ operator==(device_string lhs, device_string rhs)
{
    char* l = lhs.raw;
    char* r = rhs.raw;
    // skip equal chars
    for(; *l && *r && *l==*r;){
        ++l;
        ++r;
    }
    if(!*l && !*r){
        return THRUST_TRUE;
    }
    return THRUST_FALSE;
}

#endif //CUDADB_PROJECT_DEVICE_STRING_H