generate a copy kernel program
| 29 | |
| 30 | // generate a copy kernel program |
| 31 | compute::kernel make_copy_kernel(const compute::context& context) |
| 32 | { |
| 33 | // source for the copy_kernel program |
| 34 | const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( |
| 35 | __kernel void copy_kernel(__global const float *src, __global float *dst) |
| 36 | { |
| 37 | uint x = get_group_id(0) * TILE_DIM + get_local_id(0); |
| 38 | uint y = get_group_id(1) * TILE_DIM + get_local_id(1); |
| 39 | |
| 40 | uint width = get_num_groups(0) * TILE_DIM; |
| 41 | |
| 42 | for(uint i = 0 ; i < TILE_DIM ; i+= BLOCK_ROWS){ |
| 43 | dst[(y+i)*width +x] = src[(y+i)*width + x]; |
| 44 | } |
| 45 | } |
| 46 | ); |
| 47 | |
| 48 | // setup compilation flags for the copy program |
| 49 | std::stringstream options; |
| 50 | options << "-DTILE_DIM=" << TILE_DIM << " -DBLOCK_ROWS=" << BLOCK_ROWS; |
| 51 | |
| 52 | // create and build the copy program |
| 53 | compute::program program = |
| 54 | compute::program::build_with_source(source, context, options.str()); |
| 55 | |
| 56 | // create and return the copy kernel |
| 57 | return program.create_kernel("copy_kernel"); |
| 58 | } |
| 59 | |
| 60 | // generate a naive transpose kernel |
| 61 | compute::kernel make_naive_transpose_kernel(const compute::context& context) |