OpenCLSDKOverlay.swift

// RUN: %target-run-simple-swift
// REQUIRES: executable_test
// REQUIRES: OS=macosx

// Translated from standard OpenCL hello.c program

// Abstract:   A simple "Hello World" compute example showing basic usage of OpenCL which
//             calculates the mathematical square (X[i] = pow(X[i],2)) for a buffer of
//             floating point values.
//             
//
// 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.
//


import OpenCL
import StdlibUnittest
import Foundation
import Darwin

let KernelSource = "\n" +
"__kernel void square(                                                       \n" +
"   __global float* input,                                              \n" +
"   __global float* output,                                             \n" +
"   const unsigned int count)                                           \n" +
"{                                                                      \n" +
"   int i = get_global_id(0);                                           \n" +
"   if(i < count)                                                       \n" +
"       output[i] = input[i] * input[i];                                \n" +
"}                                                                      \n" +
"\n"


let DATA_SIZE = 1024

var tests = TestSuite("MiscSDKOverlay")

tests.test("clSetKernelArgsListAPPLE") {
  var err: cl_int                            // error code returned from api calls
  
  var data = [Float](count: DATA_SIZE, repeatedValue: 0)              // original data set given to device
  var results = [Float](count: DATA_SIZE, repeatedValue: 0)           // results returned from device
  var correct: Int               // number of correct results returned

  var global: size_t                      // global domain size for our calculation
  var local: size_t = 0                       // local domain size for our calculation

  var device_id: cl_device_id = nil             // compute device id 
  var context: cl_context                 // compute context
  var commands: cl_command_queue          // compute command queue
  var program: cl_program                 // compute program
  var kernel: cl_kernel                   // compute kernel
  
  var input: cl_mem                       // device memory used for the input array
  var output: cl_mem                      // device memory used for the output array
  
  // Fill our data set with random float values
  //
  var i = 0
  var count = DATA_SIZE
  for i = 0; i < count; i++ {
    data[i] = Float(rand()) / Float(RAND_MAX)
  }
  
  // Connect to a compute device
  //
  var gpu = 1
  err = clGetDeviceIDs(nil, cl_device_type(gpu != 0 ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU), 1, &device_id, nil)
  if (err != CL_SUCCESS)
  {
    print("Error: Failed to create a device group!")
    exit(EXIT_FAILURE)
  }
  
  // Create a compute context 
  //
  context = clCreateContext(nil, 1, &device_id, nil, nil, &err)
  if (context == nil)
  {
    print("Error: Failed to create a compute context!")
    exit(EXIT_FAILURE)
  }

  // Create a command commands
  //
  commands = clCreateCommandQueue(context, device_id, 0, &err)
  if (commands == nil)
  {
    print("Error: Failed to create a command commands!")
    exit(EXIT_FAILURE)
  }

  // Create the compute program from the source buffer
  //
  program = KernelSource.withCString {
    (var s: UnsafePointer<CChar>)->cl_program in
    withUnsafeMutablePointer(&s) {
      return clCreateProgramWithSource(context, 1, $0, nil, &err)
    }
  }
  if (program == nil)
  {
    print("Error: Failed to create compute program!")
    exit(EXIT_FAILURE)
  }

  // Build the program executable
  //
  err = clBuildProgram(program, 0, nil, nil, nil, nil)
  if (err != CL_SUCCESS)
  {
    var len: Int = 0
    
    var buffer = [CChar](count:2048, repeatedValue: 0)

    print("Error: Failed to build program executable!")
    clGetProgramBuildInfo(
      program, device_id, cl_program_build_info(CL_PROGRAM_BUILD_LOG), 2048, &buffer, &len)
    print("\(String.fromCString(buffer)!)")
    exit(1)
  }

  // Create the compute kernel in the program we wish to run
  //
  kernel = clCreateKernel(program, "square", &err)
  if (kernel == nil || err != cl_int(CL_SUCCESS))
  {
    print("Error: Failed to create compute kernel!")
    exit(1)
  }

  // Create the input and output arrays in device memory for our calculation
  //
  input = clCreateBuffer(context,  cl_mem_flags(CL_MEM_READ_ONLY),  sizeof(Float.self) * count, nil, nil)
  output = clCreateBuffer(context, cl_mem_flags(CL_MEM_WRITE_ONLY), sizeof(Float.self) * count, nil, nil)
  if (input == nil || output == nil)
  {
    print("Error: Failed to allocate device memory!")
    exit(1)
  }    
  
  // Write our data set into the input array in device memory 
  //
  err = clEnqueueWriteBuffer(commands, input, cl_bool(CL_TRUE), 0, sizeof(Float.self) * count, data, 0, nil, nil)
  if (err != CL_SUCCESS)
  {
    print("Error: Failed to write to source array!")
    exit(1)
  }

  // Set the arguments to our compute kernel
  //
  err = 0
  err = withUnsafePointers(&input, &output, &count) {
    inputPtr, outputPtr, countPtr in
    clSetKernelArgsListAPPLE(
      kernel, 3,
      0, sizeof(cl_mem.self), inputPtr,
      1, sizeof(cl_mem.self), outputPtr,
      2, sizeofValue(count), countPtr)
  }

  
  if (err != CL_SUCCESS)
  {
    print("Error: Failed to set kernel arguments! \(err)")
    exit(1)
  }

  // Get the maximum work group size for executing the kernel on the device
  //
  err = clGetKernelWorkGroupInfo(kernel, device_id, cl_kernel_work_group_info(CL_KERNEL_WORK_GROUP_SIZE), sizeofValue(local), &local, nil)
  if (err != CL_SUCCESS)
  {
    print("Error: Failed to retrieve kernel work group info! \(err)")
    exit(1)
  }

  // Execute the kernel over the entire range of our 1d input data set
  // using the maximum number of work group items for this device
  //
  global = count
  err = clEnqueueNDRangeKernel(commands, kernel, 1, nil, &global, &local, 0, nil, nil)
  if (err != 0)
  {
    print("Error: Failed to execute kernel!")
    exit(EXIT_FAILURE)
  }

  // Wait for the command commands to get serviced before reading back results
  //
  clFinish(commands)

  // Read back the results from the device to verify the output
  //
  err = clEnqueueReadBuffer( commands, output, cl_bool(CL_TRUE), 0, sizeof(Float.self) * count, &results, cl_uint(0), nil, nil )
  if (err != CL_SUCCESS)
  {
    print("Error: Failed to read output array! \(err)")
    exit(1)
  }
  
  // Validate our results
  //
  correct = 0
  for(i = 0; i < count; i++)
  {
    if(results[i] == data[i] * data[i]){
      correct++
    }
  }
  
  // Print a brief summary detailing the results
  //
  print("Computed '\(correct)/\(count)' correct values!")
  
  // Shutdown and cleanup
  //
  clReleaseMemObject(input)
  clReleaseMemObject(output)
  clReleaseProgram(program)
  clReleaseKernel(kernel)
  clReleaseCommandQueue(commands)
  clReleaseContext(context)
}

runAllTests()