compute_kernels.cl

// Version:    <1.0>
//
// Disclaimer: IMPORTANT:  This Apple software is supplied to you by Apple Inc. ("Apple")
//             in consideration of your agreement to the following terms, and your use,
//             installation, modification or redistribution of this Apple software
//             constitutes acceptance of these terms.  If you do not agree with these
//             terms, please do not use, install, modify or redistribute this Apple
//             software.
//
//             In consideration of your agreement to abide by the following terms, and
//             subject to these terms, Apple grants you a personal, non - exclusive
//             license, under Apple's copyrights in this original Apple software ( the
//             "Apple Software" ), to use, reproduce, modify and redistribute the Apple
//             Software, with or without modifications, in source and / or binary forms;
//             provided that if you redistribute the Apple Software in its entirety and
//             without modifications, you must retain this notice and the following text
//             and disclaimers in all such redistributions of the Apple Software. Neither
//             the name, trademarks, service marks or logos of Apple Inc. may be used to
//             endorse or promote products derived from the Apple Software without specific
//             prior written permission from Apple.  Except as expressly stated in this
//             notice, no other rights or licenses, express or implied, are granted by
//             Apple herein, including but not limited to any patent rights that may be
//             infringed by your derivative works or by other works in which the Apple
//             Software may be incorporated.
//
//             The Apple Software is provided by Apple on an "AS IS" basis.  APPLE MAKES NO
//             WARRANTIES, EXPRESS OR IMPLIED, INCLUDING WITHOUT LIMITATION THE IMPLIED
//             WARRANTIES OF NON - INFRINGEMENT, MERCHANTABILITY AND FITNESS FOR A
//             PARTICULAR PURPOSE, REGARDING THE APPLE SOFTWARE OR ITS USE AND OPERATION
//             ALONE OR IN COMBINATION WITH YOUR PRODUCTS.
//
//             IN NO EVENT SHALL APPLE BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL OR
//             CONSEQUENTIAL DAMAGES ( INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
//             SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
//             INTERRUPTION ) ARISING IN ANY WAY OUT OF THE USE, REPRODUCTION, MODIFICATION
//             AND / OR DISTRIBUTION OF THE APPLE SOFTWARE, HOWEVER CAUSED AND WHETHER
//             UNDER THEORY OF CONTRACT, TORT ( INCLUDING NEGLIGENCE ), STRICT LIABILITY OR
//             OTHERWISE, EVEN IF APPLE HAS BEEN ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Copyright ( C ) 2008 Apple Inc. All Rights Reserved.
//
////////////////////////////////////////////////////////////////////////////////////////////////////
 
__kernel void 
advanceintime(
    __global float4 *h0, 
    __global float2 *ht,
    const float twoPiOverWX, 
    const float twoPiOverWY, 
    const int N, 
    const int m,
    const float time)
{
    int i = get_global_id(0);
    int j = get_global_id(1);
    int index = (i << m) + j;
        
    int mi = (N - i) & (N - 1);
    int mj = (N - j) & (N - 1);
    int indexm = (mi << m) + mj;
    
    int ii = (i + (N >> 1)) & (N - 1);
    int jj = (j + (N >> 1)) & (N - 1);
    int indexo = (ii << m) + jj;
    
    int offset = 1 << (2*m);
    
    float4 hf0Wg1 = h0[index];
    float4 hf0Wg2 = h0[indexm];
    float wt = hf0Wg1.z*time;
    float cs = native_cos(wt);
    float sn = native_sin(wt);
    float2 t1 = (float2)(cs, -sn);
    float2 t2 = (float2)(sn,  cs);
    float2 H1 = (float2)(hf0Wg1.x, hf0Wg1.y);
    float2 H2 = (float2)(hf0Wg2.x, hf0Wg2.y);   
    float hor = dot(H1+H2, t1);
    float hoi = dot(H1-H2, t2);
    float2 h = (float2)(hor, hoi);
    float kx = twoPiOverWX * (i - (N >> 1));
    float ky = twoPiOverWY * (j - (N >> 1));
    float2 h2 = (float2)(h.y, -h.x);
    float4 h3 = (float4)(h2.x*kx, h2.y*kx, h2.x*ky, h2.y*ky);
 
    ht[indexo] = h;
    ht[indexo + offset] = (float2)(h3.x, h3.y);
    ht[indexo + (offset << 1)] = (float2)(h3.z, h3.w);
}