ngscopeclient · Nov 26, 2021 · Nov 26, 2021
Showing with 102 additions and 10 deletions.

+1 −1 src/glscopeclient/WaveformArea.cpp

+1 −1 src/glscopeclient/WaveformArea_rendering.cpp

+96 −4 src/glscopeclient/kernels/WaveformRendering.cl

+4 −4 src/glscopeclient/shaders/waveform-compute-core.glsl
diff --git a/src/glscopeclient/WaveformArea.cpp b/src/glscopeclient/WaveformArea.cpp
@@ -770,7 +770,7 @@ void WaveformArea::InitializeWaveformPass()
 			cl::Program::Sources sources(1, make_pair(&kernelSource[0], kernelSource.length()));
 			m_renderProgram = new cl::Program(*g_clContext, sources);
 			m_renderProgram->build(g_contextDevices);
-			m_renderDenseAnalogWaveformKernel = new cl::Kernel(*m_renderProgram, "RenderAnalogWaveform");
+			m_renderDenseAnalogWaveformKernel = new cl::Kernel(*m_renderProgram, "RenderDensePackedAnalogWaveform");
 		}
 	}
 	catch(const cl::Error& e)

diff --git a/src/glscopeclient/WaveformArea_rendering.cpp b/src/glscopeclient/WaveformArea_rendering.cpp
@@ -675,7 +675,7 @@ void WaveformArea::RenderTrace(WaveformRenderData* data)
 				//TODO: Figure out indexing
 
 				//Number of threads to use for a single column of pixels
-				const int threads_per_column = 64;
+				const int threads_per_column = 1;//64;
 
 				//Grab a few helpful variables
 				auto group = data->m_area->m_group;

diff --git a/src/glscopeclient/kernels/WaveformRendering.cl b/src/glscopeclient/kernels/WaveformRendering.cl
@@ -34,6 +34,14 @@
 //Number of threads per column of pixels
 #define ROWS_PER_BLOCK	64
 
+/**
+	@brief Linearly interpolate Y coordinates
+ */
+float InterpolateY(float leftx, float lefty, float slope, float x)
+{
+	return lefty + ( (x-leftx)*slope );
+}
+
 /**
 	@brief Render a dense-packed analog waveform
 
@@ -55,7 +63,7 @@
 	@param ypos				Y coordinates of output waveform
 	@param outbuf			Output waveform buffer
  */
-__kernel void RenderAnalogWaveform(
+__kernel void RenderDensePackedAnalogWaveform(
 	unsigned int plotRight,
 	unsigned int width,
 	unsigned int height,
@@ -81,7 +89,13 @@ __kernel void RenderAnalogWaveform(
 	__local int blockmin[ROWS_PER_BLOCK];
 	__local int blockmax[ROWS_PER_BLOCK];
 	__local bool done;
-	__local bool updating[ROWS_PER_BLOCK];
+	//__local bool updating[ROWS_PER_BLOCK];
+
+	//Abort if invalid parameters
+	if(height > MAX_HEIGHT)
+		return;
+	if(depth < 2)
+		return;
 
 	//Don't do anything if we're off the right end of the buffer
 	unsigned long x = get_global_id(0) + firstcol;
@@ -101,7 +115,85 @@ __kernel void RenderAnalogWaveform(
 		return;
 	}
 
-	//Demo fill
+	//Clear working buffer
+	//TODO: persistence
+	for(unsigned int y=0; y<height; y++)
+		workingBuffer[y] = 0;
+
+	//Main loop setup
+	done = false;
+	unsigned long istart = floor(x / xscale) + offsetSamples;
+	unsigned long iend = floor((x + 1)/xscale) + offsetSamples;
+	if(iend <= 0)
+		done = true;
+
+	//Main loop
+	unsigned long i = istart;
+	while(!done)
+	{
+		bool updating = false;
+		if(i < (depth - 1))
+		{
+			//Get raw coordinates
+			float leftx = (i + innerXoff) * xscale + xoff;
+			float lefty = (ypos[i] + yoff) * yscale + ybase;
+			float rightx = (i + 1 + innerXoff) * xscale + xoff;
+			float righty = (ypos[i+1] + yoff) * yscale + ybase;
+
+			//Skip offscreen
+			if( (rightx >= x) && (leftx <= x+1) )
+			{
+				float starty = lefty;
+				float endy = righty;
+
+				//Interpolate
+				/*
+				float slope = (righty - lefty) / (rightx - leftx);
+				if(leftx < x)
+					starty = InterpolateY(leftx, lefty, slope, x);
+				else
+					endy = InterpolateY(leftx, lefty, slope, x+1);
+				*/
+
+				//Clip to window size
+				starty = min(starty, (float)(MAX_HEIGHT-1));
+				endy = min(endy, (float)(MAX_HEIGHT-1));
+				starty = max(starty, 0.0f);
+				endy = max(endy, 0.0f);
+
+				//Sort coordinates
+				blockmin[0] = min(starty, endy);
+				blockmax[0] = max(starty, endy);
+
+				//At end of pixel? Stop
+				if(rightx > x+1)
+					done = true;
+
+				updating = true;
+			}
+
+			else
+			{
+				//nothing to do
+			}
+		}
+		else
+			done = true;
+
+		//TODO: multiple rows per block
+		i ++;
+
+		//Update the images
+		if(updating)
+		{
+			int ymin = blockmin[0];
+			int ymax = blockmax[0];
+			for(int y=ymin; y<=ymax; y++)
+				workingBuffer[y] += alpha;
+		}
+	}
+
+	//Copy working buffer to output
 	for(unsigned long y=0; y<height; y++)
-		outbuf[y*width + x] = (float)x / width;
+		outbuf[y*width + x] = workingBuffer[y];
 }
diff --git a/src/glscopeclient/shaders/waveform-compute-core.glsl b/src/glscopeclient/shaders/waveform-compute-core.glsl
@@ -132,8 +132,10 @@ void main()
 				#endif
 
 				//Clip to window size
-				starty = min(starty, MAX_HEIGHT-1);
-				endy = min(endy, MAX_HEIGHT-1);
+				starty = min(starty, MAX_HEIGHT - 1);
+				endy = min(endy, MAX_HEIGHT - 1);
+				starty = max(starty, 0);
+				endy = max(endy, 0);
 
 				//Sort Y coordinates from min to max
 				g_blockmin[gl_LocalInvocationID.y] = int(min(starty, endy));
@@ -165,8 +167,6 @@ void main()
 			{
 				//Parallel fill
 				int ymin = g_blockmin[y];
-				if(ymin < 0)
-					ymin = 0;
 				int len = g_blockmax[y] - ymin;
 				for(uint y=gl_LocalInvocationID.y; y <= len; y += ROWS_PER_BLOCK)
 				{