Skip to content

Commit 9fab4f3

Browse files
Jonas HahnfeldJonas Hahnfeld
Jonas Hahnfeld
authored and
Jonas Hahnfeld
committed
OpenCL: Add helper functions to enqueue MatVec kernels
1 parent 1ee9769 commit 9fab4f3

File tree

4 files changed

+57
-55
lines changed

4 files changed

+57
-55
lines changed

‎opencl/CGMultiOpenCL.cpp

+4-22
Original file line numberDiff line numberDiff line change
@@ -263,30 +263,12 @@ void CGMultiOpenCL::matvecKernel(Vector _x, Vector _y) {
263263

264264
switch (matrixFormat) {
265265
case MatrixFormatCRS:
266-
checkedSetKernelArg(matvecKernelCRS, 0, sizeof(cl_mem),
267-
&device.matrixCRS.ptr);
268-
checkedSetKernelArg(matvecKernelCRS, 1, sizeof(cl_mem),
269-
&device.matrixCRS.index);
270-
checkedSetKernelArg(matvecKernelCRS, 2, sizeof(cl_mem),
271-
&device.matrixCRS.value);
272-
checkedSetKernelArg(matvecKernelCRS, 3, sizeof(cl_mem), &x);
273-
checkedSetKernelArg(matvecKernelCRS, 4, sizeof(cl_mem), &y);
274-
checkedSetKernelArg(matvecKernelCRS, 5, sizeof(int), &yOffset);
275-
checkedSetKernelArg(matvecKernelCRS, 6, sizeof(int), &length);
276-
device.checkedEnqueueNDRangeKernel(matvecKernelCRS, device.globalMatvec);
266+
device.checkedEnqueueMatvecKernelCRS(matvecKernelCRS, device.matrixCRS, x,
267+
y, yOffset, length);
277268
break;
278269
case MatrixFormatELL:
279-
checkedSetKernelArg(matvecKernelELL, 0, sizeof(cl_mem),
280-
&device.matrixELL.length);
281-
checkedSetKernelArg(matvecKernelELL, 1, sizeof(cl_mem),
282-
&device.matrixELL.index);
283-
checkedSetKernelArg(matvecKernelELL, 2, sizeof(cl_mem),
284-
&device.matrixELL.data);
285-
checkedSetKernelArg(matvecKernelELL, 3, sizeof(cl_mem), &x);
286-
checkedSetKernelArg(matvecKernelELL, 4, sizeof(cl_mem), &y);
287-
checkedSetKernelArg(matvecKernelELL, 5, sizeof(int), &yOffset);
288-
checkedSetKernelArg(matvecKernelELL, 6, sizeof(int), &length);
289-
device.checkedEnqueueNDRangeKernel(matvecKernelELL, device.globalMatvec);
270+
device.checkedEnqueueMatvecKernelELL(matvecKernelELL, device.matrixELL, x,
271+
y, yOffset, length);
290272
break;
291273
default:
292274
assert(0 && "Invalid matrix format!");

‎opencl/CGOpenCL.cpp

+4-22
Original file line numberDiff line numberDiff line change
@@ -141,30 +141,12 @@ void CGOpenCL::matvecKernel(Vector _x, Vector _y) {
141141

142142
switch (matrixFormat) {
143143
case MatrixFormatCRS:
144-
checkedSetKernelArg(matvecKernelCRS, 0, sizeof(cl_mem),
145-
&device.matrixCRS.ptr);
146-
checkedSetKernelArg(matvecKernelCRS, 1, sizeof(cl_mem),
147-
&device.matrixCRS.index);
148-
checkedSetKernelArg(matvecKernelCRS, 2, sizeof(cl_mem),
149-
&device.matrixCRS.value);
150-
checkedSetKernelArg(matvecKernelCRS, 3, sizeof(cl_mem), &x);
151-
checkedSetKernelArg(matvecKernelCRS, 4, sizeof(cl_mem), &y);
152-
checkedSetKernelArg(matvecKernelCRS, 5, sizeof(int), &ZERO);
153-
checkedSetKernelArg(matvecKernelCRS, 6, sizeof(int), &N);
154-
device.checkedEnqueueNDRangeKernel(matvecKernelCRS, device.globalMatvec);
144+
device.checkedEnqueueMatvecKernelCRS(matvecKernelCRS, device.matrixCRS, x,
145+
y, ZERO, N);
155146
break;
156147
case MatrixFormatELL:
157-
checkedSetKernelArg(matvecKernelELL, 0, sizeof(cl_mem),
158-
&device.matrixELL.length);
159-
checkedSetKernelArg(matvecKernelELL, 1, sizeof(cl_mem),
160-
&device.matrixELL.index);
161-
checkedSetKernelArg(matvecKernelELL, 2, sizeof(cl_mem),
162-
&device.matrixELL.data);
163-
checkedSetKernelArg(matvecKernelELL, 3, sizeof(cl_mem), &x);
164-
checkedSetKernelArg(matvecKernelELL, 4, sizeof(cl_mem), &y);
165-
checkedSetKernelArg(matvecKernelELL, 5, sizeof(int), &ZERO);
166-
checkedSetKernelArg(matvecKernelELL, 6, sizeof(int), &N);
167-
device.checkedEnqueueNDRangeKernel(matvecKernelELL, device.globalMatvec);
148+
device.checkedEnqueueMatvecKernelELL(matvecKernelELL, device.matrixELL, x,
149+
y, ZERO, N);
168150
break;
169151
default:
170152
assert(0 && "Invalid matrix format!");

‎opencl/CGOpenCLBase.cpp

+40
Original file line numberDiff line numberDiff line change
@@ -45,6 +45,46 @@ std::vector<cl_device_id> CGOpenCLBase::getAllDevices() {
4545
return devices;
4646
}
4747

48+
void CGOpenCLBase::Device::checkedEnqueueNDRangeKernel(cl_kernel kernel,
49+
size_t global,
50+
size_t local) {
51+
if (global == 0) {
52+
global = this->global;
53+
}
54+
if (local == 0) {
55+
local = Local;
56+
}
57+
58+
checkError(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 0,
59+
NULL, NULL));
60+
}
61+
62+
void CGOpenCLBase::Device::checkedEnqueueMatvecKernelCRS(
63+
cl_kernel kernel, MatrixCRSDevice &deviceMatrix, cl_mem x, cl_mem y,
64+
int yOffset, int N) {
65+
checkedSetKernelArg(kernel, 0, sizeof(cl_mem), &deviceMatrix.ptr);
66+
checkedSetKernelArg(kernel, 1, sizeof(cl_mem), &deviceMatrix.index);
67+
checkedSetKernelArg(kernel, 2, sizeof(cl_mem), &deviceMatrix.value);
68+
checkedSetKernelArg(kernel, 3, sizeof(cl_mem), &x);
69+
checkedSetKernelArg(kernel, 4, sizeof(cl_mem), &y);
70+
checkedSetKernelArg(kernel, 5, sizeof(int), &yOffset);
71+
checkedSetKernelArg(kernel, 6, sizeof(int), &N);
72+
checkedEnqueueNDRangeKernel(kernel, globalMatvec);
73+
}
74+
75+
void CGOpenCLBase::Device::checkedEnqueueMatvecKernelELL(
76+
cl_kernel kernel, MatrixELLDevice &deviceMatrix, cl_mem x, cl_mem y,
77+
int yOffset, int N) {
78+
checkedSetKernelArg(kernel, 0, sizeof(cl_mem), &deviceMatrix.length);
79+
checkedSetKernelArg(kernel, 1, sizeof(cl_mem), &deviceMatrix.index);
80+
checkedSetKernelArg(kernel, 2, sizeof(cl_mem), &deviceMatrix.data);
81+
checkedSetKernelArg(kernel, 3, sizeof(cl_mem), &x);
82+
checkedSetKernelArg(kernel, 4, sizeof(cl_mem), &y);
83+
checkedSetKernelArg(kernel, 5, sizeof(int), &yOffset);
84+
checkedSetKernelArg(kernel, 6, sizeof(int), &N);
85+
checkedEnqueueNDRangeKernel(kernel, globalMatvec);
86+
}
87+
4888
cl_kernel CGOpenCLBase::checkedCreateKernel(const char *kernelName) {
4989
cl_int err;
5090
cl_kernel kernel = clCreateKernel(program, kernelName, &err);

‎opencl/CGOpenCLBase.h

+9-11
Original file line numberDiff line numberDiff line change
@@ -122,17 +122,15 @@ class CGOpenCLBase : public CG {
122122

123123
/// Enqueue \a kernel.
124124
void checkedEnqueueNDRangeKernel(cl_kernel kernel, size_t global = 0,
125-
size_t local = 0) {
126-
if (global == 0) {
127-
global = this->global;
128-
}
129-
if (local == 0) {
130-
local = Local;
131-
}
132-
133-
checkError(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local,
134-
0, NULL, NULL));
135-
}
125+
size_t local = 0);
126+
/// Enqueue \a kernel.
127+
void checkedEnqueueMatvecKernelCRS(cl_kernel kernel,
128+
MatrixCRSDevice &deviceMatrix, cl_mem x,
129+
cl_mem y, int yOffset, int N);
130+
/// Enqueue \a kernel.
131+
void checkedEnqueueMatvecKernelELL(cl_kernel kernel,
132+
MatrixELLDevice &deviceMatrix, cl_mem x,
133+
cl_mem y, int yOffset, int N);
136134

137135
/// Enqueue read of \a buffer.
138136
void checkedEnqueueReadBuffer(cl_mem buffer, size_t offset, size_t cb,

0 commit comments

Comments
 (0)