在 OpenCL 内核 (android) 中读取 GL_UNSIGNED_BYTE OpenGL texture2D
Posted
技术标签:
【中文标题】在 OpenCL 内核 (android) 中读取 GL_UNSIGNED_BYTE OpenGL texture2D【英文标题】：Reading GL_UNSIGNED_BYTE OpenGL texture2D in OpenCL kernel (android) 【发布时间】：2014-06-22 21:09:19 【问题描述】：我的 android 应用程序将 OpenGL texture2D 传递给我的 OpenCL 内核，但是我的内核读取的像素值超出范围 (>255)。
我像这样创建我的 OpenGL 纹理：
GLES20.glGenTextures ( 2, targetTex, 0 );
GLES20.glBindTexture(GLES20.GL_TEXTURE_2D, targetTex[0]);
GLES20.glTexParameteri(GLES20.GL_TEXTURE_2D, GLES20.GL_TEXTURE_MIN_FILTER, GLES20.GL_LINEAR);
GLES20.glTexParameteri(GLES20.GL_TEXTURE_2D, GLES20.GL_TEXTURE_MAG_FILTER, GLES20.GL_LINEAR);
GLES20.glTexImage2D(GLES20.GL_TEXTURE_2D, 0, GLES20.GL_RGBA, image_width, image_height, 0, GLES20.GL_RGBA, GLES20.GL_UNSIGNED_BYTE, null);
GLES20.glBindTexture(GLES20.GL_TEXTURE_2D, 0);
然后通过将纹理与 FBO 绑定来渲染纹理：
targetFramebuffer = IntBuffer.allocate(1);
GLES20.glGenFramebuffers(1, targetFramebuffer);
GLES20.glBindFramebuffer(GLES20.GL_FRAMEBUFFER, targetFramebuffer.get(0));
GLES20.glFramebufferTexture2D(GLES20.GL_FRAMEBUFFER, GLES20.GL_COLOR_ATTACHMENT0, GLES20.GL_TEXTURE_2D, targetTex[0], 0);
GLES20.glBindFramebuffer(GLES20.GL_FRAMEBUFFER, 0);
我这样创建 cl 内存对象：
mem_images[0] = clCreateFromGLTexture2D(m_clContext, CL_MEM_READ_ONLY, GL_TEXTURE_2D, 0, in_tex, &err);
这是我的 OpenCL 内核：
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
kernel void transfer_data(__read_only image2d_t input_image, __global float* debug)
int2 pos;
uint4 pixel;
for (pos.y = get_global_id(1); pos.y < HEIGHT; pos.y += get_global_size(1))
for (pos.x = get_global_id(0); pos.x < WIDTH; pos.x += get_global_size(0))
pixel = read_imageui(input_image, sampler, pos);
debug[(pos.x + pos.y*WIDTH)*NUM_CHANNELS + 0] = pixel.x;
debug[(pos.x + pos.y*WIDTH)*NUM_CHANNELS + 1] = pixel.y;
debug[(pos.x + pos.y*WIDTH)*NUM_CHANNELS + 2] = pixel.z;
这就是我对内核进行排队的方式：
local2Dsize[0] = 4;
local2Dsize[1] = 4;
global2Dsize[0] = clamp(image_width, 0, max_work_items[0]);
global2Dsize[1] = clamp(image_height, 0, max_work_items[1]);
global2Dsize[0] = ceil((float)global2Dsize[0]/(float)local2Dsize[0])*local2Dsize[0];
global2Dsize[1] = ceil((float)global2Dsize[1]/(float)local2Dsize[1])*local2Dsize[1];
twoDlocal_sizes["transfer_data"] = local2Dsize;
twoDglobal_sizes["transfer_data"] = global2Dsize;
kernels["transfer_data"] = clCreateKernel(m_program, "transfer_data", &err);
err = clSetKernelArg(kernels["transfer_data"], 0, sizeof(cl_mem), &mem_images[0]);
err |= clSetKernelArg(kernels["transfer_data"], 1, sizeof(cl_mem), &mems["image"]);
err = clEnqueueAcquireGLObjects(m_queue, 1, &mem_images[0], 0, 0, 0);
CHECK_ERROR_OCL(err, "acquiring GL objects", return false);
err = clEnqueueNDRangeKernel(m_queue, kernels["transfer_data"], 2, NULL, twoDglobal_sizes["transfer_data"], twoDlocal_sizes["transfer_data"], 0, NULL, NULL);
err = clFinish(m_queue);
err = clEnqueueReleaseGLObjects(m_queue, 1, &mem_images[0], 0, 0, 0);
CHECK_ERROR_OCL(err, "releasing GL objects", return false);
当我打印出这些像素值（来自数组调试）时，现在回到主机，它们超出了界限，我不明白为什么会这样。
如果您需要更多见解：
我的项目的总体目标是获取 OpenGL 纹理形式的相机帧，使用 OpenCL 处理它们并将输出渲染回屏幕。 然而，从 android 相机获取的纹理只能绑定到 GL_TEXTURE_EXTERNAL_OES (http://developer.android.com/reference/android/graphics/SurfaceTexture.html)，这不是创建 OpenCL 内存对象的有效纹理。 因此，我将相机输出渲染到 GL_TEXTURE_2D 并将其传递给 OpenCL。
我确信像素被正确渲染到纹理，因为当我在屏幕上显示纹理时（不涉及任何 OpenCL），它会正确显示图像。
我通过创建纹理（而不是从相机获取数据）并将其传递给 opencl 进行了一些调试。所以这些是我得到的映射：
0 -> 0
1 -> 7172
2 -> 8196
3 -> 8710
4 -> 9220
5 -> 9477
6 -> 9734
7 -> 9991
8 -> 10244
9 -> 10372
10 -> 10501
11 -> 10629
12 -> 10758
13 -> 10886
14 -> 11015
15 -> 11143
16 -> 11268
17 -> 11332
18 -> 11396
19 -> 11460
20 -> 11525
21 -> 11589
22 -> 11653
23 -> 11717
24 -> 11782
25 -> 11846
26 -> 11910
27 -> 11974
28 -> 12039
29 -> 12103
30 -> 12167
31 -> 12231
32 -> 12292
33 -> 12324
34 -> 12356
35 -> 12388
36 -> 12420
37 -> 12452
38 -> 12484
39 -> 12516
40 -> 12549
41 -> 12581
42 -> 12613
43 -> 12645
44 -> 12677
45 -> 12709
46 -> 12741
47 -> 12773
48 -> 12806
49 -> 12838
50 -> 12870
51 -> 12902
52 -> 12934
53 -> 12966
54 -> 12998
55 -> 13030
56 -> 13063
57 -> 13095
58 -> 13127
59 -> 13159
60 -> 13191
61 -> 13223
62 -> 13255
63 -> 13287
64 -> 13316
65 -> 13332
66 -> 13348
67 -> 13364
68 -> 13380
69 -> 13396
70 -> 13412
71 -> 13428
72 -> 13444
73 -> 13460
74 -> 13476
75 -> 13492
76 -> 13508
77 -> 13524
78 -> 13540
79 -> 13556
80 -> 13573
81 -> 13589
82 -> 13605
83 -> 13621
84 -> 13637
85 -> 13653
86 -> 13669
87 -> 13685
88 -> 13701
89 -> 13717
90 -> 13733
91 -> 13749
92 -> 13765
93 -> 13781
94 -> 13797
95 -> 13813
96 -> 13830
97 -> 13846
98 -> 13862
99 -> 13878
100 -> 13894
101 -> 13910
102 -> 13926
103 -> 13942
104 -> 13958
105 -> 13974
106 -> 13990
107 -> 14006
108 -> 14022
109 -> 14038
110 -> 14054
111 -> 14070
112 -> 14087
113 -> 14103
114 -> 14119
115 -> 14135
116 -> 14151
117 -> 14167
118 -> 14183
119 -> 14199
120 -> 14215
121 -> 14231
122 -> 14247
123 -> 14263
124 -> 14279
125 -> 14295
126 -> 14311
127 -> 14327
128 -> 14340
129 -> 14348
130 -> 14356
131 -> 14364
132 -> 14372
133 -> 14380
134 -> 14388
135 -> 14396
136 -> 14404
137 -> 14412
138 -> 14420
139 -> 14428
140 -> 14436
141 -> 14444
142 -> 14452
143 -> 14460
144 -> 14468
145 -> 14476
146 -> 14484
147 -> 14492
148 -> 14500
149 -> 14508
150 -> 14516
151 -> 14524
152 -> 14532
153 -> 14540
154 -> 14548
155 -> 14556
156 -> 14564
157 -> 14572
158 -> 14580
159 -> 14588
160 -> 14597
161 -> 14605
162 -> 14613
163 -> 14621
164 -> 14629
165 -> 14637
166 -> 14645
167 -> 14653
168 -> 14661
169 -> 14669
170 -> 14677
171 -> 14685
172 -> 14693
173 -> 14701
174 -> 14709
175 -> 14717
176 -> 14725
177 -> 14733
178 -> 14741
179 -> 14749
180 -> 14757
181 -> 14765
182 -> 14773
183 -> 14781
184 -> 14789
185 -> 14797
186 -> 14805
187 -> 14813
188 -> 14821
189 -> 14829
190 -> 14837
191 -> 14845
192 -> 14854
193 -> 14862
194 -> 14870
195 -> 14878
196 -> 14886
197 -> 14894
198 -> 14902
199 -> 14910
200 -> 14918
201 -> 14926
202 -> 14934
203 -> 14942
204 -> 14950
205 -> 14958
206 -> 14966
207 -> 14974
208 -> 14982
209 -> 14990
210 -> 14998
211 -> 15006
212 -> 15014
213 -> 15022
214 -> 15030
215 -> 15038
216 -> 15046
217 -> 15054
218 -> 15062
219 -> 15070
220 -> 15078
221 -> 15086
222 -> 15094
223 -> 15102
224 -> 15111
225 -> 15119
226 -> 15127
227 -> 15135
228 -> 15143
229 -> 15151
230 -> 15159
231 -> 15167
232 -> 15175
233 -> 15183
234 -> 15191
235 -> 15199
236 -> 15207
237 -> 15215
238 -> 15223
239 -> 15231
240 -> 15239
241 -> 15247
242 -> 15255
243 -> 15263
244 -> 15271
245 -> 15279
246 -> 15287
247 -> 15295
248 -> 15303
249 -> 15311
250 -> 15319
251 -> 15327
252 -> 15335
253 -> 15343
254 -> 15351
255 -> 15359
左边是我在 OpenGL 纹理中输入的颜色值，左边是我在 OpenCL 中读取值时得到的对应值。
【问题讨论】：
您是否记得在调用内核之前调用 glFinish 和 clEnqueueAcquireGLObjects？另外宽度和宽度有什么区别？ 是的，我在获取对象之前调用了 glFinish。和宽度和宽度是相同的。我将编辑问题。 在内核和 clFinish 之后你是否也调用了 release GL 对象？只是检查。您必须调用内核的代码也很有用。 是的，我也在调用 clFinish。我刚刚包含了内核执行的过程。 @amirhbp 你说你在 cmets 中调用它们，但我在你的代码中看不到 clEnqueueAcquireGLObjects() 和 clEnqueueReleaseGLObjects()。否则，互操作性将永远无法发挥作用。 【参考方案1】：GL_UNSIGNED_BYTE 纹理作为CL_UNORM_INT8 映射到 OpenCL。您需要使用 read_imagef 而不是 read_imageui 来读取这些图像。使用read_imageui 时看到的值是内部浮点格式的原始位。
【讨论】：
【参考方案2】：在访问 GL 对象之前，您不会获取和释放它们。这会导致内核不读取 GL 内部缓冲区中的数据，而是读取它的本地 CL 副本。
正确的代码：（顺便说一句，您应该检查那些“err”值是否有错误）
local2Dsize[0] = 4;
local2Dsize[1] = 4;
global2Dsize[0] = clamp(image_width, 0, max_work_items[0]);
global2Dsize[1] = clamp(image_height, 0, max_work_items[1]);
global2Dsize[0] = ceil((float)global2Dsize[0]/(float)local2Dsize[0])*local2Dsize[0];
global2Dsize[1] = ceil((float)global2Dsize[1]/(float)local2Dsize[1])*local2Dsize[1];
twoDlocal_sizes["transfer_data"] = local2Dsize;
twoDglobal_sizes["transfer_data"] = global2Dsize;
kernels["transfer_data"] = clCreateKernel(m_program, "transfer_data", &err);
err = clSetKernelArg(kernels["transfer_data"], 0, sizeof(cl_mem), &mem_images[0]);
err |= clSetKernelArg(kernels["transfer_data"], 1, sizeof(cl_mem), &mems["image"]);
err = clEnqueueAcquireGLObjects(m_queue, 1, mem_images, NULL, NULL, NULL);
err = clEnqueueNDRangeKernel(m_queue, kernels["transfer_data"], 2, NULL, twoDglobal_sizes["transfer_data"], twoDlocal_sizes["transfer_data"], 0, NULL, NULL);
err = clEnqueueReleaseGLObjects (m_queue, 1, mem_images, NULL, NULL, NULL);
err = clFinish(m_queue);
【讨论】：
【参考方案3】：好的，我认为这是 OpenCL 实现中的一个错误。
在 Wolfram Alpha 的帮助下，我创建了一个函数来反转上述映射并获得 0 到 255 范围内的值。
float GL_to_CL(uint val)
if (val >= 14340) return round(0.1245790*val - 1658.44); //>=128
if (val >= 13316) return round(0.0622869*val - 765.408); //>=64
if (val >= 12292) return round(0.0311424*val - 350.800); //>=32
if (val >= 11268) return round(0.0155702*val - 159.443); //>=16
float v = (float) val;
return round(0.0000000000000125922*pow(v,4.f) - 0.00000000026729*pow(v,3.f) + 0.00000198135*pow(v,2.f) - 0.00496681*v - 0.0000808829);
所以 GL_to_CL() 是 4 个线性函数和一个四次函数的组合。 我尝试使用多项式插值仅创建一个函数，但是多项式的次数太大，因此比上面提出的 5 个函数的组合解决的计算成本更高。
另一个替代方法是使用 15k 大小的数组来实现恒定时间，但这需要我将大约 15k 字节上传到 GPU 的全局内存。 考虑到我正在使用内核进行图像处理，我会稍微推动它。 此外，从 OpenCL 中的全局内存访问通常比执行一些简单的计算更昂贵。
【讨论】：
我不会把它作为正确的答案，因为对我来说你在这里遗漏了一些东西。也许来自 GL 端的纹理具有不同的按位表示或字节序。破解的解决方案永远不应该是问题的正确答案。我想说的是，在创建纹理时 GL_UNSIGNED_BYTE 与 OpenCL 规范中的含义不同。 @DarkZeros，GL_UNSIGNED_BYTE 几乎是唯一一种可以在 OpenGL ES 2.0 中创建并由 OpenCL 读取的纹理（最接近的 OpenCL 对应关系是 CL_UNSIGNED_INT8）。但是，我承认绕过有缺陷的 OpenCL 实现是一种黑客行为，因此我不接受这个答案。但是，我会把它留在这里，以防有人遇到类似的问题并且只是在寻找有用的东西。以上是关于在 OpenCL 内核 (android) 中读取 GL_UNSIGNED_BYTE OpenGL texture2D的主要内容，如果未能解决你的问题，请参考以下文章
如何使用 PyOpenCL 将带有数组和变量的 C 结构传递给 OpenCL 内核