Skip to content
Permalink
Branch: master
Find file Copy path
Find file Copy path
Fetching contributors…
Cannot retrieve contributors at this time
87 lines (78 sloc) 3.07 KB
/***************************************************************************
*
* Copyright (C) 2016 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
*
* simple-vector-add.cpp
*
* Description:
* Example of a vector addition in SYCL.
*
**************************************************************************/
/* This example is a very small one designed to show how compact SYCL code
* can be. That said, it includes no error checking and is rather terse. */
#include <CL/sycl.hpp>
#include <array>
#include <iostream>
constexpr cl::sycl::access::mode sycl_read = cl::sycl::access::mode::read;
constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write;
/* This is the class used to name the kernel for the runtime.
* This must be done when the kernel is expressed as a lambda. */
template <typename T>
class SimpleVadd;
template <typename T, size_t N>
void simple_vadd(const std::array<T, N>& VA, const std::array<T, N>& VB,
std::array<T, N>& VC) {
cl::sycl::queue deviceQueue;
cl::sycl::range<1> numOfItems{N};
cl::sycl::buffer<T, 1> bufferA(VA.data(), numOfItems);
cl::sycl::buffer<T, 1> bufferB(VB.data(), numOfItems);
cl::sycl::buffer<T, 1> bufferC(VC.data(), numOfItems);
deviceQueue.submit([&](cl::sycl::handler& cgh) {
auto accessorA = bufferA.template get_access<sycl_read>(cgh);
auto accessorB = bufferB.template get_access<sycl_read>(cgh);
auto accessorC = bufferC.template get_access<sycl_write>(cgh);
auto kern = [=](cl::sycl::id<1> wiID) {
accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
};
cgh.parallel_for<class SimpleVadd<T>>(numOfItems, kern);
});
}
int main() {
const size_t array_size = 4;
std::array<cl::sycl::cl_int, array_size> A = {{1, 2, 3, 4}},
B = {{1, 2, 3, 4}}, C;
std::array<cl::sycl::cl_float, array_size> D = {{1.f, 2.f, 3.f, 4.f}},
E = {{1.f, 2.f, 3.f, 4.f}}, F;
simple_vadd(A, B, C);
simple_vadd(D, E, F);
for (unsigned int i = 0; i < array_size; i++) {
if (C[i] != A[i] + B[i]) {
std::cout << "The results are incorrect (element " << i << " is " << C[i]
<< "!\n";
return 1;
}
if (F[i] != D[i] + E[i]) {
std::cout << "The results are incorrect (element " << i << " is " << F[i]
<< "!\n";
return 1;
}
}
std::cout << "The results are correct!\n";
return 0;
}
You can’t perform that action at this time.