Skip to content

Commit 2e2459a

Browse files
author
fynv
committed
Adding demos
1 parent 6c80494 commit 2e2459a

File tree

5 files changed

+230
-3
lines changed

5 files changed

+230
-3
lines changed

‎TRTCContext.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -113,6 +113,7 @@ TRTCContext::TRTCContext()
113113

114114
this->add_code_block("#define DEVICE_ONLY\n");
115115
this->add_inlcude_filename("cstdint");
116+
this->add_inlcude_filename("cfloat");
116117
this->add_inlcude_filename("built_in.h");
117118
}
118119

‎internal/cuda_inline_headers_global.hpp‎

Lines changed: 44 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,9 @@
1+
/*
2+
These header implementations are copied from the project "jitify".
3+
https://github.com/NVIDIA/jitify
4+
Thanks for the work.
5+
*/
6+
17
static const char s_safe_header_climits[] =
28
"#pragma once\n"
39
"\n"
@@ -96,15 +102,50 @@ static const char s_safe_header_cstdint[] =
96102
"#define PTRDIFF_MAX INTPTR_MAX\n"
97103
"#define SIZE_MAX UINT64_MAX\n";
98104

105+
static const char s_safe_header_cfloat[] =
106+
"#pragma once\n"
107+
"\n"
108+
"inline __host__ __device__ float jitify_int_as_float(int i) "
109+
"{ union FloatInt { float f; int i; } fi; fi.i = i; return fi.f; }\n"
110+
"inline __host__ __device__ double jitify_longlong_as_double(long long i) "
111+
"{ union DoubleLongLong { double f; long long i; } fi; fi.i = i; return "
112+
"fi.f; }\n"
113+
"#define FLT_RADIX 2\n"
114+
"#define FLT_MANT_DIG 24\n"
115+
"#define DBL_MANT_DIG 53\n"
116+
"#define FLT_DIG 6\n"
117+
"#define DBL_DIG 15\n"
118+
"#define FLT_MIN_EXP -125\n"
119+
"#define DBL_MIN_EXP -1021\n"
120+
"#define FLT_MIN_10_EXP -37\n"
121+
"#define DBL_MIN_10_EXP -307\n"
122+
"#define FLT_MAX_EXP 128\n"
123+
"#define DBL_MAX_EXP 1024\n"
124+
"#define FLT_MAX_10_EXP 38\n"
125+
"#define DBL_MAX_10_EXP 308\n"
126+
"#define FLT_MAX jitify_int_as_float(2139095039)\n"
127+
"#define DBL_MAX jitify_longlong_as_double(9218868437227405311)\n"
128+
"#define FLT_EPSILON jitify_int_as_float(872415232)\n"
129+
"#define DBL_EPSILON jitify_longlong_as_double(4372995238176751616)\n"
130+
"#define FLT_MIN jitify_int_as_float(8388608)\n"
131+
"#define DBL_MIN jitify_longlong_as_double(4503599627370496)\n"
132+
"#define FLT_ROUNDS 1\n"
133+
"#if defined __cplusplus && __cplusplus >= 201103L\n"
134+
"#define FLT_EVAL_METHOD 0\n"
135+
"#define DECIMAL_DIG 21\n"
136+
"#endif\n";
137+
99138

100-
static int s_num_headers_global = 2;
139+
static int s_num_headers_global = 3;
101140
static const char* s_name_headers_global[] = {
102141
"climits",
103-
"cstdint"
142+
"cstdint",
143+
"cfloat"
104144
};
105145

106146
static const char* s_content_headers_global[] = {
107147
s_safe_header_climits,
108-
s_safe_header_cstdint
148+
s_safe_header_cstdint,
149+
s_safe_header_cfloat
109150
};
110151

‎python/CMakeLists.txt‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -62,6 +62,9 @@ set(PYTHON_TEST
6262
test/test_transform_iter.py
6363
test/test_zipped.py
6464
test/test_custom_vector.py
65+
66+
demo/histogram.py
67+
demo/k-means.py
6568
)
6669

6770
install(FILES ${PYTHON_TEST} DESTINATION test_python )

‎python/demo/histogram.py‎

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
import numpy as np
2+
import ThrustRTC as trtc
3+
import matplotlib.pyplot as plt
4+
5+
6+
def demo_histogram(ctx, d_data):
7+
8+
# sort data to bring equal elements together
9+
trtc.Sort(ctx, d_data)
10+
11+
# caculate 20 bins from 0~200
12+
# 1 extra to exclude possible negative values
13+
d_cumulative_histogram = trtc.device_vector(ctx, "int32_t", 21)
14+
15+
d_counter = trtc.DVCounter(ctx, trtc.DVFloat(0.0), 21)
16+
d_range_ends = trtc.DVTransform(ctx, d_counter, "float", trtc.Functor(ctx, {}, ['x'], ' return x*10.0;\n' ))
17+
18+
trtc.Upper_Bound_V(ctx, d_data, d_range_ends, d_cumulative_histogram)
19+
20+
d_histogram = trtc.device_vector(ctx, "int32_t", 21)
21+
trtc.Adjacent_Difference(ctx, d_cumulative_histogram, d_histogram)
22+
23+
h_histogram = d_histogram.to_host(1, 21)
24+
25+
# plot the histogram
26+
x_axis = [str(x) for x in np.arange(5, 200, 10)]
27+
positions = np.arange(len(x_axis))
28+
plt.bar(positions, h_histogram, align='center', alpha=0.5)
29+
plt.xticks(positions, x_axis)
30+
plt.ylabel('Count')
31+
plt.title('Histogram')
32+
33+
plt.show()
34+
35+
36+
37+
if __name__ == '__main__':
38+
39+
ctx = trtc.Context()
40+
41+
h_data = np.random.randn(2000).astype(np.float32) * 30.0 +100.0
42+
d_data = trtc.device_vector_from_numpy(ctx, h_data)
43+
44+
demo_histogram(ctx, d_data)
45+
46+

‎python/demo/k-means.py‎

Lines changed: 136 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,136 @@
1+
import numpy as np
2+
import ThrustRTC as trtc
3+
import matplotlib.pyplot as plt
4+
from matplotlib import collections as mc
5+
6+
def demo_k_means(ctx, d_x, d_y, k):
7+
8+
n = d_x.size()
9+
10+
# create a zipped vector for convenience
11+
d_points = trtc.DVZipped(ctx, [d_x, d_y], ['x','y'])
12+
13+
# operations
14+
point_plus = trtc.Functor(ctx, { }, ['pos1', "pos2"],
15+
'''
16+
return decltype(pos1)({pos1.x + pos2.x, pos1.y + pos2.y});
17+
''')
18+
19+
point_div = trtc.Functor(ctx, { }, ['pos', "count"],
20+
'''
21+
return decltype(pos)({pos.x/(float)count, pos.y/(float)count});
22+
''')
23+
24+
# initialize centers
25+
center_ids = [0] * k
26+
d_min_dis = trtc.device_vector(ctx, "float", n)
27+
28+
for i in range(1, k):
29+
d_count = trtc.DVInt32(i)
30+
d_center_ids = trtc.device_vector_from_list(ctx, center_ids[0:i], 'int32_t')
31+
calc_min_dis = trtc.Functor(ctx, {"points": d_points, "center_ids": d_center_ids, "count": d_count }, ['pos'],
32+
'''
33+
float minDis = FLT_MAX;
34+
for (int i=0; i<count; i++)
35+
{
36+
int j = center_ids[i];
37+
float dis = (pos.x - points[j].x)*(pos.x - points[j].x);
38+
dis+= (pos.y - points[j].y)*(pos.y - points[j].y);
39+
if (dis<minDis) minDis = dis;
40+
}
41+
return minDis;
42+
''')
43+
trtc.Transform(ctx, d_points, d_min_dis, calc_min_dis)
44+
center_ids[i] = trtc.Max_Element(ctx, d_min_dis)
45+
46+
d_count = trtc.DVInt32(k)
47+
d_center_ids = trtc.device_vector_from_list(ctx, center_ids, 'int32_t')
48+
49+
# initialize group-average values
50+
d_group_aves_x = trtc.device_vector(ctx, "float", k)
51+
d_group_aves_y = trtc.device_vector(ctx, "float", k)
52+
d_group_aves = trtc.DVZipped(ctx, [d_group_aves_x, d_group_aves_y], ['x','y'])
53+
54+
trtc.Gather(ctx, d_center_ids, d_points, d_group_aves)
55+
56+
# initialize labels
57+
d_labels = trtc.device_vector(ctx, "int32_t", n)
58+
trtc.Fill(ctx, d_labels, trtc.DVInt32(-1))
59+
60+
# buffer for new-calculated lables
61+
d_labels_new = trtc.device_vector(ctx, "int32_t", n)
62+
63+
d_labels_sink = trtc.DVDiscard(ctx, "int32_t", k)
64+
d_group_sums = trtc.device_vector(ctx, d_points.name_elem_cls(), k)
65+
d_group_cumulate_counts = trtc.device_vector(ctx, "int32_t", k)
66+
d_group_counts = trtc.device_vector(ctx, "int32_t", k)
67+
68+
d_counter = trtc.DVCounter(ctx, trtc.DVInt32(0), k)
69+
70+
# iterations
71+
while True:
72+
# calculate new labels
73+
calc_new_labels = trtc.Functor(ctx, {"aves": d_group_aves, "count": d_count }, ['pos'],
74+
'''
75+
float minDis = FLT_MAX;
76+
int label = -1;
77+
for (int i=0; i<count; i++)
78+
{
79+
float dis = (pos.x - aves[i].x)*(pos.x - aves[i].x);
80+
dis+= (pos.y - aves[i].y)*(pos.y - aves[i].y);
81+
if (dis<minDis)
82+
{
83+
minDis = dis;
84+
label = i;
85+
}
86+
}
87+
return label;
88+
''')
89+
trtc.Transform(ctx, d_points, d_labels_new, calc_new_labels)
90+
if trtc.Equal(ctx, d_labels, d_labels_new):
91+
break
92+
trtc.Copy(ctx, d_labels_new, d_labels)
93+
94+
# recalculate group-average values
95+
trtc.Sort_By_Key(ctx, d_labels, d_points)
96+
trtc.Reduce_By_Key(ctx, d_labels, d_points, d_labels_sink, d_group_sums, trtc.EqualTo(), point_plus)
97+
trtc.Upper_Bound_V(ctx, d_labels, d_counter, d_group_cumulate_counts)
98+
trtc.Adjacent_Difference(ctx, d_group_cumulate_counts, d_group_counts)
99+
trtc.Transform_Binary(ctx, d_group_sums, d_group_counts, d_group_aves, point_div)
100+
101+
h_x = d_x.to_host()
102+
h_y = d_y.to_host()
103+
h_labels = d_labels.to_host()
104+
h_group_aves_x = d_group_aves_x.to_host()
105+
h_group_aves_y = d_group_aves_y.to_host()
106+
h_group_counts = d_group_counts.to_host()
107+
108+
lines = []
109+
110+
for i in range(n):
111+
label = h_labels[i]
112+
lines.append([(h_x[i], h_y[i]), (h_group_aves_x[label], h_group_aves_y[label]) ] )
113+
114+
lc = mc.LineCollection(lines)
115+
116+
fig, ax = plt.subplots()
117+
ax.set_xlim((0, 1000))
118+
ax.set_ylim((0, 1000))
119+
120+
ax.add_collection(lc)
121+
122+
plt.show()
123+
124+
125+
126+
if __name__ == '__main__':
127+
128+
ctx = trtc.Context()
129+
130+
h_x = np.random.rand(1000).astype(np.float32)*1000.0
131+
h_y = np.random.rand(1000).astype(np.float32)*1000.0
132+
133+
d_x = trtc.device_vector_from_numpy(ctx, h_x)
134+
d_y = trtc.device_vector_from_numpy(ctx, h_y)
135+
136+
demo_k_means(ctx, d_x, d_y, 50)

0 commit comments

Comments
 (0)