libceed-sys 0.12.0

Low-level bindings for libCEED library.
Documentation 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126

// Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other
// CEED contributors. All Rights Reserved. See the top-level LICENSE and NOTICE
// files for details.
//
// SPDX-License-Identifier: BSD-2-Clause
//
// This file is part of CEED:  http://github.com/ceed

#include <ceed/ceed.h>
#include <sycl/sycl.hpp>

//------------------------------------------------------------------------------
// Kernel for set value on device
//------------------------------------------------------------------------------
__global__ static void setValueK(CeedScalar *__restrict__ vec, CeedInt size, CeedScalar val) {
  int index = threadIdx.x + blockDim.x * blockIdx.x;

  if (index >= size) return;
  vec[index] = val;
}

//------------------------------------------------------------------------------
// Set value on device memory
//------------------------------------------------------------------------------
extern "C" int CeedDeviceSetValue_Sycl(CeedScalar *d_array, CeedInt length, CeedScalar val) {
  const int block_size = 512;
  const int vec_size   = length;
  int       grid_size  = vec_size / block_size;

  if (block_size * grid_size < vec_size) grid_size += 1;
  setValueK<<<grid_size, block_size>>>(d_array, length, val);
  return 0;
}

//------------------------------------------------------------------------------
// Kernel for taking reciprocal
//------------------------------------------------------------------------------
__global__ static void rcpValueK(CeedScalar *__restrict__ vec, CeedInt size) {
  int index = threadIdx.x + blockDim.x * blockIdx.x;

  if (index >= size) return;
  if (fabs(vec[index]) > 1E-16) vec[index] = 1. / vec[index];
}

//------------------------------------------------------------------------------
// Take vector reciprocal in device memory
//------------------------------------------------------------------------------
extern "C" int CeedDeviceReciprocal_Sycl(CeedScalar *d_array, CeedInt length) {
  const int block_size = 512;
  const int vec_size   = length;
  int       grid_size  = vec_size / block_size;

  if (block_size * grid_size < vec_size) grid_size += 1;
  rcpValueK<<<grid_size, block_size>>>(d_array, length);
  return 0;
}

//------------------------------------------------------------------------------
// Kernel for scale
//------------------------------------------------------------------------------
__global__ static void scaleValueK(CeedScalar *__restrict__ x, CeedScalar alpha, CeedInt size) {
  int index = threadIdx.x + blockDim.x * blockIdx.x;

  if (index >= size) return;
  x[index] *= alpha;
}

//------------------------------------------------------------------------------
// Compute x = alpha x on device
//------------------------------------------------------------------------------
extern "C" int CeedDeviceScale_Sycl(CeedScalar *x_array, CeedScalar alpha, CeedInt length) {
  const int block_size = 512;
  const int vec_size   = length;
  int       grid_size  = vec_size / block_size;

  if (block_size * grid_size < vec_size) grid_size += 1;
  scaleValueK<<<grid_size, block_size>>>(x_array, alpha, length);
  return 0;
}

//------------------------------------------------------------------------------
// Kernel for axpy
//------------------------------------------------------------------------------
__global__ static void axpyValueK(CeedScalar *__restrict__ y, CeedScalar alpha, CeedScalar *__restrict__ x, CeedInt size) {
  int index = threadIdx.x + blockDim.x * blockIdx.x;
  if (index >= size) return;
  y[index] += alpha * x[index];
}

//------------------------------------------------------------------------------
// Compute y = alpha x + y on device
//------------------------------------------------------------------------------
extern "C" int CeedDeviceAXPY_Sycl(CeedScalar *y_array, CeedScalar alpha, CeedScalar *x_array, CeedInt length) {
  const int block_size = 512;
  const int vec_size   = length;
  int       grid_size  = vec_size / block_size;

  if (block_size * grid_size < vec_size) grid_size += 1;
  axpyValueK<<<grid_size, block_size>>>(y_array, alpha, x_array, length);
  return 0;
}

//------------------------------------------------------------------------------
// Kernel for pointwise mult
//------------------------------------------------------------------------------
__global__ static void pointwiseMultValueK(CeedScalar *__restrict__ w, CeedScalar *x, CeedScalar *__restrict__ y, CeedInt size) {
  int index = threadIdx.x + blockDim.x * blockIdx.x;

  if (index >= size) return;
  w[index] = x[index] * y[index];
}

//------------------------------------------------------------------------------
// Compute the pointwise multiplication w = x .* y on device
//------------------------------------------------------------------------------
extern "C" int CeedDevicePointwiseMult_Sycl(CeedScalar *w_array, CeedScalar *x_array, CeedScalar *y_array, CeedInt length) {
  const int block_size = 512;
  const int vec_size   = length;
  int       grid_size  = vec_size / block_size;

  if (block_size * grid_size < vec_size) grid_size += 1;
  pointwiseMultValueK<<<grid_size, block_size>>>(w_array, x_array, y_array, length);
  return 0;
}

//------------------------------------------------------------------------------