/***************************************************************************
 *
 *  Copyright (C) 2018 Codeplay Software Limited
 *  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
 *
 *  For your convenience, a copy of the License has been included in this
 *  repository.
 *
 *  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.
 *
 *  Codeplay's ComputeCpp SDK
 *
 *  reinterpret.cpp
 *
 *  Description:
 *    Sample code showing the reinterpret buffer feature of SYCL 1.2.1
 *
 **************************************************************************/

#include 

class mult;

int main() {
  cl::sycl::range r(128);
  cl::sycl::buffer buf_float(r);
  cl::sycl::queue q;

  {
    auto acc = buf_float.get_access();
    for (auto i = 0u; i < r.size(); i++) {
      acc[i] = i + 1;
    }
  }

  /* buf_int is a new SYCL buffer, with the same total size as buf_float,
   * but will provide uint32_t elements instead. However, the device memory
   * is the *same*.
   * This is actually true of the host memory, too, though there isn't any
   * host memory explicitly used here. In effect the reinterpret is a way
   * to obtain a differently-shaped handle over the top of SYCL's internal
   * data structures. */
  auto buf_int = buf_float.reinterpret(r);
  q.submit([&](cl::sycl::handler& cgh) {
    auto acc = buf_int.get_access(cgh);
    /* This kernel will multiply IEEE-754 32-bit floats by two, by manipulating
     * the exponent directly */
    cgh.parallel_for(r, [=](cl::sycl::item i) {
      constexpr auto mask = 0x7F800000u;
      constexpr auto mantissa_shift = 23u;
      auto& elem = acc[i.get_id()];
      auto exponent = (elem & mask) >> mantissa_shift;
      exponent++;
      elem &= ~mask;
      elem |= (exponent