Skip to content

Commit 9207570

Browse files
rok-cesnovarrok-cesnovar
authored
Merge pull request #2471 from bstatcomp/opencl_prim_gp
Add OpenCL prim gp_*_cov function
2 parents c0de977 + 2a92d6a commit 9207570

12 files changed

Lines changed: 1237 additions & 0 deletions

‎stan/math/opencl/kernels/gp_exponential_cov.hpp‎

Lines changed: 105 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,105 @@
1+
#ifndef STAN_MATH_OPENCL_KERNELS_GP_EXPONENTIAL_COV_HPP
2+
#define STAN_MATH_OPENCL_KERNELS_GP_EXPONENTIAL_COV_HPP
3+
#ifdef STAN_OPENCL
4+
5+
#include <stan/math/opencl/kernel_cl.hpp>
6+
#include <string>
7+
8+
namespace stan {
9+
namespace math {
10+
namespace opencl_kernels {
11+
// \cond
12+
static const std::string gp_exponential_cov_kernel_code = STRINGIFY(
13+
// \endcond
14+
/** \ingroup opencl_kernels
15+
* GPU part of calculation of Matern exponential kernel.
16+
*
17+
* @param[in] x input vector or matrix
18+
* @param[out] res squared distances between elements of x
19+
* @param sigma_sq squared standard deviation
20+
* @param neg_inv_l -1.0 / length_scale
21+
* @param size number of elements in x
22+
* @param element_size the number of doubles that make one element of x
23+
*/
24+
__kernel void gp_exponential_cov(
25+
const __global double* x, __global double* res, const double sigma_sq,
26+
const double neg_inv_l, const int size, const int element_size) {
27+
const int i = get_global_id(0);
28+
const int j = get_global_id(1);
29+
if (i < size && j < size) {
30+
if (i > j) {
31+
double sum = 0;
32+
for (int k = 0; k < element_size; k++) {
33+
double d = x[i * element_size + k] - x[j * element_size + k];
34+
sum += d * d;
35+
}
36+
double a = sigma_sq * exp(neg_inv_l * sqrt(sum));
37+
res[j * size + i] = a;
38+
res[i * size + j] = a;
39+
} else if (i == j) {
40+
res[j * size + i] = sigma_sq;
41+
}
42+
}
43+
}
44+
// \cond
45+
);
46+
// \endcond
47+
48+
/** \ingroup opencl_kernels
49+
* See the docs for \link kernels/gp_exponential_cov.hpp gp_exponential_cov()
50+
* \endlink
51+
*/
52+
const kernel_cl<in_buffer, out_buffer, double, double, int, int>
53+
gp_exponential_cov("gp_exponential_cov", {gp_exponential_cov_kernel_code});
54+
55+
// \cond
56+
static const std::string gp_exponential_cov_cross_kernel_code = STRINGIFY(
57+
// \endcond
58+
/** \ingroup opencl_kernels
59+
* GPU part of calculation of Matern exponential kernel.
60+
*
61+
* This function is for the cross covariance
62+
* matrix needed to compute the posterior predictive density.
63+
*
64+
* @param[in] x1 first input vector or matrix
65+
* @param[in] x2 second input vector or matrix
66+
* @param[out] res squared distances between elements of x
67+
* @param sigma_sq squared standard deviation
68+
* @param neg_inv_l -1.0 / length_scale
69+
* @param size1 number of elements in x1
70+
* @param size2 number of elements in x2
71+
* @param element_size the number of doubles that make one element of x and
72+
* y
73+
*/
74+
__kernel void gp_exponential_cov_cross(
75+
const __global double* x1, const __global double* x2,
76+
__global double* res, const double sigma_sq, const double neg_inv_l,
77+
const int size1, const int size2, const int element_size) {
78+
const int i = get_global_id(0);
79+
const int j = get_global_id(1);
80+
if (i < size1 && j < size2) {
81+
double sum = 0;
82+
for (int k = 0; k < element_size; k++) {
83+
double d = x1[i * element_size + k] - x2[j * element_size + k];
84+
sum += d * d;
85+
}
86+
res[j * size1 + i] = sigma_sq * exp(neg_inv_l * sqrt(sum));
87+
}
88+
}
89+
// \cond
90+
);
91+
// \endcond
92+
93+
/** \ingroup opencl_kernels
94+
* See the docs for \link kernels/gp_exponential_cov.hpp
95+
* gp_exponential_cov_cross() \endlink
96+
*/
97+
const kernel_cl<in_buffer, in_buffer, out_buffer, double, double, int, int, int>
98+
gp_exponential_cov_cross("gp_exponential_cov_cross",
99+
{gp_exponential_cov_cross_kernel_code});
100+
101+
} // namespace opencl_kernels
102+
} // namespace math
103+
} // namespace stan
104+
#endif
105+
#endif

‎stan/math/opencl/kernels/gp_matern32_cov.hpp‎

Lines changed: 108 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,108 @@
1+
#ifndef STAN_MATH_OPENCL_KERNELS_GP_MATERN32_COV_HPP
2+
#define STAN_MATH_OPENCL_KERNELS_GP_MATERN32_COV_HPP
3+
#ifdef STAN_OPENCL
4+
5+
#include <stan/math/opencl/kernel_cl.hpp>
6+
#include <string>
7+
8+
namespace stan {
9+
namespace math {
10+
namespace opencl_kernels {
11+
// \cond
12+
static const std::string gp_matern32_cov_kernel_code = STRINGIFY(
13+
// \endcond
14+
/** \ingroup opencl_kernels
15+
* GPU part of calculation of Matern 3/2 kernel.
16+
*
17+
* @param[in] x input vector or matrix
18+
* @param[out] res squared distances between elements of x
19+
* @param sigma_sq squared standard deviation
20+
* @param root_3_inv_l sqrt(3.0) / length_scale
21+
* @param size number of elements in x
22+
* @param element_size the number of doubles that make one element of x
23+
*/
24+
__kernel void gp_matern32_cov(
25+
const __global double* x, __global double* res, const double sigma_sq,
26+
const double root_3_inv_l, const int size, const int element_size) {
27+
const int i = get_global_id(0);
28+
const int j = get_global_id(1);
29+
if (i < size && j < size) {
30+
if (i > j) {
31+
double sum = 0;
32+
for (int k = 0; k < element_size; k++) {
33+
double d = x[i * element_size + k] - x[j * element_size + k];
34+
sum += d * d;
35+
}
36+
double dist = sqrt(sum);
37+
double a = sigma_sq * (1.0 + root_3_inv_l * dist)
38+
* exp(-root_3_inv_l * dist);
39+
res[j * size + i] = a;
40+
res[i * size + j] = a;
41+
} else if (i == j) {
42+
res[j * size + i] = sigma_sq;
43+
}
44+
}
45+
}
46+
// \cond
47+
);
48+
// \endcond
49+
50+
/** \ingroup opencl_kernels
51+
* See the docs for \link kernels/gp_matern32_cov.hpp gp_matern32_cov() \endlink
52+
*/
53+
const kernel_cl<in_buffer, out_buffer, double, double, int, int>
54+
gp_matern32_cov("gp_matern32_cov", {gp_matern32_cov_kernel_code});
55+
56+
// \cond
57+
static const std::string gp_matern32_cov_cross_kernel_code = STRINGIFY(
58+
// \endcond
59+
/** \ingroup opencl_kernels
60+
* GPU part of calculation of Matern 3/2 kernel.
61+
*
62+
* This function is for the cross covariance
63+
* matrix needed to compute the posterior predictive density.
64+
*
65+
* @param[in] x1 first input vector or matrix
66+
* @param[in] x2 second input vector or matrix
67+
* @param[out] res squared distances between elements of x
68+
* @param sigma_sq squared standard deviation
69+
* @param root_3_inv_l sqrt(3.0) / length_scale
70+
* @param size1 number of elements in x1
71+
* @param size2 number of elements in x2
72+
* @param element_size the number of doubles that make one element of x and
73+
* y
74+
*/
75+
__kernel void gp_matern32_cov_cross(
76+
const __global double* x1, const __global double* x2,
77+
__global double* res, const double sigma_sq, const double root_3_inv_l,
78+
const int size1, const int size2, const int element_size) {
79+
const int i = get_global_id(0);
80+
const int j = get_global_id(1);
81+
if (i < size1 && j < size2) {
82+
double sum = 0;
83+
for (int k = 0; k < element_size; k++) {
84+
double d = x1[i * element_size + k] - x2[j * element_size + k];
85+
sum += d * d;
86+
}
87+
double dist = sqrt(sum);
88+
res[j * size1 + i] = sigma_sq * (1.0 + root_3_inv_l * dist)
89+
* exp(-root_3_inv_l * dist);
90+
}
91+
}
92+
// \cond
93+
);
94+
// \endcond
95+
96+
/** \ingroup opencl_kernels
97+
* See the docs for \link kernels/gp_matern32_cov.hpp gp_matern32_cov_cross()
98+
* \endlink
99+
*/
100+
const kernel_cl<in_buffer, in_buffer, out_buffer, double, double, int, int, int>
101+
gp_matern32_cov_cross("gp_matern32_cov_cross",
102+
{gp_matern32_cov_cross_kernel_code});
103+
104+
} // namespace opencl_kernels
105+
} // namespace math
106+
} // namespace stan
107+
#endif
108+
#endif

‎stan/math/opencl/kernels/gp_matern52_cov.hpp‎

Lines changed: 114 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,114 @@
1+
#ifndef STAN_MATH_OPENCL_KERNELS_gp_MATERN52_COV_HPP
2+
#define STAN_MATH_OPENCL_KERNELS_gp_MATERN52_COV_HPP
3+
#ifdef STAN_OPENCL
4+
5+
#include <stan/math/opencl/kernel_cl.hpp>
6+
#include <string>
7+
8+
namespace stan {
9+
namespace math {
10+
namespace opencl_kernels {
11+
// \cond
12+
static const std::string gp_matern52_cov_kernel_code = STRINGIFY(
13+
// \endcond
14+
/** \ingroup opencl_kernels
15+
* GPU part of calculation of Matern 5/2 kernel.
16+
*
17+
* @param[in] x input vector or matrix
18+
* @param[out] res squared distances between elements of x
19+
* @param sigma_sq squared standard deviation
20+
* @param root_5_inv_l sqrt(5.0) / length_scale
21+
* @param inv_l_sq_5_3 5.0 / 3.0 / square(length_scale)
22+
* @param size number of elements in x
23+
* @param element_size the number of doubles that make one element of x
24+
*/
25+
__kernel void gp_matern52_cov(
26+
const __global double* x, __global double* res, const double sigma_sq,
27+
const double root_5_inv_l, const double inv_l_sq_5_3, const int size,
28+
const int element_size) {
29+
const int i = get_global_id(0);
30+
const int j = get_global_id(1);
31+
if (i < size && j < size) {
32+
if (i > j) {
33+
double sum = 0;
34+
for (int k = 0; k < element_size; k++) {
35+
double d = x[i * element_size + k] - x[j * element_size + k];
36+
sum += d * d;
37+
}
38+
double dist = sqrt(sum);
39+
double a = sigma_sq * (1.0 + root_5_inv_l * dist + inv_l_sq_5_3 * sum)
40+
* exp(-root_5_inv_l * dist);
41+
res[j * size + i] = a;
42+
res[i * size + j] = a;
43+
} else if (i == j) {
44+
res[j * size + i] = sigma_sq;
45+
}
46+
}
47+
}
48+
// \cond
49+
);
50+
// \endcond
51+
52+
/** \ingroup opencl_kernels
53+
* See the docs for \link kernels/gp_matern52_cov.hpp gp_matern52_cov() \endlink
54+
*/
55+
const kernel_cl<in_buffer, out_buffer, double, double, double, int, int>
56+
gp_matern52_cov("gp_matern52_cov", {gp_matern52_cov_kernel_code});
57+
58+
// \cond
59+
static const std::string gp_matern52_cov_cross_kernel_code = STRINGIFY(
60+
// \endcond
61+
/** \ingroup opencl_kernels
62+
* GPU part of calculation of Matern 5/2 kernel.
63+
*
64+
* This function is for the cross covariance
65+
* matrix needed to compute the posterior predictive density.
66+
*
67+
* @param[in] x1 first input vector or matrix
68+
* @param[in] x2 second input vector or matrix
69+
* @param[out] res squared distances between elements of x
70+
* @param sigma_sq squared standard deviation
71+
* @param root_5_inv_l sqrt(5.0) / length_scale
72+
* @param inv_l_sq_5_3 5.0 / 3.0 / square(length_scale)
73+
* @param size1 number of elements in x1
74+
* @param size2 number of elements in x2
75+
* @param element_size the number of doubles that make one element of x and
76+
* y
77+
*/
78+
__kernel void gp_matern52_cov_cross(
79+
const __global double* x1, const __global double* x2,
80+
__global double* res, const double sigma_sq, const double root_5_inv_l,
81+
const double inv_l_sq_5_3, const int size1, const int size2,
82+
const int element_size) {
83+
const int i = get_global_id(0);
84+
const int j = get_global_id(1);
85+
if (i < size1 && j < size2) {
86+
double sum = 0;
87+
for (int k = 0; k < element_size; k++) {
88+
double d = x1[i * element_size + k] - x2[j * element_size + k];
89+
sum += d * d;
90+
}
91+
double dist = sqrt(sum);
92+
res[j * size1 + i] = sigma_sq
93+
* (1.0 + root_5_inv_l * dist + inv_l_sq_5_3 * sum)
94+
* exp(-root_5_inv_l * dist);
95+
}
96+
}
97+
// \cond
98+
);
99+
// \endcond
100+
101+
/** \ingroup opencl_kernels
102+
* See the docs for \link kernels/gp_matern52_cov.hpp gp_matern52_cov_cross()
103+
* \endlink
104+
*/
105+
const kernel_cl<in_buffer, in_buffer, out_buffer, double, double, double, int,
106+
int, int>
107+
gp_matern52_cov_cross("gp_matern52_cov_cross",
108+
{gp_matern52_cov_cross_kernel_code});
109+
110+
} // namespace opencl_kernels
111+
} // namespace math
112+
} // namespace stan
113+
#endif
114+
#endif

‎stan/math/opencl/prim.hpp‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -151,7 +151,11 @@
151151
#include <stan/math/opencl/prim/frechet_lcdf.hpp>
152152
#include <stan/math/opencl/prim/frechet_lpdf.hpp>
153153
#include <stan/math/opencl/prim/gamma_lpdf.hpp>
154+
#include <stan/math/opencl/prim/gp_dot_prod_cov.hpp>
155+
#include <stan/math/opencl/prim/gp_exponential_cov.hpp>
154156
#include <stan/math/opencl/prim/gp_exp_quad_cov.hpp>
157+
#include <stan/math/opencl/prim/gp_matern32_cov.hpp>
158+
#include <stan/math/opencl/prim/gp_matern52_cov.hpp>
155159
#include <stan/math/opencl/prim/gumbel_cdf.hpp>
156160
#include <stan/math/opencl/prim/gumbel_lccdf.hpp>
157161
#include <stan/math/opencl/prim/gumbel_lcdf.hpp>

0 commit comments

Comments
 (0)