Skip to content

Commit a8da48c

Browse files
smistadsmistad
committed
Added some Python OpenCL tests
1 parent 3886557 commit a8da48c

File tree

3 files changed

+163
-0
lines changed

3 files changed

+163
-0
lines changed

‎source/FAST/Python/ProcessObjects.i.in‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,7 @@
77
%feature("director") PythonRandomAccessStreamer;
88
%shared_ptr(fast::PythonRandomAccessStreamer)
99

10+
%template(setBufferArg) fast::Kernel::setArg<fast::OpenCLBuffer>;
1011
%template(setIntArg) fast::Kernel::setArg<int>;
1112
%template(setShortArg) fast::Kernel::setArg<short>;
1213
%template(setCharArg) fast::Kernel::setArg<char>;
@@ -332,6 +333,8 @@ def getQueue(self, device: "fast::OpenCLDevice::pointer"=None) -> "fast::Queue":
332333
data = list(mem)
333334
if isinstance(data, fast.Tensor):
334335
self.setTensorArg(id, data, fast.ACCESS_READ_WRITE)
336+
elif isinstance(data, fast.OpenCLBuffer):
337+
self.setBufferArg(id, data)
335338
elif isinstance(data, list):
336339
type = ''.join(char for char in arg.type if char.isalpha())
337340
if type == 'float':

‎source/FAST/Python/Tests/opencl_kernel.cl‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
__const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST | CLK_ADDRESS_NONE;
2+
__kernel void invert(
3+
__read_only image2d_t input,
4+
__write_only image2d_t output,
5+
__private uint max
6+
) {
7+
int2 pos = {get_global_id(0), get_global_id(1)};
8+
int value = read_imageui(input, sampler, pos).x;
9+
write_imageui(output, pos, max - value);
10+
}

‎source/FAST/Python/Tests/test_opencl.py‎

Lines changed: 150 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,150 @@
1+
import fast
2+
import pytest
3+
import numpy as np
4+
import os
5+
6+
7+
class OpenCLProcessObject(fast.PythonProcessObject):
8+
def __init__(self, useIndex=False, missingArgument=False, loadFromFile=False):
9+
super().__init__()
10+
self.createInputPort(0)
11+
self.createOutputPort(0)
12+
self.useIndex = useIndex
13+
self.missingArgument = missingArgument
14+
15+
# Create an image invert OpenCL kernel
16+
if loadFromFile:
17+
self.createOpenCLProgram('opencl_kernel.cl')
18+
else:
19+
self.createInlineOpenCLProgram(
20+
'''
21+
__const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST | CLK_ADDRESS_NONE;
22+
__kernel void invert(
23+
__read_only image2d_t input,
24+
__write_only image2d_t output,
25+
__private uint max
26+
) {
27+
int2 pos = {get_global_id(0), get_global_id(1)};
28+
int value = read_imageui(input, sampler, pos).x;
29+
write_imageui(output, pos, max - value);
30+
}
31+
'''
32+
)
33+
34+
def execute(self):
35+
# Get input image and create output image:
36+
image = self.getInputData()
37+
output = fast.Image.createFromImage(image)
38+
39+
# Get kernel from OpenCL program
40+
kernel = self.getKernel('invert')
41+
# Provide arguments to the kernel
42+
ID1 = 0 if self.useIndex else 'input'
43+
ID2 = 1 if self.useIndex else 'output'
44+
ID3 = 2 if self.useIndex else 'max'
45+
kernel.setArg(ID1, image)
46+
kernel.setArg(ID2, output)
47+
if not self.missingArgument:
48+
kernel.setArg(ID3, 255)
49+
50+
# Add the kernel to the command queue
51+
self.getQueue().add(kernel, image.getSize())
52+
53+
# Pass on the output image
54+
self.addOutputData(0, output)
55+
56+
57+
def test_simple_image_kernel():
58+
importer = fast.ImageFileImporter.create(fast.Config.getTestDataPath() + 'US/Heart/ApicalFourChamber/US-2D_0.mhd')
59+
inverter = fast.ImageInverter.create(min=0, max=255).connect(importer)
60+
result2 = inverter.runAndGetOutputData()
61+
62+
PO = OpenCLProcessObject.create(useIndex=False).connect(importer)
63+
result = PO.runAndGetOutputData()
64+
assert (np.asarray(result) == np.asarray(result2)).all()
65+
66+
PO = OpenCLProcessObject.create(useIndex=True).connect(importer)
67+
result = PO.runAndGetOutputData()
68+
assert (np.asarray(result) == np.asarray(result2)).all()
69+
70+
PO = OpenCLProcessObject.create(loadFromFile=True).connect(importer)
71+
result = PO.runAndGetOutputData()
72+
assert (np.asarray(result) == np.asarray(result2)).all()
73+
74+
75+
def test_missing_argument():
76+
importer = fast.ImageFileImporter.create(fast.Config.getTestDataPath() + 'US/Heart/ApicalFourChamber/US-2D_0.mhd')
77+
PO = OpenCLProcessObject.create(missingArgument=True).connect(importer)
78+
with pytest.raises(RuntimeError):
79+
result = PO.runAndGetOutputData()
80+
81+
82+
class OpenCLProcessObject2(fast.PythonProcessObject):
83+
def __init__(self, useNumpy=False, createBuffer=False):
84+
super().__init__()
85+
self.createInputPort(0)
86+
self.createOutputPort(0)
87+
88+
self.useNumpy = useNumpy
89+
self.useCustomBuffer = createBuffer
90+
self.data = list(range(255, -1, -1))
91+
92+
self.createInlineOpenCLProgram(
93+
'''
94+
__const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST | CLK_ADDRESS_NONE;
95+
__kernel void map(
96+
__read_only image2d_t input,
97+
__write_only image2d_t output,
98+
__global uint* data
99+
) {
100+
int2 pos = {get_global_id(0), get_global_id(1)};
101+
int value = read_imageui(input, sampler, pos).x;
102+
write_imageui(output, pos, data[value]);
103+
}
104+
'''
105+
)
106+
107+
def execute(self):
108+
# Get input image and create output image:
109+
image = self.getInputData()
110+
output = fast.Image.createFromImage(image)
111+
112+
# Get kernel from OpenCL program
113+
kernel = self.getKernel('map')
114+
115+
kernel.setArg('input', image)
116+
kernel.setArg('output', output)
117+
if self.useNumpy:
118+
kernel.setArg('data', np.array(self.data, dtype=np.uint8))
119+
elif self.useCustomBuffer:
120+
buffer = self.createUIntBuffer(self.data)
121+
kernel.setArg('data', buffer)
122+
else:
123+
kernel.setArg('data', self.data)
124+
125+
# Add the kernel to the command queue
126+
self.getQueue().add(kernel, image.getSize())
127+
128+
# Pass on the output image
129+
self.addOutputData(0, output)
130+
131+
132+
def test_image_kernel_with_buffer():
133+
importer = fast.ImageFileImporter.create(fast.Config.getTestDataPath() + 'US/Heart/ApicalFourChamber/US-2D_0.mhd')
134+
inverter = fast.ImageInverter.create(min=0, max=255).connect(importer)
135+
result2 = inverter.runAndGetOutputData()
136+
137+
PO = OpenCLProcessObject2.create().connect(importer)
138+
result = PO.runAndGetOutputData()
139+
assert (np.asarray(result) == np.asarray(result2)).all()
140+
141+
PO = OpenCLProcessObject2.create(useNumpy=True).connect(importer)
142+
result = PO.runAndGetOutputData()
143+
assert (np.asarray(result) == np.asarray(result2)).all()
144+
145+
PO = OpenCLProcessObject2.create(createBuffer=True).connect(importer)
146+
result = PO.runAndGetOutputData()
147+
assert (np.asarray(result) == np.asarray(result2)).all()
148+
149+
150+
# TODO Test: setImageArg, setTensorArg

0 commit comments

Comments
 (0)