92 lines
3.1 KiB
Common Lisp
92 lines
3.1 KiB
Common Lisp
|
/*
|
||
|
|
* Copyright (c) 2017, Intel Corporation
|
||
|
|
*
|
||
|
|
* Permission is hereby granted, free of charge, to any person obtaining a
|
||
|
|
* copy of this software and associated documentation files (the "Software"),
|
||
|
|
* to deal in the Software without restriction, including without limitation
|
||
|
|
* the rights to use, copy, modify, merge, publish, distribute, sublicense,
|
||
|
|
* and/or sell copies of the Software, and to permit persons to whom the
|
||
|
|
* Software is furnished to do so, subject to the following conditions:
|
||
|
|
*
|
||
|
|
* The above copyright notice and this permission notice shall be included
|
||
|
|
* in all copies or substantial portions of the Software.
|
||
|
|
*
|
||
|
|
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
|
||
|
|
* OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||
|
|
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
|
||
|
|
* THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
|
||
|
|
* OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
|
||
|
|
* ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
|
||
|
|
* OTHER DEALINGS IN THE SOFTWARE.
|
||
|
|
*/
|
||
|
|
|
||
|
|
void block_fn(size_t tid, int mul, __global int* res)
|
||
|
|
{
|
||
|
|
res[tid] = mul * 7 - 21;
|
||
|
|
}
|
||
|
|
|
||
|
|
kernel void simple_block_kernel(__global int* res)
|
||
|
|
{
|
||
|
|
int multiplier = 3;
|
||
|
|
size_t tid = get_global_id(0);
|
||
|
|
|
||
|
|
void (^kernelBlock)(void) = ^{ block_fn(tid, multiplier, res); };
|
||
|
|
|
||
|
|
res[tid] = -1;
|
||
|
|
queue_t def_q = get_default_queue();
|
||
|
|
ndrange_t ndrange = ndrange_1D(1);
|
||
|
|
int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);
|
||
|
|
if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }
|
||
|
|
}
|
||
|
|
|
||
|
|
void block_reflection(size_t scalar, __global uint* buffer, read_only image3d_t img, sampler_t sampler)
|
||
|
|
{
|
||
|
|
float4 color;
|
||
|
|
int4 coord;
|
||
|
|
buffer[0] = scalar;
|
||
|
|
coord.x = scalar;
|
||
|
|
coord.y = 0;
|
||
|
|
coord.z = 0;
|
||
|
|
|
||
|
|
int width = get_image_width( img );
|
||
|
|
int heigth = get_image_height( img );
|
||
|
|
int depth = get_image_depth( img );
|
||
|
|
|
||
|
|
int order = get_image_channel_order( img );
|
||
|
|
int type = get_image_channel_data_type( img );
|
||
|
|
|
||
|
|
color = read_imagef( img, sampler, coord );
|
||
|
|
buffer[1] = (uint)color.x;
|
||
|
|
buffer[2] = (uint)width;
|
||
|
|
buffer[3] = (uint)heigth;
|
||
|
|
buffer[4] = (uint)depth;
|
||
|
|
buffer[5] = (uint)order;
|
||
|
|
buffer[6] = (uint)type;
|
||
|
|
|
||
|
|
sampler_t samplerA = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
|
||
|
|
color = read_imagef( img, samplerA, coord );
|
||
|
|
|
||
|
|
buffer[7] = (uint)color.y;
|
||
|
|
|
||
|
|
queue_t def_q = get_default_queue();
|
||
|
|
ndrange_t ndrange = ndrange_1D(1);
|
||
|
|
if( scalar > 2 ){
|
||
|
|
scalar--;
|
||
|
|
int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, ^{block_reflection(scalar, buffer, img, sampler);});
|
||
|
|
}
|
||
|
|
}
|
||
|
|
|
||
|
|
kernel void kernel_reflection(sampler_t sampler, read_only image3d_t img, __global uint* buffer, size_t scalar )
|
||
|
|
{
|
||
|
|
size_t tid = get_global_id(0);
|
||
|
|
|
||
|
|
void (^kernelBlock)(void) = ^{ block_reflection(scalar, buffer, img, sampler); };
|
||
|
|
|
||
|
|
queue_t def_q = get_default_queue();
|
||
|
|
ndrange_t ndrange = ndrange_1D(1);
|
||
|
|
if( scalar > 0 ){
|
||
|
|
int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);
|
||
|
|
if(enq_res != CLK_SUCCESS) { buffer[2] = 1; }
|
||
|
|
}
|
||
|
|
}
|