//
// Copyright (c) 2017 The Khronos Group Inc.
// 
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//    http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "harness/compat.h"

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>

#include "procs.h"

const char *kernel_code =
"__kernel void test_kernel(\n"
"char%s c, uchar%s uc, short%s s, ushort%s us, int%s i, uint%s ui, float%s f,\n"
"__global float%s *result)\n"
"{\n"
"  result[0] = %s(c);\n"
"  result[1] = %s(uc);\n"
"  result[2] = %s(s);\n"
"  result[3] = %s(us);\n"
"  result[4] = %s(i);\n"
"  result[5] = %s(ui);\n"
"  result[6] = f;\n"
"}\n";

const char *kernel_code_long =
"__kernel void test_kernel_long(\n"
"long%s l, ulong%s ul,\n"
"__global float%s *result)\n"
"{\n"
"  result[0] = %s(l);\n"
"  result[1] = %s(ul);\n"
"}\n";

int test_parameter_types_long(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
  clMemWrapper results;
  int error;
  size_t global[3] = {1, 1, 1};
  float results_back[2*16];
  int count, index;
  const char* types[] = { "long", "ulong" };
  char kernel_string[8192];
  int sizes[] = {1, 2, 4, 8, 16};
  const char* size_strings[] = {"", "2", "4", "8", "16"};
  float expected;
  int total_errors = 0;
  int size_to_test;
  char *ptr;
  char convert_string[1024];
  size_t max_parameter_size;

  // We don't really care about the contents since we're just testing that the types work.
  cl_long l[16]={-21,-1,2,-3,4,-5,6,-7,8,-9,10,-11,12,-13,14,-15};
  cl_ulong ul[16]={22,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15};

  // Calculate how large our paramter size is to the kernel
  size_t parameter_size = sizeof(cl_long) + sizeof(cl_ulong);

  // Init our strings.
  kernel_string[0] = '\0';
  convert_string[0] = '\0';

  // Get the maximum parameter size allowed
  error = clGetDeviceInfo( device, CL_DEVICE_MAX_PARAMETER_SIZE, sizeof( max_parameter_size ), &max_parameter_size, NULL );
    test_error( error, "Unable to get max parameter size from device" );

  // Create the results buffer
  results = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_float)*2*16, NULL, &error);
  test_error(error, "clCreateBuffer failed");

  // Go over all the vector sizes
  for (size_to_test = 0; size_to_test < 5; size_to_test++) {
    clProgramWrapper program;
    clKernelWrapper kernel;

    size_t total_parameter_size = parameter_size*sizes[size_to_test] + sizeof(cl_mem);
    if (total_parameter_size > max_parameter_size) {
      log_info("Can not test with vector size %d because it would exceed the maximum allowed parameter size to the kernel. (%d > %d)\n",
               (int)sizes[size_to_test], (int)total_parameter_size, (int)max_parameter_size);
      continue;
    }

    log_info("Testing vector size %d\n", sizes[size_to_test]);

    // If size is > 1, then we need a explicit convert call.
    if (sizes[size_to_test] > 1) {
      sprintf(convert_string, "convert_float%s",  size_strings[size_to_test]);
    } else {
      sprintf(convert_string, " ");
    }

    // Build the kernel
    sprintf(kernel_string, kernel_code_long,
            size_strings[size_to_test], size_strings[size_to_test], size_strings[size_to_test],
            convert_string, convert_string
    );

    ptr = kernel_string;
    error = create_single_kernel_helper(context, &program, &kernel, 1, (const char **)&ptr, "test_kernel_long");
    test_error(error, "create single kernel failed");

    // Set the arguments
    for (count = 0; count < 2; count++) {
      switch (count) {
        case 0: error = clSetKernelArg(kernel, count, sizeof(cl_long)*sizes[size_to_test], &l); break;
        case 1: error = clSetKernelArg(kernel, count, sizeof(cl_ulong)*sizes[size_to_test], &ul); break;
        default: log_error("Test error"); break;
      }
      if (error)
        log_error("Setting kernel arg %d %s%s: ", count, types[count], size_strings[size_to_test]);
      test_error(error, "clSetKernelArgs failed");
    }
    error = clSetKernelArg(kernel, 2, sizeof(cl_mem), &results);
    test_error(error, "clSetKernelArgs failed");

    // Execute
    error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global, NULL, 0, NULL, NULL);
    test_error(error, "clEnqueueNDRangeKernel failed");

    // Read back the results
    error = clEnqueueReadBuffer(queue, results, CL_TRUE, 0, sizeof(cl_float)*2*16, results_back, 0, NULL, NULL);
    test_error(error, "clEnqueueReadBuffer failed");

    // Verify the results
    for (count = 0; count < 2; count++) {
      for (index=0; index < sizes[size_to_test]; index++) {
        switch (count) {
          case 0: expected = (float)l[index]; break;
          case 1: expected = (float)ul[index]; break;
          default: log_error("Test error"); break;
        }

        if (results_back[count*sizes[size_to_test]+index] != expected) {
          total_errors++;
          log_error("Conversion from %s%s failed: index %d got %g, expected %g.\n", types[count], size_strings[size_to_test],
                    index, results_back[count*sizes[size_to_test]+index], expected);
        }
      }
    }
  }

  return total_errors;
}

int test_parameter_types(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
  clMemWrapper results;
  int error;
  size_t global[3] = {1, 1, 1};
  float results_back[7*16];
  int count, index;
  const char* types[] = {"char", "uchar", "short", "ushort", "int", "uint", "float"};
  char kernel_string[8192];
  int sizes[] = {1, 2, 4, 8, 16};
  const char* size_strings[] = {"", "2", "4", "8", "16"};
  float expected;
  int total_errors = 0;
  int size_to_test;
  char *ptr;
  char convert_string[1024];
  size_t max_parameter_size;

  // We don't really care about the contents since we're just testing that the types work.
  cl_char c[16]={0,-1,2,-3,4,-5,6,-7,8,-9,10,-11,12,-13,14,-15};
  cl_uchar uc[16]={16,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15};
  cl_short s[16]={-17,-1,2,-3,4,-5,6,-7,8,-9,10,-11,12,-13,14,-15};
  cl_ushort us[16]={18,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15};
  cl_int i[16]={-19,-1,2,-3,4,-5,6,-7,8,-9,10,-11,12,-13,14,-15};
  cl_uint ui[16]={20,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15};
  cl_float f[16]={-23,-1,2,-3,4,-5,6,-7,8,-9,10,-11,12,-13,14,-15};

  // Calculate how large our paramter size is to the kernel
  size_t parameter_size = sizeof(cl_char) + sizeof(cl_uchar) +
  sizeof(cl_short) +sizeof(cl_ushort) +
  sizeof(cl_int) +sizeof(cl_uint) +
  sizeof(cl_float);

  // Init our strings.
  kernel_string[0] = '\0';
  convert_string[0] = '\0';

  // Get the maximum parameter size allowed
  error = clGetDeviceInfo( device, CL_DEVICE_MAX_PARAMETER_SIZE, sizeof( max_parameter_size ), &max_parameter_size, NULL );
    test_error( error, "Unable to get max parameter size from device" );

  // Create the results buffer
  results = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_float)*7*16, NULL, &error);
  test_error(error, "clCreateBuffer failed");

  // Go over all the vector sizes
  for (size_to_test = 0; size_to_test < 5; size_to_test++) {
    clProgramWrapper program;
    clKernelWrapper kernel;

    size_t total_parameter_size = parameter_size*sizes[size_to_test] + sizeof(cl_mem);
    if (total_parameter_size > max_parameter_size) {
      log_info("Can not test with vector size %d because it would exceed the maximum allowed parameter size to the kernel. (%d > %d)\n",
               (int)sizes[size_to_test], (int)total_parameter_size, (int)max_parameter_size);
      continue;
    }

    log_info("Testing vector size %d\n", sizes[size_to_test]);

    // If size is > 1, then we need a explicit convert call.
    if (sizes[size_to_test] > 1) {
      sprintf(convert_string, "convert_float%s",  size_strings[size_to_test]);
    } else {
      sprintf(convert_string, " ");
    }

    // Build the kernel
    sprintf(kernel_string, kernel_code,
            size_strings[size_to_test], size_strings[size_to_test], size_strings[size_to_test],
            size_strings[size_to_test], size_strings[size_to_test], size_strings[size_to_test],
            size_strings[size_to_test], size_strings[size_to_test],
            convert_string, convert_string, convert_string,
            convert_string, convert_string, convert_string
    );

    ptr = kernel_string;
    error = create_single_kernel_helper(context, &program, &kernel, 1, (const char **)&ptr, "test_kernel");
    test_error(error, "create single kernel failed");

    // Set the arguments
    for (count = 0; count < 7; count++) {
      switch (count) {
        case 0: error = clSetKernelArg(kernel, count, sizeof(cl_char)*sizes[size_to_test], &c); break;
        case 1: error = clSetKernelArg(kernel, count, sizeof(cl_uchar)*sizes[size_to_test], &uc); break;
        case 2: error = clSetKernelArg(kernel, count, sizeof(cl_short)*sizes[size_to_test], &s); break;
        case 3: error = clSetKernelArg(kernel, count, sizeof(cl_ushort)*sizes[size_to_test], &us); break;
        case 4: error = clSetKernelArg(kernel, count, sizeof(cl_int)*sizes[size_to_test], &i); break;
        case 5: error = clSetKernelArg(kernel, count, sizeof(cl_uint)*sizes[size_to_test], &ui); break;
        case 6: error = clSetKernelArg(kernel, count, sizeof(cl_float)*sizes[size_to_test], &f); break;
        default: log_error("Test error"); break;
      }
      if (error)
        log_error("Setting kernel arg %d %s%s: ", count, types[count], size_strings[size_to_test]);
      test_error(error, "clSetKernelArgs failed");
    }
    error = clSetKernelArg(kernel, 7, sizeof(cl_mem), &results);
    test_error(error, "clSetKernelArgs failed");

    // Execute
    error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global, NULL, 0, NULL, NULL);
    test_error(error, "clEnqueueNDRangeKernel failed");

    // Read back the results
    error = clEnqueueReadBuffer(queue, results, CL_TRUE, 0, sizeof(cl_float)*7*16, results_back, 0, NULL, NULL);
    test_error(error, "clEnqueueReadBuffer failed");

    // Verify the results
    for (count = 0; count < 7; count++) {
      for (index=0; index < sizes[size_to_test]; index++) {
        switch (count) {
          case 0: expected = (float)c[index]; break;
          case 1: expected = (float)uc[index]; break;
          case 2: expected = (float)s[index]; break;
          case 3: expected = (float)us[index]; break;
          case 4: expected = (float)i[index]; break;
          case 5: expected = (float)ui[index]; break;
          case 6: expected = (float)f[index]; break;
          default: log_error("Test error"); break;
        }

        if (results_back[count*sizes[size_to_test]+index] != expected) {
          total_errors++;
          log_error("Conversion from %s%s failed: index %d got %g, expected %g.\n", types[count], size_strings[size_to_test],
                    index, results_back[count*sizes[size_to_test]+index], expected);
        }
      }
    }
  }

  if (gHasLong) {
    log_info("Testing long types...\n");
    total_errors += test_parameter_types_long( device, context, queue, num_elements );
  }
  else {
    log_info("Longs unsupported, skipping.");
  }

  return total_errors;
}