@joej

What I have noticed and is strange, while rendering a cube I get 44FPS, but while rendering the cow, I get +100 FPS, how is that possible ?

@JoeJ

I have implemented the sorting by area method.

What I have seen is a cube is rendered at 5FPS, while the COW is at 145 FPS, no difference for the cow, but it's very slow for the cube.

Why is that ?

Also I see some flickering over the cow, some triangles are flickering.

@JoeJ

I'm currently working on that GPU, it's a custom based GPGPU, and I'm doing the rasterizer for it as hobby task.

They wrote on their research paper that they did it, so I wanted to take the challenge and do it on my own too ?

vortexgpgpu/vortex (github.com)

Regarding the sorting by area, I put a complex model, still it's slower than the plain scan line.

I will try the binning by nVidia research paper and will tell you the progress.

But Many thanks for your cooperation and your thoughts, you always come to rescue and give ideas to my work as usual.

@JoeJ

Hello,

hope you're doing fine. The basic scanline triangle works on the simulator of the GPU (PC Simulator) it takes 1 second for a 512*512 framebuffer, for a cube. and for cow it takes 13 seconds…

Would you think that tilemap rendering enhance the performance? If not, what would enhance the performance to overcome at least 20ms rendering ?

Also another issue,

Increasing the number of threads and warps, make the performance worth in the simulator… is there a solution for that to get benefits of multiple threads ?
-DNUM_CORES=1 -DNUM_WARPS=4 -DNUM_THREADS=1 -DL2_ENABLE=1 -DL3_ENABLE=1

for example using 1 thread, its very fast but no rendering comes out

Here is how I execute the kernel

global_size[0] = 800;

global_size[1] = 800;

auto time_start = std::chrono::high_resolution_clock::now();

CL_CHECK(clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL, global_size, NULL, 0, NULL, NULL));

here is the kernel:

#pragma OPENCL EXTENSION cl_amd_printf : enable                                


struct AABB {
	float min[3];
	float max[3];
};

struct  VertexIn {
	float pos[3];
	float nor[3];
	float col[3];
};

struct  VertexOut {
	float pos[3];
	float nor[3];
	float col[3];
	float mpos[3];
};
struct Triangle_ {
	struct VertexOut v[3];
	struct  AABB box;
	bool isPoint;
	bool isLine;
	bool isValidGeom;
	float signedArea;
	float minDepth;
};

struct Fragment {
	float pos[3];
	float nor[3];
	float col[3];
	bool isCovered;
};


//LOOK: checks if a barycentric coordinate is within the boundaries of a triangle
bool isBarycentricCoordInBounds(float3 barycentricCoord) {
	return barycentricCoord.x >= 0.0 && barycentricCoord.x <= 1.0 &&
		barycentricCoord.y >= 0.0 && barycentricCoord.y <= 1.0 &&
		barycentricCoord.z >= 0.0 && barycentricCoord.z <= 1.0;
}
float edgeFunction(float3 a, float3 b, float3 c)
{
	return (c.x - a.x) * (b.y - a.y) - (c.y - a.y) * (b.x - a.x);
}


__kernel void VertexShaderKernel(__global float* vbo, int vbosize
)
{
	size_t threadIdX = get_local_id(0);
	size_t threadIdY = get_local_id(1);
	size_t blockIdx = get_group_id(0);
	size_t blockIdy = get_group_id(1);
	size_t blockDimX = get_local_size(0);
	size_t blockDimY = get_local_size(1);


	int i = (blockIdx * blockDimY) + threadIdX;


}
 bool PointInTriangle(float3 p, float3 p0, float3 p1, float3 p2)
{
	float s = (p0.x - p2.x) * (p.y - p2.y) - (p0.y - p2.y) * (p.x - p2.x);
	float t = (p1.x - p0.x) * (p.y - p0.y) - (p1.y - p0.y) * (p.x - p0.x);

	if ((s < 0) != (t < 0) && s != 0 && t != 0)
		return false;

	float d = (p2.x - p1.x) * (p.y - p1.y) - (p2.y - p1.y) * (p.x - p1.x);
	return d == 0 || (d < 0) == (s + t <= 0);
}
 float min3(const float  a, const float  b, const float c)
 {
	 return fmin(a, fmin(b, c));
 }

 float max3(const float a, const float b, const float c)
 {
	 return fmax(a, fmax(b, c));
 }

__kernel void sendImageToPBO(__global uchar4* dst_buffer, __global struct Triangle_* triangles, int triCount)

{
	size_t blockIdx = get_group_id(0);
	size_t blockIdy = get_group_id(1);
	size_t blockDimX = get_local_size(0);
	size_t blockDimY = get_local_size(1);
	size_t threadIdX = get_local_id(0);
	size_t threadIdY = get_local_id(1);

	int xt = (blockIdx * blockDimX) + threadIdX;
	int yt = (blockIdy * blockDimY) + threadIdY;
    
	int imageWidth = 800;
	int imageHeight = 800;
	int index = xt + (yt * imageWidth);
    
	
	float3 c0 = { 1, 0, 0 };
	float3 c1 = { 0, 1, 0 };
	float3 c2 = { 0, 0, 1 };

	int x_pos = get_global_id(0);
	int y_pos = get_global_id(1);

	if (index < triCount)
	{
	

		float3 v0Raster = (float3)(triangles[index].v[0].pos[0], triangles[index].v[0].pos[1], 0);
		float3 v1Raster = (float3)(triangles[index].v[1].pos[0], triangles[index].v[1].pos[1], 0);
		float3 v2Raster = (float3)(triangles[index].v[2].pos[0], triangles[index].v[2].pos[1], 0);
		float xmin = min3(v0Raster.x, v1Raster.x, v2Raster.x);
		float ymin = min3(v0Raster.y, v1Raster.y, v2Raster.y);
		float xmax = max3(v0Raster.x, v1Raster.x, v2Raster.x);
		float ymax = max3(v0Raster.y, v1Raster.y, v2Raster.y);


		// the triangle is out of screen
		if (xmin < imageWidth - 1 || xmax > 0 || ymin < imageHeight - 1 || ymax > 0)
		{

			// be careful xmin/xmax/ymin/ymax can be negative. Don't cast to uint32_t
			unsigned int x0 = max(0, (int)(floor(xmin)));
			unsigned int x1 = min((int)(imageWidth)-1, (int)(floor(xmax)));
			unsigned int y0 = max(0, (int)(floor(ymin)));
			unsigned int y1 = min((int)(imageHeight)-1, (int)(floor(ymax)));
			for (unsigned int y = y0; y <= y1; ++y) {
				for (unsigned int x = x0; x <= x1; ++x) {
					float3 p = { x + 0.5f, y + 0.5f, 0 };

					float w0 = edgeFunction(v1Raster, v2Raster, p);
					float w1 = edgeFunction(v2Raster, v0Raster, p);
					float w2 = edgeFunction(v0Raster, v1Raster, p);

					if (w0 >= 0 && w1 >= 0 && w2 >= 0) {

						float area = edgeFunction(v0Raster, v1Raster, v2Raster);

						float r = w0 * c0.x + w1 * c1.x + w2 * c2.x;
						float g = w0 * c0.y + w1 * c1.y + w2 * c2.y;
						float b = w0 * c0.z + w1 * c1.z + w2 * c2.z;

						w0 /= area;
						w1 /= area;
						w2 /= area;
						float z = 1 / (w0 * v0Raster.z + w1 * v1Raster.z + w2 * v2Raster.z);
						r *= z, g *= z, b *= z;


						dst_buffer[y * get_global_size(0) + x] = (uchar4)(0,  255, 0, 255);
					}
				}
			}
		}
	}

}