Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 5e9e234

Browse files
cperkinsintelcperkinsintel
authored
Sampler Lit Tests (#152)
Between the coordinate_normalization_mode , addressing_mode and filtering_mode there are 16 valid combinations that a sampler should support for RGBA images. Here we are adding each combination as its own test. Some of the devices do not support particular combinations. For example, the CPU device is not correctly supporting linear interpolation at this time. In those cases, I have simply skipped the //RUN: directive for the not-yet-supporting device and left a comment in its place. Signed-off-by: Chris Perkins <[email protected]>
1 parent d565b8a commit 5e9e234

17 files changed

+2923
-0
lines changed

‎SYCL/Sampler/basic-rw.cpp

Lines changed: 131 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,131 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
2+
// RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
5+
6+
/*
7+
This file sets up an image, initializes it with data, and verifies that the
8+
data can be read directly.
9+
10+
Use it as a base file for testing any condition.
11+
12+
clang++ -fsycl -sycl-std=121 -o binx.bin basic-rw.cpp
13+
14+
SYCL_DEVICE_FILTER=opencl:gpu ./binx.bin
15+
SYCL_DEVICE_FILTER=level_zero:gpu ./binx.bin
16+
SYCL_DEVICE_FILTER=opencl:cpu ./binx.bin
17+
18+
SYCL_DEVICE_FILTER=opencl:host ./binx.bin
19+
SYCL_DEVICE_FILTER=opecl:acc ../binx.bin <-- does not support image
20+
operations at this time.
21+
22+
*/
23+
24+
#include <CL/sycl.hpp>
25+
26+
using namespace cl::sycl;
27+
28+
// pixel data-type for RGBA operations (which is the minimum image type)
29+
using pixelT = sycl::uint4;
30+
31+
// will output a pixel as {r,g,b,a}. provide override if a different pixelT is
32+
// defined.
33+
void outputPixel(sycl::uint4 somePixel) {
34+
std::cout << "{" << somePixel[0] << "," << somePixel[1] << "," << somePixel[2]
35+
<< "," << somePixel[3] << "} ";
36+
}
37+
38+
// 4 pixels on a side. 1D at the moment
39+
constexpr long width = 4;
40+
41+
void test_rw(image_channel_order ChanOrder, image_channel_type ChanType) {
42+
int numTests = 4; // drives the size of the testResults buffer, and the number
43+
// of report iterations. Kludge.
44+
45+
// we'll use these four pixels for our image. Makes it easy to measure
46+
// interpolation and spot "off-by-one" probs.
47+
pixelT leftEdge{1, 2, 3, 4};
48+
pixelT body{49, 48, 47, 46};
49+
pixelT bony{59, 58, 57, 56};
50+
pixelT rightEdge{11, 12, 13, 14};
51+
52+
queue Q;
53+
const sycl::range<1> ImgRange_1D(width);
54+
{ // closure
55+
// - create an image
56+
image<1> image_1D(ChanOrder, ChanType, ImgRange_1D);
57+
event E_Setup = Q.submit([&](handler &cgh) {
58+
auto image_acc = image_1D.get_access<pixelT, access::mode::write>(cgh);
59+
cgh.single_task<class setupUnormLinear>([=]() {
60+
image_acc.write(0, leftEdge);
61+
image_acc.write(1, body);
62+
image_acc.write(2, bony);
63+
image_acc.write(3, rightEdge);
64+
});
65+
});
66+
E_Setup.wait();
67+
68+
// use a buffer to report back test results.
69+
buffer<pixelT, 1> testResults((range<1>(numTests)));
70+
71+
event E_Test = Q.submit([&](handler &cgh) {
72+
auto image_acc = image_1D.get_access<pixelT, access::mode::read>(cgh);
73+
auto test_acc = testResults.get_access<access::mode::write>(cgh);
74+
75+
cgh.single_task<class im1D_Unorm_Linear>([=]() {
76+
int i = 0; // the index for writing into the testResult buffer.
77+
78+
// verify our four pixels were set up correctly.
79+
// 0-3 read four pixels. no sampler
80+
test_acc[i++] = image_acc.read(0); // {1,2,3,4}
81+
test_acc[i++] = image_acc.read(1); // {49,48,47,46}
82+
test_acc[i++] = image_acc.read(2); // {59,58,57,56}
83+
test_acc[i++] = image_acc.read(3); // {11,12,13,14}
84+
85+
// Add more tests below. Just be sure to increase the numTests counter
86+
// at the beginning of this function
87+
});
88+
});
89+
E_Test.wait();
90+
91+
// REPORT RESULTS
92+
auto test_acc = testResults.get_access<access::mode::read>();
93+
for (int i = 0, idx = 0; i < numTests; i++, idx++) {
94+
if (i == 0) {
95+
idx = 0;
96+
std::cout << "read four pixels, no sampler" << std::endl;
97+
}
98+
99+
pixelT testPixel = test_acc[i];
100+
std::cout << i << /* " -- " << idx << */ ": ";
101+
outputPixel(testPixel);
102+
std::cout << std::endl;
103+
}
104+
} // ~image / ~buffer
105+
}
106+
107+
int main() {
108+
109+
queue Q;
110+
device D = Q.get_device();
111+
112+
if (D.has(aspect::image)) {
113+
// the _int8 channels are one byte per channel, or four bytes per pixel (for
114+
// RGBA) the _int16/fp16 channels are two bytes per channel, or eight bytes
115+
// per pixel (for RGBA) the _int32/fp32 channels are four bytes per
116+
// channel, or sixteen bytes per pixel (for RGBA).
117+
// CUDA has limited support for image_channel_type, so the tests use
118+
// unsigned_int32
119+
test_rw(image_channel_order::rgba, image_channel_type::unsigned_int32);
120+
} else {
121+
std::cout << "device does not support image operations" << std::endl;
122+
}
123+
124+
return 0;
125+
}
126+
127+
// CHECK: read four pixels, no sampler
128+
// CHECK-NEXT: 0: {1,2,3,4}
129+
// CHECK-NEXT: 1: {49,48,47,46}
130+
// CHECK-NEXT: 2: {59,58,57,56}
131+
// CHECK-NEXT: 3: {11,12,13,14}

‎SYCL/Sampler/normalized-clamp-linear.cpp

Lines changed: 187 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,187 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
2+
// RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
5+
// XFAIL: gpu && (level_zero || opencl || cuda)
6+
// XFAIL: cpu
7+
8+
// GPU does not correctly interpolate when using clamp. Waiting on fix.
9+
// Both OCL and LevelZero have this issue.
10+
// CPU failing all linear interpolation at moment. Waiting on fix.
11+
// CUDA fails all linear interpolation. Waiting on fix.
12+
13+
/*
14+
This file sets up an image, initializes it with data,
15+
and verifies that the data is sampled correctly with a
16+
sampler configured NORMALIZED coordinate_normalization_mode
17+
CLAMP address_mode and LINEAR filter_mode
18+
19+
*/
20+
21+
#include <CL/sycl.hpp>
22+
23+
using namespace cl::sycl;
24+
25+
// pixel data-type for RGBA operations (which is the minimum image type)
26+
using pixelT = sycl::uint4;
27+
28+
// will output a pixel as {r,g,b,a}. provide override if a different pixelT is
29+
// defined.
30+
void outputPixel(sycl::uint4 somePixel) {
31+
std::cout << "{" << somePixel[0] << "," << somePixel[1] << "," << somePixel[2]
32+
<< "," << somePixel[3] << "} ";
33+
}
34+
35+
// some constants.
36+
37+
// 4 pixels on a side. 1D at the moment
38+
constexpr long width = 4;
39+
40+
constexpr auto normalized = coordinate_normalization_mode::normalized;
41+
constexpr auto linear = filtering_mode::linear;
42+
43+
void test_normalized_clamp_linear_sampler(image_channel_order ChanOrder,
44+
image_channel_type ChanType) {
45+
int numTests = 9; // drives the size of the testResults buffer, and the number
46+
// of report iterations. Kludge.
47+
48+
// we'll use these four pixels for our image. Makes it easy to measure
49+
// interpolation and spot "off-by-one" probs.
50+
pixelT leftEdge{1, 2, 3, 4};
51+
pixelT body{49, 48, 47, 46};
52+
pixelT bony{59, 58, 57, 56};
53+
pixelT rightEdge{11, 12, 13, 14};
54+
55+
queue Q;
56+
const sycl::range<1> ImgRange_1D(width);
57+
{ // closure
58+
// - create an image
59+
image<1> image_1D(ChanOrder, ChanType, ImgRange_1D);
60+
event E_Setup = Q.submit([&](handler &cgh) {
61+
auto image_acc = image_1D.get_access<pixelT, access::mode::write>(cgh);
62+
cgh.single_task<class setupUnormLinear>([=]() {
63+
image_acc.write(0, leftEdge);
64+
image_acc.write(1, body);
65+
image_acc.write(2, bony);
66+
image_acc.write(3, rightEdge);
67+
});
68+
});
69+
E_Setup.wait();
70+
71+
// use a buffer to report back test results.
72+
buffer<pixelT, 1> testResults((range<1>(numTests)));
73+
74+
// sampler
75+
auto Norm_Clamp_Linear_sampler =
76+
sampler(normalized, addressing_mode::clamp, linear);
77+
78+
event E_Test = Q.submit([&](handler &cgh) {
79+
auto image_acc = image_1D.get_access<pixelT, access::mode::read>(cgh);
80+
auto test_acc = testResults.get_access<access::mode::write>(cgh);
81+
82+
cgh.single_task<class im1D_norm_linear>([=]() {
83+
int i = 0; // the index for writing into the testResult buffer.
84+
85+
// clang-format off
86+
// Normalized Pixel Locations.
87+
// .125 .375 .625 .875 <-- exact center
88+
// |-----^-----|-----^-----|-----^-----|-----^-----
89+
//[0.0 .25 .50 .75 (1) <-- low boundary (included in pixel)
90+
// upper boundary inexact. (e.g. .2499999)
91+
// clang-format on
92+
93+
// 0-6 read seven pixels at 'boundary' locations, starting out of
94+
// bounds, sample: Normalized + Clamp + Linear
95+
test_acc[i++] =
96+
image_acc.read(-0.25f, Norm_Clamp_Linear_sampler); // {0,0,0,0}
97+
test_acc[i++] = image_acc.read(
98+
0.00f,
99+
Norm_Clamp_Linear_sampler); // {0,1,2,2} // interpolating with bg
100+
// color. consistent with unnormalized.
101+
// Doesn't seem 100% correct to me, but
102+
// don't ahve anything to compare
103+
// against presnetly
104+
test_acc[i++] =
105+
image_acc.read(0.25f, Norm_Clamp_Linear_sampler); // {25,25,25,25}
106+
test_acc[i++] =
107+
image_acc.read(0.50f, Norm_Clamp_Linear_sampler); // {54,53,52,51}
108+
test_acc[i++] =
109+
image_acc.read(0.75f, Norm_Clamp_Linear_sampler); // {35,35,35,35}
110+
test_acc[i++] = image_acc.read(
111+
1.00f,
112+
Norm_Clamp_Linear_sampler); // {6,6,6,7} // interpolating with bg
113+
test_acc[i++] =
114+
image_acc.read(1.25f, Norm_Clamp_Linear_sampler); // {0,0,0,0}
115+
116+
// 7-8 read two pixels on either side of first pixel. float coordinates.
117+
// CLAMP
118+
// on GPU CLAMP is apparently stopping the interpolation. ( values on
119+
// right are expected value)
120+
test_acc[i++] =
121+
image_acc.read(0.2499f, Norm_Clamp_Linear_sampler); // {25,25,25,25}
122+
test_acc[i++] =
123+
image_acc.read(0.2501f, Norm_Clamp_Linear_sampler); // {25,25,25,25}
124+
});
125+
});
126+
E_Test.wait();
127+
128+
// REPORT RESULTS
129+
auto test_acc = testResults.get_access<access::mode::read>();
130+
for (int i = 0, idx = 0; i < numTests; i++, idx++) {
131+
if (i == 0) {
132+
idx = -1;
133+
std::cout << "read six pixels at 'boundary' locations, starting out of "
134+
"bounds, sample: Normalized + Clamp + Linear"
135+
<< std::endl;
136+
}
137+
if (i == 7) {
138+
idx = 1;
139+
std::cout << "read two pixels on either side of first pixel. float "
140+
"coordinates. Normalized + Clamp + Linear"
141+
<< std::endl;
142+
}
143+
if (i == 8) {
144+
idx = 1;
145+
}
146+
pixelT testPixel = test_acc[i];
147+
std::cout << i << " -- " << idx << ": ";
148+
outputPixel(testPixel);
149+
std::cout << std::endl;
150+
}
151+
} // ~image / ~buffer
152+
}
153+
154+
int main() {
155+
156+
queue Q;
157+
device D = Q.get_device();
158+
159+
if (D.has(aspect::image)) {
160+
// the _int8 channels are one byte per channel, or four bytes per pixel (for
161+
// RGBA) the _int16/fp16 channels are two bytes per channel, or eight bytes
162+
// per pixel (for RGBA) the _int32/fp32 channels are four bytes per
163+
// channel, or sixteen bytes per pixel (for RGBA).
164+
// CUDA has limited support for image_channel_type, so the tests use
165+
// unsigned_int32
166+
test_normalized_clamp_linear_sampler(image_channel_order::rgba,
167+
image_channel_type::unsigned_int32);
168+
} else {
169+
std::cout << "device does not support image operations" << std::endl;
170+
}
171+
172+
return 0;
173+
}
174+
175+
// clang-format off
176+
// CHECK: read six pixels at 'boundary' locations, starting out of bounds, sample: Normalized + Clamp + Linear
177+
// CHECK-NEXT: 0 -- -1: {0,0,0,0}
178+
// CHECK-NEXT: 1 -- 0: {0,1,2,2}
179+
// CHECK-NEXT: 2 -- 1: {25,25,25,25}
180+
// CHECK-NEXT: 3 -- 2: {54,53,52,51}
181+
// CHECK-NEXT: 4 -- 3: {35,35,35,35}
182+
// CHECK-NEXT: 5 -- 4: {6,6,6,7}
183+
// CHECK-NEXT: 6 -- 5: {0,0,0,0}
184+
// CHECK-NEXT: read two pixels on either side of first pixel. float coordinates. Normalized + Clamp + Linear
185+
// CHECK-NEXT: 7 -- 1: {25,25,25,25}
186+
// CHECK-NEXT: 8 -- 1: {25,25,25,25}
187+
// clang-format on

0 commit comments

Comments
 (0)