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

Commit f7036ed

Browse files
cperkinsintelcperkinsintel
authored
[SYCL] Update Sampler Linear Filter tests (#179)
* Linear Filtering is limited to floats for both accessor and underlying channel representation. This updates the tests to account for this. Signed-off-by: Chris Perkins <[email protected]>
1 parent 4219664 commit f7036ed

9 files changed

+619
-316
lines changed

‎SYCL/Sampler/basic-rw-float.cpp

Lines changed: 120 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,120 @@
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+
#include <CL/sycl.hpp>
7+
8+
using namespace cl::sycl;
9+
10+
using pixelT = sycl::float4;
11+
12+
// will output a pixel as {r,g,b,a}. provide override if a different pixelT is
13+
// defined.
14+
void outputPixel(sycl::float4 somePixel) {
15+
std::cout << "{" << somePixel[0] << "," << somePixel[1] << "," << somePixel[2]
16+
<< "," << somePixel[3] << "} ";
17+
}
18+
19+
// 4 pixels on a side. 1D at the moment
20+
constexpr long width = 4;
21+
22+
void test_rw(image_channel_order ChanOrder, image_channel_type ChanType) {
23+
int numTests = 4; // drives the size of the testResults buffer, and the number
24+
// of report iterations. Kludge.
25+
26+
// we'll use these four pixels for our image. Makes it easy to measure
27+
// interpolation and spot "off-by-one" probs.
28+
// These values will work consistently with different levels of float
29+
// precision (like unorm_int8 vs. fp32)
30+
pixelT leftEdge{0.2f, 0.4f, 0.6f, 0.8f};
31+
pixelT body{0.6f, 0.4f, 0.2f, 0.0f};
32+
pixelT bony{0.2f, 0.4f, 0.6f, 0.8f};
33+
pixelT rightEdge{0.6f, 0.4f, 0.2f, 0.0f};
34+
35+
queue Q;
36+
const sycl::range<1> ImgRange_1D(width);
37+
{ // closure
38+
// - create an image
39+
image<1> image_1D(ChanOrder, ChanType, ImgRange_1D);
40+
event E_Setup = Q.submit([&](handler &cgh) {
41+
auto image_acc = image_1D.get_access<pixelT, access::mode::write>(cgh);
42+
cgh.single_task<class setupUnormLinear>([=]() {
43+
image_acc.write(0, leftEdge);
44+
image_acc.write(1, body);
45+
image_acc.write(2, bony);
46+
image_acc.write(3, rightEdge);
47+
});
48+
});
49+
E_Setup.wait();
50+
51+
// use a buffer to report back test results.
52+
buffer<pixelT, 1> testResults((range<1>(numTests)));
53+
54+
event E_Test = Q.submit([&](handler &cgh) {
55+
auto image_acc = image_1D.get_access<pixelT, access::mode::read>(cgh);
56+
auto test_acc = testResults.get_access<access::mode::write>(cgh);
57+
58+
cgh.single_task<class im1D_Unorm_Linear>([=]() {
59+
int i = 0; // the index for writing into the testResult buffer.
60+
61+
// verify our four pixels were set up correctly.
62+
// 0-3 read four pixels. no sampler
63+
test_acc[i++] = image_acc.read(0); // {1,2,3,4}
64+
test_acc[i++] = image_acc.read(1); // {49,48,47,46}
65+
test_acc[i++] = image_acc.read(2); // {59,58,57,56}
66+
test_acc[i++] = image_acc.read(3); // {11,12,13,14}
67+
68+
// Add more tests below. Just be sure to increase the numTests counter
69+
// at the beginning of this function
70+
});
71+
});
72+
E_Test.wait();
73+
74+
// REPORT RESULTS
75+
auto test_acc = testResults.get_access<access::mode::read>();
76+
for (int i = 0, idx = 0; i < numTests; i++, idx++) {
77+
if (i == 0) {
78+
idx = 0;
79+
std::cout << "read four pixels, no sampler" << std::endl;
80+
}
81+
82+
pixelT testPixel = test_acc[i];
83+
std::cout << i << /* " -- " << idx << */ ": ";
84+
outputPixel(testPixel);
85+
std::cout << std::endl;
86+
}
87+
} // ~image / ~buffer
88+
}
89+
90+
int main() {
91+
92+
queue Q;
93+
device D = Q.get_device();
94+
95+
if (D.has(aspect::image)) {
96+
// the _int8 channels are one byte per channel, or four bytes per pixel (for
97+
// RGBA) the _int16/fp16 channels are two bytes per channel, or eight bytes
98+
// per pixel (for RGBA) the _int32/fp32 channels are four bytes per
99+
// channel, or sixteen bytes per pixel (for RGBA).
100+
101+
std::cout << "fp32 -------------" << std::endl;
102+
test_rw(image_channel_order::rgba, image_channel_type::fp32);
103+
104+
// CUDA, strangely, does not support 8-bit channels. Turning this off for
105+
// now.
106+
// std::cout << "unorm_int8 -------" << std::endl;
107+
// test_rw(image_channel_order::rgba, image_channel_type::unorm_int8);
108+
} else {
109+
std::cout << "device does not support image operations" << std::endl;
110+
}
111+
112+
return 0;
113+
}
114+
115+
// CHECK: fp32 -------------
116+
// CHECK-NEXT: read four pixels, no sampler
117+
// CHECK-NEXT: 0: {0.2,0.4,0.6,0.8}
118+
// CHECK-NEXT: 1: {0.6,0.4,0.2,0}
119+
// CHECK-NEXT: 2: {0.2,0.4,0.6,0.8}
120+
// CHECK-NEXT: 3: {0.6,0.4,0.2,0}

‎SYCL/Sampler/normalized-clamp-linear.cpp renamed to ‎SYCL/Sampler/normalized-clamp-linear-float.cpp

Lines changed: 65 additions & 43 deletions
Original file line numberDiff line numberDiff line change
@@ -2,13 +2,14 @@
22
// RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER
33
// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
44
// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
5-
// XFAIL: gpu && (level_zero || opencl || cuda)
6-
// XFAIL: cpu
75

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.
6+
// UNSUPPORTED: level_zero && windows
7+
// XFAIL: cuda
8+
9+
// LevelZero on Windows hangs with normalized coordinates. Waiting on fix.
10+
11+
// CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
12+
// type (such as unorm_int8)
1213

1314
/*
1415
This file sets up an image, initializes it with data,
@@ -23,11 +24,11 @@
2324
using namespace cl::sycl;
2425

2526
// pixel data-type for RGBA operations (which is the minimum image type)
26-
using pixelT = sycl::uint4;
27+
using pixelT = sycl::float4;
2728

2829
// will output a pixel as {r,g,b,a}. provide override if a different pixelT is
2930
// defined.
30-
void outputPixel(sycl::uint4 somePixel) {
31+
void outputPixel(sycl::float4 somePixel) {
3132
std::cout << "{" << somePixel[0] << "," << somePixel[1] << "," << somePixel[2]
3233
<< "," << somePixel[3] << "} ";
3334
}
@@ -47,10 +48,12 @@ void test_normalized_clamp_linear_sampler(image_channel_order ChanOrder,
4748

4849
// we'll use these four pixels for our image. Makes it easy to measure
4950
// 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};
51+
// These values will work consistently with different levels of float
52+
// precision (like unorm_int8 vs. fp32)
53+
pixelT leftEdge{0.2f, 0.4f, 0.6f, 0.8f};
54+
pixelT body{0.6f, 0.4f, 0.2f, 0.0f};
55+
pixelT bony{0.2f, 0.4f, 0.6f, 0.8f};
56+
pixelT rightEdge{0.6f, 0.4f, 0.2f, 0.0f};
5457

5558
queue Q;
5659
const sycl::range<1> ImgRange_1D(width);
@@ -96,31 +99,32 @@ void test_normalized_clamp_linear_sampler(image_channel_order ChanOrder,
9699
image_acc.read(-0.25f, Norm_Clamp_Linear_sampler); // {0,0,0,0}
97100
test_acc[i++] = image_acc.read(
98101
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}
102+
Norm_Clamp_Linear_sampler); // {0.1,0.2,0.3,0.4} // interpolating
103+
// with bg color. consistent with
104+
// unnormalized. Doesn't seem 100%
105+
// correct to me, but don't ahve
106+
// anything to compare against presnetly
110107
test_acc[i++] = image_acc.read(
111-
1.00f,
112-
Norm_Clamp_Linear_sampler); // {6,6,6,7} // interpolating with bg
108+
0.25f, Norm_Clamp_Linear_sampler); // {0.4,0.4,0.4,0.4}
109+
test_acc[i++] = image_acc.read(
110+
0.50f, Norm_Clamp_Linear_sampler); // {0.4,0.4,0.4,0.4}
111+
test_acc[i++] = image_acc.read(
112+
0.75f, Norm_Clamp_Linear_sampler); // {0.4,0.4,0.4,0.4}
113+
test_acc[i++] =
114+
image_acc.read(1.00f,
115+
Norm_Clamp_Linear_sampler); // {0.3,0.2,0.1,0} //
116+
// interpolating with bg
113117
test_acc[i++] =
114118
image_acc.read(1.25f, Norm_Clamp_Linear_sampler); // {0,0,0,0}
115119

116120
// 7-8 read two pixels on either side of first pixel. float coordinates.
117121
// CLAMP
118122
// on GPU CLAMP is apparently stopping the interpolation. ( values on
119123
// 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+
test_acc[i++] = image_acc.read(
125+
0.2499999f, Norm_Clamp_Linear_sampler); // {0.4,0.4,0.4,0.4}
126+
test_acc[i++] = image_acc.read(
127+
0.2500001f, Norm_Clamp_Linear_sampler); // {0.4,0.4,0.4,0.4}
124128
});
125129
});
126130
E_Test.wait();
@@ -130,9 +134,10 @@ void test_normalized_clamp_linear_sampler(image_channel_order ChanOrder,
130134
for (int i = 0, idx = 0; i < numTests; i++, idx++) {
131135
if (i == 0) {
132136
idx = -1;
133-
std::cout << "read six pixels at 'boundary' locations, starting out of "
134-
"bounds, sample: Normalized + Clamp + Linear"
135-
<< std::endl;
137+
std::cout
138+
<< "read seven pixels at 'boundary' locations, starting out of "
139+
"bounds, sample: Normalized + Clamp + Linear"
140+
<< std::endl;
136141
}
137142
if (i == 7) {
138143
idx = 1;
@@ -161,10 +166,14 @@ int main() {
161166
// RGBA) the _int16/fp16 channels are two bytes per channel, or eight bytes
162167
// per pixel (for RGBA) the _int32/fp32 channels are four bytes per
163168
// channel, or sixteen bytes per pixel (for RGBA).
164-
// CUDA has limited support for image_channel_type, so the tests use
165-
// unsigned_int32
169+
170+
std::cout << "fp32 -------------" << std::endl;
166171
test_normalized_clamp_linear_sampler(image_channel_order::rgba,
167-
image_channel_type::unsigned_int32);
172+
image_channel_type::fp32);
173+
174+
std::cout << "unorm_int8 -------" << std::endl;
175+
test_normalized_clamp_linear_sampler(image_channel_order::rgba,
176+
image_channel_type::unorm_int8);
168177
} else {
169178
std::cout << "device does not support image operations" << std::endl;
170179
}
@@ -173,15 +182,28 @@ int main() {
173182
}
174183

175184
// clang-format off
176-
// CHECK: read six pixels at 'boundary' locations, starting out of bounds, sample: Normalized + Clamp + Linear
185+
// CHECK: fp32 -------------
186+
// CHECK-NEXT: read seven pixels at 'boundary' locations, starting out of bounds, sample: Normalized + Clamp + Linear
187+
// CHECK-NEXT: 0 -- -1: {0,0,0,0}
188+
// CHECK-NEXT: 1 -- 0: {0.1,0.2,0.3,0.4}
189+
// CHECK-NEXT: 2 -- 1: {0.4,0.4,0.4,0.4}
190+
// CHECK-NEXT: 3 -- 2: {0.4,0.4,0.4,0.4}
191+
// CHECK-NEXT: 4 -- 3: {0.4,0.4,0.4,0.4}
192+
// CHECK-NEXT: 5 -- 4: {0.3,0.2,0.1,0}
193+
// CHECK-NEXT: 6 -- 5: {0,0,0,0}
194+
// CHECK-NEXT: read two pixels on either side of first pixel. float coordinates. Normalized + Clamp + Linear
195+
// CHECK-NEXT: 7 -- 1: {0.4,0.4,0.4,0.4}
196+
// CHECK-NEXT: 8 -- 1: {0.4,0.4,0.4,0.4}
197+
// CHECK-NEXT: unorm_int8 -------
198+
// CHECK-NEXT: read seven pixels at 'boundary' locations, starting out of bounds, sample: Normalized + Clamp + Linear
177199
// 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}
200+
// CHECK-NEXT: 1 -- 0: {0.1,0.2,0.3,0.4}
201+
// CHECK-NEXT: 2 -- 1: {0.4,0.4,0.4,0.4}
202+
// CHECK-NEXT: 3 -- 2: {0.4,0.4,0.4,0.4}
203+
// CHECK-NEXT: 4 -- 3: {0.4,0.4,0.4,0.4}
204+
// CHECK-NEXT: 5 -- 4: {0.3,0.2,0.1,0}
183205
// CHECK-NEXT: 6 -- 5: {0,0,0,0}
184206
// 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}
207+
// CHECK-NEXT: 7 -- 1: {0.4,0.4,0.4,0.4}
208+
// CHECK-NEXT: 8 -- 1: {0.4,0.4,0.4,0.4}
187209
// clang-format on

0 commit comments

Comments
 (0)