Sha256: 667e3a6b14d910579d7a83ac372f6ef0686c64add09e20c6d6d0da25bb9d4993
Contents?: true
Size: 1.88 KB
Versions: 17
Compression:
Stored size: 1.88 KB
Contents
<% unless defined?($cumo_narray_gen_tmpl_accum_index_kernel_included) %>
<% $cumo_narray_gen_tmpl_accum_index_kernel_included = 1 %>
<% unless type_name == 'robject' %>
<% [64,32].each do |i| %>
#define idx_t int<%=i%>_t
#if defined(__cplusplus)
#if 0
{ /* satisfy cc-mode */
#endif
} /* extern "C" { */
#endif
struct cumo_<%=type_name%>_min_index_int<%=i%>_impl {
struct MinAndArgMin {
dtype min;
idx_t argmin;
};
__device__ MinAndArgMin Identity(idx_t index) { return {DATA_MAX, index}; }
__device__ MinAndArgMin MapIn(dtype in, idx_t index) { return {in, index}; }
__device__ void Reduce(MinAndArgMin next, MinAndArgMin& accum) {
if (accum.min > next.min) {
accum = next;
}
}
__device__ idx_t MapOut(MinAndArgMin accum) { return accum.argmin; }
};
struct cumo_<%=type_name%>_max_index_int<%=i%>_impl {
struct MaxAndArgMax {
dtype max;
idx_t argmax;
};
__device__ MaxAndArgMax Identity(idx_t index) { return {DATA_MIN, index}; }
__device__ MaxAndArgMax MapIn(dtype in, idx_t index) { return {in, index}; }
__device__ void Reduce(MaxAndArgMax next, MaxAndArgMax& accum) {
if (accum.max < next.max) {
accum = next;
}
}
__device__ idx_t MapOut(MaxAndArgMax accum) { return accum.argmax; }
};
#if defined(__cplusplus)
extern "C" {
#if 0
} /* satisfy cc-mode */
#endif
#endif
void cumo_<%=type_name%>_min_index_int<%=i%>_kernel_launch(cumo_na_reduction_arg_t* arg)
{
cumo_reduce<dtype, idx_t, cumo_<%=type_name%>_min_index_int<%=i%>_impl>(*arg, cumo_<%=type_name%>_min_index_int<%=i%>_impl{});
}
void cumo_<%=type_name%>_max_index_int<%=i%>_kernel_launch(cumo_na_reduction_arg_t* arg)
{
cumo_reduce<dtype, idx_t, cumo_<%=type_name%>_max_index_int<%=i%>_impl>(*arg, cumo_<%=type_name%>_max_index_int<%=i%>_impl{});
}
#undef idx_t
<% end %>
<% end %>
<% end %>
Version data entries
17 entries across 17 versions & 1 rubygems