Skip to content

Commit 480bae9

Browse files
BeirdoBeirdo
committed
Add negative offset to X-correlate
1 parent df627c1 commit 480bae9

File tree

5 files changed

+88
-21
lines changed

5 files changed

+88
-21
lines changed

‎mythtv/programs/mythgpucommflag/openclvideoprocessor.cpp‎

Lines changed: 68 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,4 @@
1+
#include <QFile>
12
#include <QMap>
23
#include <strings.h>
34

@@ -1000,6 +1001,8 @@ void OpenCLHistogram64(OpenCLDevice *dev, VideoSurface *in, VideoHistogram *out)
10001001
}
10011002
}
10021003

1004+
VideoSurface binned(dev, kSurfaceRGB, in->m_width, in->m_height);
1005+
10031006
// for histogram64
10041007
LOG(VB_GPUVIDEO, LOG_INFO, QString("Histogram64: %1x%2 -> %3x%4")
10051008
.arg(totDims[0]) .arg(totDims[1]) .arg(reduceWorkDims[0])
@@ -1009,9 +1012,11 @@ void OpenCLHistogram64(OpenCLDevice *dev, VideoSurface *in, VideoHistogram *out)
10091012
&in->m_clBuffer[0]);
10101013
ciErrNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem),
10111014
&in->m_clBuffer[1]);
1012-
ciErrNum |= clSetKernelArg(kernel, 2, sizeof(cl_uint) * bins, NULL);
1013-
ciErrNum |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &memBufs->m_bufs[0]);
1014-
ciErrNum |= clSetKernelArg(kernel, 4, sizeof(cl_int2), totDims);
1015+
ciErrNum |= clSetKernelArg(kernel, 2, sizeof(cl_mem),
1016+
&binned.m_clBuffer[0]);
1017+
ciErrNum |= clSetKernelArg(kernel, 3, sizeof(cl_uint) * bins, NULL);
1018+
ciErrNum |= clSetKernelArg(kernel, 4, sizeof(cl_mem), &memBufs->m_bufs[0]);
1019+
ciErrNum |= clSetKernelArg(kernel, 5, sizeof(cl_int2), totDims);
10151020

10161021
if (ciErrNum != CL_SUCCESS)
10171022
{
@@ -1126,13 +1131,14 @@ void OpenCLHistogram64(OpenCLDevice *dev, VideoSurface *in, VideoHistogram *out)
11261131

11271132
static int dump = 2;
11281133

1129-
if (dump)
1134+
if (1 || dump)
11301135
{
11311136
float *outBins = new float[bins];
11321137

11331138
// Read back the results (finally!)
11341139
ciErrNum = clEnqueueReadBuffer(dev->m_commandQ, out->m_buf,
1135-
CL_TRUE, 0, bins * sizeof(cl_float),
1140+
CL_TRUE, 0,
1141+
bins * sizeof(cl_float),
11361142
outBins, 0, NULL, NULL);
11371143
if (ciErrNum != CL_SUCCESS)
11381144
{
@@ -1143,19 +1149,48 @@ void OpenCLHistogram64(OpenCLDevice *dev, VideoSurface *in, VideoHistogram *out)
11431149
return;
11441150
}
11451151

1152+
static int frameNum = 0;
1153+
1154+
#define DEBUG_VIDEO
1155+
#ifdef DEBUG_VIDEO
1156+
in->Dump("rgb", frameNum);
1157+
binned.Dump("binned", frameNum);
1158+
#undef DEBUG_VIDEO
1159+
#endif
1160+
1161+
#if 0
11461162
float tot = 0.0;
1163+
#endif
1164+
1165+
QString filename = QString("out/histogram-%1.gnuplot").arg(frameNum++);
1166+
QFile outfile(filename);
1167+
outfile.open(QIODevice::WriteOnly);
11471168

11481169
for (size_t i = 0; i < bins; i++)
11491170
{
1171+
#if 0
11501172
LOG(VB_GENERAL, LOG_INFO,
11511173
QString("Index %1 (%2-%3): Count = %4")
11521174
.arg(i) .arg(4 * i) .arg((4 * (i + 1)) - 1)
11531175
.arg(outBins[i]));
11541176
tot += outBins[i];
1177+
#endif
1178+
QString line = QString("%1 %2\n").arg(i) .arg(outBins[i]);
1179+
outfile.write(line.toLocal8Bit());
11551180
}
1181+
outfile.close();
11561182

1183+
#if 0
11571184
LOG(VB_GENERAL, LOG_INFO, QString("Total: Count = %1")
11581185
.arg(tot));
1186+
#endif
1187+
1188+
QFile accumfile("out/histogram0.gnuplot");
1189+
accumfile.open(QIODevice::Append);
1190+
QString line = QString("%1 %2\n").arg(frameNum-1) .arg(outBins[0]);
1191+
accumfile.write(line.toLocal8Bit());
1192+
accumfile.close();
1193+
11591194

11601195
delete [] outBins;
11611196
dump--;
@@ -1189,7 +1224,7 @@ void OpenCLCrossCorrelate(OpenCLDevice *dev, VideoHistogram *prev,
11891224
// Setup the kernel arguments
11901225
int bins = prev->m_binCount;
11911226

1192-
size_t globalWorkDim = bins;
1227+
size_t globalWorkDim = (2 * bins) - 1;
11931228

11941229
static OpenCLBufferPtr memBufs = NULL;
11951230

@@ -1203,7 +1238,7 @@ void OpenCLCrossCorrelate(OpenCLDevice *dev, VideoHistogram *prev,
12031238
}
12041239

12051240
memBufs->m_bufs[0] = clCreateBuffer(dev->m_context, CL_MEM_READ_WRITE,
1206-
sizeof(cl_float) * bins,
1241+
sizeof(cl_float) * globalWorkDim,
12071242
NULL, &ciErrNum);
12081243
if (ciErrNum != CL_SUCCESS)
12091244
{
@@ -1242,7 +1277,8 @@ void OpenCLCrossCorrelate(OpenCLDevice *dev, VideoHistogram *prev,
12421277

12431278
// Read back the results (finally!)
12441279
ciErrNum = clEnqueueReadBuffer(dev->m_commandQ, memBufs->m_bufs[0],
1245-
CL_TRUE, 0, bins * sizeof(cl_float),
1280+
CL_TRUE, 0,
1281+
globalWorkDim * sizeof(cl_float),
12461282
results, 0, NULL, NULL);
12471283
if (ciErrNum != CL_SUCCESS)
12481284
{
@@ -1254,14 +1290,34 @@ void OpenCLCrossCorrelate(OpenCLDevice *dev, VideoHistogram *prev,
12541290

12551291
static bool dumped = false;
12561292

1257-
if (!dumped)
1293+
if (1 || !dumped)
12581294
{
1259-
for (int i = 0; i < bins; i++)
1295+
static int frameNum = 1;
1296+
1297+
QString filename = QString("out/correlate-%1.gnuplot").arg(frameNum++);
1298+
QFile outfile(filename);
1299+
outfile.open(QIODevice::WriteOnly);
1300+
1301+
for (int i = 1 - bins; i < bins; i++)
12601302
{
1303+
#if 0
12611304
LOG(VB_GENERAL, LOG_INFO,
12621305
QString("Delay %1: Value = %4")
1263-
.arg(i) .arg(results[i]));
1306+
.arg(i) .arg(results[i + bins - 1]));
1307+
#endif
1308+
QString line = QString("%1 %2\n").arg(i) .arg(results[i + bins -1]);
1309+
outfile.write(line.toLocal8Bit());
1310+
12641311
}
1312+
1313+
outfile.close();
1314+
1315+
QFile accumfile("out/correlation0.gnuplot");
1316+
accumfile.open(QIODevice::Append);
1317+
QString line = QString("%1 %2\n").arg(frameNum-1) .arg(results[bins-1]);
1318+
accumfile.write(line.toLocal8Bit());
1319+
accumfile.close();
1320+
12651321
dumped = true;
12661322
}
12671323

@@ -1371,7 +1427,7 @@ FlagResults *OpenCLSceneChangeDetect(OpenCLDevice *dev, AVFrame *frame,
13711427
{
13721428
int bins = videoPacket->m_prevHistogram->m_binCount;
13731429

1374-
float *results = new float[bins];
1430+
float *results = new float[(2 * bins) - 1];
13751431

13761432
// Calculate the cross-correlation between previous and current
13771433
// color histograms

‎mythtv/programs/mythgpucommflag/videoHistogram.cl‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,7 @@
22
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
33
__kernel
44
void videoHistogram64(__read_only image2d_t fTop, __read_only image2d_t fBot,
5+
__write_only image2d_t fBin,
56
__local uint *histLoc, __global uint *histOut,
67
int2 total)
78
{
@@ -37,6 +38,11 @@ void videoHistogram64(__read_only image2d_t fTop, __read_only image2d_t fBot,
3738

3839
atom_inc( histLoc + indexT );
3940
atom_inc( histLoc + indexB );
41+
42+
uint4 binint = offsetT << 6;
43+
float4 binned = convert_float4(binint) / 255.0;
44+
binned.w = 1.0;
45+
write_imagef(fBin, (int2)(x, y), binned);
4046
}
4147

4248
barrier(CLK_LOCAL_MEM_FENCE);

‎mythtv/programs/mythgpucommflag/videoXCorrelate.cl‎

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -6,16 +6,19 @@ void videoCrossCorrelation(__global float *a, __global float *b,
66
__global float *out)
77
{
88
int l = get_global_id(0); // eg. 0-63
9-
int N = get_global_size(0); // eg. 64
9+
int N = get_global_size(0); // eg. 127
10+
int mid = ((N + 1) / 2) - 1; // eg. 63
11+
float numN = convert_float(mid + 1);
12+
float mean = 1.0 / numN;
1013

1114
// Do the a[n] * b[n - l] windowed for this value of l
1215
float locData = 0.0;
1316

14-
for (int n = l; n < N; n++)
17+
for (int n = max(l - mid, 0); n <= min(l, mid); n++)
1518
{
16-
float aVal = a[n];
17-
float bVal = b[n - l];
18-
locData += aVal * bVal;
19+
float aVal = a[n] - mean;
20+
float bVal = b[n - l + mid] - mean;
21+
locData += (aVal * bVal);
1922
}
2023

2124
out[l] = locData;

‎mythtv/programs/mythgpucommflag/videoconsumer.cpp‎

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -200,15 +200,15 @@ void VideoConsumer::ProcessPacket(Packet *packet)
200200
prevWavelet, prevHistogram);
201201
LOG(VB_GENERAL, LOG_INFO, "Finished transferring frame to OpenCL");
202202

203-
#define DEBUG_VIDEO
204203
#ifdef DEBUG_VIDEO
205-
if ((count <= 100) && videoFrame)
204+
if ((count == 1000 || count == 45000) && videoFrame)
206205
{
207206
videoFrame->m_frameRaw->Dump("frame", count);
208207
videoFrame->m_frameYUVSNORM->Dump("yuv", count);
209208
videoFrame->m_wavelet->Dump("wavelet", count);
210-
videoFrame->m_frameRGBDump("rgb", count);
209+
videoFrame->m_frameRGB->Dump("rgb", count);
211210

211+
#if 0
212212
VideoSurface yuv(m_dev, kSurfaceYUV2,
213213
videoFrame->m_frameRaw->m_width,
214214
videoFrame->m_frameRaw->m_height);
@@ -220,6 +220,7 @@ void VideoConsumer::ProcessPacket(Packet *packet)
220220
OpenCLYUVFromSNORM(m_dev, &yuv, &yuv2);
221221
OpenCLYUVToRGB(m_dev, &yuv2, &rgb);
222222
rgb.Dump("unwaveletRGB", count);
223+
#endif
223224
}
224225
#endif
225226
}

‎mythtv/programs/mythgpucommflag/videosurface.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -334,7 +334,8 @@ VideoHistogram::VideoHistogram(OpenCLDevice *dev, int binCount) :
334334
cl_int ciErrNum;
335335

336336
m_buf = clCreateBuffer(dev->m_context, CL_MEM_READ_WRITE,
337-
sizeof(cl_float) * m_binCount, NULL, &ciErrNum);
337+
sizeof(cl_float) * m_binCount,
338+
NULL, &ciErrNum);
338339
if (ciErrNum != CL_SUCCESS)
339340
{
340341
LOG(VB_GPU, LOG_ERR, QString("Error creating histogram: %1 (%2)")

0 commit comments

Comments
 (0)