36 #include <OpenCL/cl.h>
47 #include "dirichlet_kernel.cl.h"
60 "dirichlet_apply_scalar_kernel", &err);
63 CL_CHECK(clSetKernelArg(kernel, 0,
sizeof(cl_mem), (
void *) &msk));
64 CL_CHECK(clSetKernelArg(kernel, 1,
sizeof(cl_mem), (
void *) &
x));
66 CL_CHECK(clSetKernelArg(kernel, 3,
sizeof(
int), m));
68 const int nb = ((*m) + 256 - 1) / 256;
69 const size_t global_item_size = 256 * nb;
70 const size_t local_item_size = 256;
73 NULL, &global_item_size, &local_item_size,
81 void *z,
real *g,
int *m) {
88 "dirichlet_apply_vector_kernel", &err);
91 CL_CHECK(clSetKernelArg(kernel, 0,
sizeof(cl_mem), (
void *) &msk));
92 CL_CHECK(clSetKernelArg(kernel, 1,
sizeof(cl_mem), (
void *) &
x));
93 CL_CHECK(clSetKernelArg(kernel, 2,
sizeof(cl_mem), (
void *) &y));
94 CL_CHECK(clSetKernelArg(kernel, 3,
sizeof(cl_mem), (
void *) &z));
96 CL_CHECK(clSetKernelArg(kernel, 5,
sizeof(
int), m));
98 const int nb = ((*m) + 256 - 1) / 256;
99 const size_t global_item_size = 256 * nb;
100 const size_t local_item_size = 256;
103 NULL, &global_item_size, &local_item_size,
__global__ void const T *__restrict__ x
void opencl_dirichlet_apply_vector(void *msk, void *x, void *y, void *z, real *g, int *m)
void opencl_dirichlet_apply_scalar(void *msk, void *x, real *g, int *m)
void opencl_kernel_jit(const char *kernel, cl_program *program)