import ThrustRTC as trtc
d_keys = trtc.device_vector_from_list([0, 0, 0, 1, 1, 2, 3, 3, 3, 3],
'int32_t')
d_values = trtc.device_vector_from_list([1, 1, 1, 1, 1, 1, 1, 1, 1, 1],
'int32_t')
trtc.Inclusive_Scan_By_Key(d_keys, d_values, d_values)
print(d_values.to_host())
d_values = trtc.device_vector_from_list([1, 1, 1, 1, 1, 1, 1, 1, 1, 1],
'int32_t')
trtc.Exclusive_Scan_By_Key(d_keys, d_values, d_values)
print(d_values.to_host())
d_values = trtc.device_vector_from_list([1, 1, 1, 1, 1, 1, 1, 1, 1, 1],
'int32_t')
trtc.Exclusive_Scan_By_Key(d_keys, d_values, d_values, trtc.DVInt32(5))
print(d_values.to_host())
d_int_in = trtc.device_vector_from_list([0, 1, 2, 3, 4], 'int32_t')
d_float_in = trtc.device_vector_from_list([0.0, 10.0, 20.0, 30.0, 40.0],
'float')
d_int_out = trtc.device_vector('int32_t', 5)
d_float_out = trtc.device_vector('float', 5)
zipped_in = trtc.DVZipped([d_int_in, d_float_in], ['a', 'b'])
zipped_out = trtc.DVZipped([d_int_out, d_float_out], ['a', 'b'])
trtc.Copy(zipped_in, zipped_out)
print(d_int_out.to_host())
print(d_float_out.to_host())
d_int_in = trtc.DVCounter(trtc.DVInt32(0), 5)
d_float_in = trtc.DVTransform(
d_int_in, "float",
trtc.Functor({}, ['i'], ' return (float)i*10.0f +10.0f;\n'))
zipped_in = trtc.DVZipped([d_int_in, d_float_in], ['a', 'b'])
trtc.Copy(zipped_in, zipped_out)
print(d_int_out.to_host())
print(d_float_out.to_host())
const_in = trtc.DVConstant(
trtc.DVTuple({
'a': trtc.DVInt32(123),
'b': trtc.DVFloat(456.0)
}), 5)
trtc.Copy(const_in, zipped_out)
print(d_int_out.to_host())
import ThrustRTC as trtc
is_less_than_zero = trtc.Functor({}, ['x'], '''
return x<0;
''')
darr1 = trtc.device_vector_from_list([1, 2, 3, 1, 2], 'int32_t')
trtc.Replace(darr1, trtc.DVInt32(1), trtc.DVInt32(99))
print(darr1.to_host())
darr2 = trtc.device_vector_from_list([1, -2, 3, -4, 5], 'int32_t')
trtc.Replace_If(darr2, is_less_than_zero, trtc.DVInt32(0))
print(darr2.to_host())
darr3_in = trtc.device_vector_from_list([1, 2, 3, 1, 2], 'int32_t')
darr3_out = trtc.device_vector('int32_t', 5)
trtc.Replace_Copy(darr3_in, darr3_out, trtc.DVInt32(1), trtc.DVInt32(99))
print(darr3_out.to_host())
darr4_in = trtc.device_vector_from_list([1, -2, 3, -4, 5], 'int32_t')
darr4_out = trtc.device_vector('int32_t', 5)
trtc.Replace_Copy_If(darr4_in, darr4_out, is_less_than_zero, trtc.DVInt32(0))
print(darr4_out.to_host())
import ThrustRTC as trtc d_input = trtc.device_vector_from_list([0, 2, 5, 7, 8], 'int32_t') print(trtc.Lower_Bound(d_input, trtc.DVInt32(0))) print(trtc.Lower_Bound(d_input, trtc.DVInt32(1))) print(trtc.Lower_Bound(d_input, trtc.DVInt32(2))) print(trtc.Lower_Bound(d_input, trtc.DVInt32(3))) print(trtc.Lower_Bound(d_input, trtc.DVInt32(8))) print(trtc.Lower_Bound(d_input, trtc.DVInt32(9))) print() print(trtc.Upper_Bound(d_input, trtc.DVInt32(0))) print(trtc.Upper_Bound(d_input, trtc.DVInt32(1))) print(trtc.Upper_Bound(d_input, trtc.DVInt32(2))) print(trtc.Upper_Bound(d_input, trtc.DVInt32(3))) print(trtc.Upper_Bound(d_input, trtc.DVInt32(8))) print(trtc.Upper_Bound(d_input, trtc.DVInt32(9))) print() print(trtc.Binary_Search(d_input, trtc.DVInt32(0))) print(trtc.Binary_Search(d_input, trtc.DVInt32(1))) print(trtc.Binary_Search(d_input, trtc.DVInt32(2))) print(trtc.Binary_Search(d_input, trtc.DVInt32(3))) print(trtc.Binary_Search(d_input, trtc.DVInt32(8))) print(trtc.Binary_Search(d_input, trtc.DVInt32(9)))
import ThrustRTC as trtc
negate = trtc.Functor( {}, ['x'],
'''
return -x;
''')
darr = trtc.device_vector('int32_t', 10)
trtc.Transform(trtc.DVCounter(trtc.DVInt32(5), 10), darr, trtc.Negate())
print (darr.to_host())
import ThrustRTC as trtc darr = trtc.device_vector_from_list([3, 7, 2, 5 ], 'int32_t') trtc.Transform_Binary(darr, trtc.DVConstant(trtc.DVInt32(10)), darr, trtc.Plus()) print (darr.to_host())
import ThrustRTC as trtc
trtc.Set_Verbose()
trtc.Transform(trtc.DVCounter(trtc.DVInt32(5), 10), trtc.DVDiscard("int32_t"),
trtc.Negate())
import ThrustRTC as trtc
darr = trtc.device_vector('int32_t', 5)
trtc.Fill(darr, trtc.DVInt32(123))
print(darr.to_host())
trtc.Fill(darr.range(1, 3), trtc.DVInt32(456))
print(darr.to_host())
import ThrustRTC as trtc
d_values = trtc.device_vector_from_list([0, 5, 3, 7], 'int32_t')
print(trtc.Find(d_values, trtc.DVInt32(3)))
print(trtc.Find(d_values, trtc.DVInt32(5)))
print(trtc.Find(d_values, trtc.DVInt32(9)))
print(trtc.Find_If(d_values, trtc.Functor({}, ['x'], ' return x>4;\n')))
print(trtc.Find_If(d_values, trtc.Functor({}, ['x'],
' return x>10;\n')))
print(
trtc.Find_If_Not(d_values, trtc.Functor({}, ['x'],
' return x>4;\n')))
print(
trtc.Find_If_Not(d_values, trtc.Functor({}, ['x'],
' return x>10;\n')))
import ThrustRTC as trtc
d_value = trtc.device_vector_from_list([3, 1, 4, 1, 5, 9], 'int32_t')
count = trtc.Remove(d_value, trtc.DVInt32(1))
print(d_value.to_host(0, count))
d_in = trtc.device_vector_from_list([-2, 0, -1, 0, 1, 2], 'int32_t')
d_out = trtc.device_vector('int32_t', 6)
count = trtc.Remove_Copy(d_in, d_out, trtc.DVInt32(0))
print(d_out.to_host(0, count))
is_even = trtc.Functor({}, ['x'], '''
return x % 2 == 0;
''')
d_value = trtc.device_vector_from_list([1, 4, 2, 8, 5, 7], 'int32_t')
count = trtc.Remove_If(d_value, is_even)
print(d_value.to_host(0, count))
d_in = trtc.device_vector_from_list([-2, 0, -1, 0, 1, 2], 'int32_t')
d_out = trtc.device_vector('int32_t', 6)
count = trtc.Remove_Copy_If(d_in, d_out, is_even)
print(d_out.to_host(0, count))
d_value = trtc.device_vector_from_list([1, 4, 2, 8, 5, 7], 'int32_t')
d_stencil = trtc.device_vector_from_list([0, 1, 1, 1, 0, 0], 'int32_t')
count = trtc.Remove_If_Stencil(d_value, d_stencil, trtc.Identity())
print(d_value.to_host(0, count))
d_in = trtc.device_vector_from_list([-2, 0, -1, 0, 1, 2], 'int32_t')
d_stencil = trtc.device_vector_from_list([1, 1, 0, 1, 0, 1], 'int32_t')
import ThrustRTC as trtc
darr = trtc.device_vector_from_list([1, 0, 2, 2, 1, 3], 'int32_t')
trtc.Transform_Inclusive_Scan(darr, darr, trtc.Negate(), trtc.Plus())
print(darr.to_host())
darr = trtc.device_vector_from_list([1, 0, 2, 2, 1, 3], 'int32_t')
trtc.Transform_Exclusive_Scan(darr, darr, trtc.Negate(), trtc.DVInt32(4),
trtc.Plus())
print(darr.to_host())
import ThrustRTC as trtc
darr = trtc.device_vector('int32_t', 10)
trtc.Sequence(darr)
print(darr.to_host())
trtc.Sequence(darr, trtc.DVInt32(1))
print(darr.to_host())
trtc.Sequence(darr, trtc.DVInt32(1), trtc.DVInt32(3))
print(darr.to_host())
import ThrustRTC as trtc
darr = trtc.device_vector_from_list([1, 0, 2, 2, 1, 3], 'int32_t')
print(trtc.Reduce(darr))
print(trtc.Reduce(darr, trtc.DVInt32(1)))
print(trtc.Reduce(darr, trtc.DVInt32(-1), trtc.Maximum()))
d_keys_in = trtc.device_vector_from_list([1, 3, 3, 3, 2, 2, 1], 'int32_t')
d_value_in = trtc.device_vector_from_list([9, 8, 7, 6, 5, 4, 3], 'int32_t')
d_keys_out = trtc.device_vector('int32_t', 7)
d_values_out = trtc.device_vector('int32_t', 7)
count = trtc.Reduce_By_Key(d_keys_in, d_value_in, d_keys_out, d_values_out)
print (d_keys_out.to_host(0, count))
print (d_values_out.to_host(0, count))
count = trtc.Reduce_By_Key(d_keys_in, d_value_in, d_keys_out, d_values_out, trtc.EqualTo())
print (d_keys_out.to_host(0, count))
print (d_values_out.to_host(0, count))
count = trtc.Reduce_By_Key(d_keys_in, d_value_in, d_keys_out, d_values_out, trtc.EqualTo(), trtc.Plus())
print (d_keys_out.to_host(0, count))
print (d_values_out.to_host(0, count))
harr = np.array([1.0, 2.0, 3.0, 4.0, 5.0], dtype='float32')
darr = trtc.device_vector_from_numpy(harr)
print(darr.to_host())
# C data type
print(darr.name_view_cls())
harr2 = np.array([6,7,8,9,10], dtype='int32')
darr2 = trtc.device_vector_from_numpy(harr2)
# kernel with auto parameters, launched twice with different types
kernel = trtc.Kernel(['arr_in', 'arr_out', 'k'],
'''
size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= arr_in.size()) return;
arr_out[idx] = arr_in[idx]*k;
''')
darr_out = trtc.device_vector('float', 5)
kernel.launch(1,128, [darr, darr_out, trtc.DVFloat(10.0)])
print (darr_out.to_host())
darr_out = trtc.device_vector('int32_t', 5)
kernel.launch(1,128, [darr2, darr_out, trtc.DVInt32(5)])
print (darr_out.to_host())
# create a vector from python list with C type specified
darr3 = trtc.device_vector_from_list([3.0, 5.0, 7.0, 9.0 , 11.0], 'float')
print(darr3.to_host())
def demo_k_means(d_x, d_y, k):
n = d_x.size()
# create a zipped vector for convenience
d_points = trtc.DVZipped([d_x, d_y], ['x','y'])
# operations
point_plus = trtc.Functor({ }, ['pos1', "pos2"],
'''
return decltype(pos1)({pos1.x + pos2.x, pos1.y + pos2.y});
''')
point_div = trtc.Functor({ }, ['pos', "count"],
'''
return decltype(pos)({pos.x/(float)count, pos.y/(float)count});
''')
# initialize centers
center_ids = [0] * k
d_min_dis = trtc.device_vector("float", n)
for i in range(1, k):
d_count = trtc.DVInt32(i)
d_center_ids = trtc.device_vector_from_list(center_ids[0:i], 'int32_t')
calc_min_dis = trtc.Functor({"points": d_points, "center_ids": d_center_ids, "count": d_count }, ['pos'],
'''
float minDis = FLT_MAX;
for (int i=0; i<count; i++)
{
int j = center_ids[i];
float dis = (pos.x - points[j].x)*(pos.x - points[j].x);
dis+= (pos.y - points[j].y)*(pos.y - points[j].y);
if (dis<minDis) minDis = dis;
}
return minDis;
''')
trtc.Transform(d_points, d_min_dis, calc_min_dis)
center_ids[i] = trtc.Max_Element(d_min_dis)
d_count = trtc.DVInt32(k)
d_center_ids = trtc.device_vector_from_list(center_ids, 'int32_t')
# initialize group-average values
d_group_aves_x = trtc.device_vector("float", k)
d_group_aves_y = trtc.device_vector("float", k)
d_group_aves = trtc.DVZipped([d_group_aves_x, d_group_aves_y], ['x','y'])
trtc.Gather(d_center_ids, d_points, d_group_aves)
# initialize labels
d_labels = trtc.device_vector("int32_t", n)
trtc.Fill(d_labels, trtc.DVInt32(-1))
# buffer for new-calculated lables
d_labels_new = trtc.device_vector("int32_t", n)
d_labels_sink = trtc.DVDiscard("int32_t", k)
d_group_sums = trtc.device_vector(d_points.name_elem_cls(), k)
d_group_cumulate_counts = trtc.device_vector("int32_t", k)
d_group_counts = trtc.device_vector("int32_t", k)
d_counter = trtc.DVCounter(trtc.DVInt32(0), k)
# iterations
while True:
# calculate new labels
calc_new_labels = trtc.Functor({"aves": d_group_aves, "count": d_count }, ['pos'],
'''
float minDis = FLT_MAX;
int label = -1;
for (int i=0; i<count; i++)
{
float dis = (pos.x - aves[i].x)*(pos.x - aves[i].x);
dis+= (pos.y - aves[i].y)*(pos.y - aves[i].y);
if (dis<minDis)
{
minDis = dis;
label = i;
}
}
return label;
''')
trtc.Transform(d_points, d_labels_new, calc_new_labels)
if trtc.Equal(d_labels, d_labels_new):
break
trtc.Copy(d_labels_new, d_labels)
# recalculate group-average values
trtc.Sort_By_Key(d_labels, d_points)
trtc.Reduce_By_Key(d_labels, d_points, d_labels_sink, d_group_sums, trtc.EqualTo(), point_plus)
trtc.Upper_Bound_V(d_labels, d_counter, d_group_cumulate_counts)
trtc.Adjacent_Difference(d_group_cumulate_counts, d_group_counts)
trtc.Transform_Binary(d_group_sums, d_group_counts, d_group_aves, point_div)
h_x = d_x.to_host()
h_y = d_y.to_host()
h_labels = d_labels.to_host()
h_group_aves_x = d_group_aves_x.to_host()
h_group_aves_y = d_group_aves_y.to_host()
h_group_counts = d_group_counts.to_host()
lines = []
for i in range(n):
label = h_labels[i]
lines.append([(h_x[i], h_y[i]), (h_group_aves_x[label], h_group_aves_y[label]) ] )
lc = mc.LineCollection(lines)
fig, ax = plt.subplots()
ax.set_xlim((0, 1000))
ax.set_ylim((0, 1000))
ax.add_collection(lc)
plt.show()
import ThrustRTC as trtc
op = trtc.Functor( {}, ['x'],
'''
return x % 100;
''')
darr = trtc.device_vector('int32_t', 2000)
trtc.Transform(trtc.DVCounter(trtc.DVInt32(0), 2000), darr, op)
print(trtc.Count(darr, trtc.DVInt32(47)))
op2 = trtc.Functor({}, ['x'],
'''
return (x % 100)==47;
''')
trtc.Sequence(darr)
print(trtc.Count_If(darr, op2))
import ThrustRTC as trtc
darr = trtc.device_vector_from_list([ -1, 0, -2, -2, 1, -3], 'int32_t')
absolute_value = trtc.Functor( {}, ['x'],
'''
return x<(decltype(x))0 ? -x : x;
''')
print(trtc.Transform_Reduce(darr, absolute_value, trtc.DVInt32(0), trtc.Maximum()))
import ThrustRTC as trtc darr = trtc.device_vector_from_list([1, 0, 2, 2, 1, 3], 'int32_t') trtc.Inclusive_Scan(darr, darr) print (darr.to_host()) darr = trtc.device_vector_from_list([-5, 0, 2, -3, 2, 4, 0, -1, 2, 8], 'int32_t') trtc.Inclusive_Scan(darr, darr, trtc.Maximum()) print (darr.to_host()) darr = trtc.device_vector_from_list([1, 0, 2, 2, 1, 3], 'int32_t') trtc.Exclusive_Scan(darr, darr) print (darr.to_host()) darr = trtc.device_vector_from_list([1, 0, 2, 2, 1, 3], 'int32_t') trtc.Exclusive_Scan(darr, darr, trtc.DVInt32(4)) print (darr.to_host()) darr = trtc.device_vector_from_list([-5, 0, 2, -3, 2, 4, 0, -1, 2, 8], 'int32_t') trtc.Exclusive_Scan(darr, darr, trtc.DVInt32(1), trtc.Maximum()) print (darr.to_host())