jszuppe

#include <vector>
#include <iostream>
#include <string>
#include <algorithm>
#include <memory>
#include <random>

#include <nvrtc.h>
#include <cuda.h>
#include <cuda_runtime_api.h>

jszuppe

#define boost_pair_type(t1, t2) _pair_ ## t1 ## _ ## t2 ## _t
#define boost_pair_get(x, n) (n == 0 ? x.first ## x.second)
#define boost_make_pair(t1, x, t2, y) (boost_pair_type(t1, t2)) { x, y }
#define boost_tuple_get(x, n) (x.v ## n)
inline int ret42(){ return 42; }
__kernel void copy(__global int* _buf0, const uint count)
{
uint index = get_local_id(0) + (512 * get_group_id(0));
for(uint i = 0; i < 4; i++){
    if(index < count){

jszuppe

// This kernel DOES NOT WORK
#define boost_pair_type(t1, t2) _pair_ ## t1 ## _ ## t2 ## _t
#define boost_pair_get(x, n) (n == 0 ? x.first ## x.second)
#define boost_make_pair(t1, x, t2, y) (boost_pair_type(t1, t2)) { x, y }
#define boost_tuple_get(x, n) (x.v ## n)

inline int ret42(){ return 42; }

__kernel void copy(__global int* _buf0, const uint count)
{

jszuppe

// This is for not packed structs but _NOT_ for cl_typeN structs
#define BOOST_COMPUTE_ADAPT_STRUCT_NOT_PACKED(type, name, members) \
    BOOST_COMPUTE_TYPE_NAME(type, name) \
    namespace boost { namespace compute { \
    template<> \
    inline std::string type_definition<type>() \
    { \
        std::stringstream declaration; \
        declaration << "typedef struct {\n" \
                    BOOST_PP_SEQ_FOR_EACH( \

jszuppe

template<class HostIterator, class DeviceIterator>
inline DeviceIterator copy(HostIterator first,
                           HostIterator last,
                           DeviceIterator result,
                           boost::compute::command_queue &queue)
{
...
    // map result buffer to host
    value_type *pointer = static_cast<value_type*>(
        queue.enqueue_map_buffer(

jszuppe

namespace detail {

template<class InputIterator, class OutputIterator, class BinaryFunction>
inline void dispatch_gpu_reduce(InputIterator first,
                                InputIterator last,
                                OutputIterator result,
                                BinaryFunction function,
                                command_queue &queue)
{
 ... 

jszuppe

uint x = 0;
        while(l_idx < r_idx && x < 32)
        {
            int mid_idx = (l_idx + r_idx) / 2;            
            int mid_value = input[mid_idx];
            if(x > 11)
            	printf("gid: %d, mine: %d, l_dix: %d, r_idx: %d, mid_idx: %d, mid_value: %d, x: %d\n", gid, mine_value, l_idx, r_idx, mid_idx, mid_value, x);
            if(mid_value > mine_value)  // 1st division
            {
                r_idx = mid_idx;