//
// 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 "common.h"

const char *SVMPointerPassing_test_kernel[] = {
  "__kernel void verify_char(__global uchar* pChar, volatile __global uint* num_correct, uchar expected)\n"
  "{\n"
  "    if(0 == get_global_id(0))\n"
  "    {\n"
  "        *num_correct = 0;\n"
  "        if(*pChar == expected)\n"
  "        {\n"
  "                    *num_correct=1;\n"
  "        }\n"
  "    }\n"
  "}\n"
};


// Test that arbitrarily aligned char pointers into shared buffers can be passed directly to a kernel.
// This iterates through a buffer passing a pointer to each location to the kernel.
// The buffer is initialized to known values at each location.
// The kernel checks that it finds the expected value at each location.
// TODO: possibly make this work across all base types (including typeN?), also check ptr arithmetic ++,--.
int test_svm_pointer_passing(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements)
{
  clContextWrapper    context = NULL;
  clProgramWrapper    program = NULL;
  cl_uint     num_devices = 0;
  cl_int      error = CL_SUCCESS;
  clCommandQueueWrapper queues[MAXQ];

  error = create_cl_objects(deviceID, &SVMPointerPassing_test_kernel[0], &context, &program, &queues[0], &num_devices, CL_DEVICE_SVM_COARSE_GRAIN_BUFFER);
  if(error) return -1;

  clKernelWrapper kernel_verify_char = clCreateKernel(program, "verify_char", &error);
  test_error(error,"clCreateKernel failed");

  size_t bufSize = 256;
  cl_uchar *pbuf_svm_alloc = (cl_uchar*) clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_uchar)*bufSize, 0);

  cl_int *pNumCorrect = NULL;
  pNumCorrect = (cl_int*) clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_int), 0);

  {
    clMemWrapper buf = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, sizeof(cl_uchar)*bufSize, pbuf_svm_alloc, &error);
    test_error(error, "clCreateBuffer failed.");

    clMemWrapper num_correct = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, sizeof(cl_int), pNumCorrect, &error);
    test_error(error, "clCreateBuffer failed.");

    error = clSetKernelArg(kernel_verify_char, 1, sizeof(void*), (void *) &num_correct);
    test_error(error, "clSetKernelArg failed");

    // put values into buf so that we can expect to see these values in the kernel when we pass a pointer to them.
    cl_command_queue cmdq = queues[0];
    cl_uchar* pbuf_map_buffer = (cl_uchar*) clEnqueueMapBuffer(cmdq, buf, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_uchar)*bufSize, 0, NULL,NULL, &error);
    test_error2(error, pbuf_map_buffer, "clEnqueueMapBuffer failed");
    for(int i = 0; i<(int)bufSize; i++)
    {
      pbuf_map_buffer[i]= (cl_uchar)i;
    }
    error = clEnqueueUnmapMemObject(cmdq, buf, pbuf_map_buffer, 0,NULL,NULL);
    test_error(error, "clEnqueueUnmapMemObject failed.");

    for (cl_uint ii = 0; ii<num_devices; ++ii)  // iterate over all devices in the platform.
    {
      cmdq = queues[ii];
      for(int i = 0; i<(int)bufSize; i++)
      {
        cl_uchar* pChar = &pbuf_svm_alloc[i];
        error = clSetKernelArgSVMPointer(kernel_verify_char, 0, pChar); // pass a pointer to a location within the buffer
        test_error(error, "clSetKernelArg failed");
        error = clSetKernelArg(kernel_verify_char, 2, sizeof(cl_uchar), (void *) &i );  // pass the expected value at the above location.
        test_error(error, "clSetKernelArg failed");
        error = clEnqueueNDRangeKernel(cmdq, kernel_verify_char, 1, NULL, &bufSize, NULL, 0, NULL, NULL);
        test_error(error,"clEnqueueNDRangeKernel failed");

        pNumCorrect = (cl_int*) clEnqueueMapBuffer(cmdq, num_correct, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_int), 0, NULL,NULL, &error);
        test_error2(error, pNumCorrect, "clEnqueueMapBuffer failed");
        cl_int correct_count = *pNumCorrect;
        error = clEnqueueUnmapMemObject(cmdq, num_correct, pNumCorrect, 0,NULL,NULL);
        test_error(error, "clEnqueueUnmapMemObject failed.");

        if(correct_count != 1)
        {
          log_error("Passing pointer directly to kernel for byte #%d failed on device %d\n", i, ii);
          return -1;
        }
      }
    }

    error = clFinish(cmdq);
    test_error(error, "clFinish failed");
  }


  clSVMFree(context, pbuf_svm_alloc);
  clSVMFree(context, pNumCorrect);

  return 0;
}