Made it human-friendly
This commit is contained in:
+145
-145
@@ -1,145 +1,145 @@
|
||||
#include "kernels.h"
|
||||
|
||||
const char* vertex_shader_kernel_source =
|
||||
"__kernel \n"
|
||||
"void vertex_shader( __global float* m, \n"
|
||||
" __global float* VertexBuffer, \n"
|
||||
" __global float* NewVertexBuffer) \n"
|
||||
"{ \n"
|
||||
" int local_index = get_local_id(0); \n"
|
||||
" int global_index = get_group_id(0); \n"
|
||||
" NewVertexBuffer[4*global_index+local_index] = \n"
|
||||
" m[local_index*4]*VertexBuffer[3*global_index] \n"
|
||||
" + m[local_index*4 + 1]*VertexBuffer[3*global_index+1] \n"
|
||||
" + m[local_index*4 + 2]*VertexBuffer[3*global_index+2] \n"
|
||||
" + m[local_index*4 + 3]; \n"
|
||||
"} \n";
|
||||
|
||||
const char* fragment_shader_kernel_source =
|
||||
"float3 barycentric(float3* pts, float3 P) \n"
|
||||
"{ \n"
|
||||
" float3 u = cross( \n"
|
||||
" (float3){pts[0][2] - pts[0][0], pts[0][1] - pts[0][0], pts[0][0] - P[0]}, // AC_x, AB_x, distance_x \n"
|
||||
" (float3){pts[1][2] - pts[1][0], pts[1][1] - pts[1][0], pts[1][0] - P[1]} // AC_y, AB_y, distance_y \n"
|
||||
" ); \n"
|
||||
" if (fabs(u[2]) < 1) return (float3){-1, 1, 1}; \n"
|
||||
" return (float3){1.f - (u[0] + u[1]) / u[2], u[1] / u[2], u[0] / u[2]}; \n"
|
||||
"} \n"
|
||||
" \n"
|
||||
"__kernel void fragment_shader ( \n"
|
||||
" __global int3* faces, \n"
|
||||
" __global float* vertices, \n"
|
||||
" __global int* pixels, \n"
|
||||
" __global int* screen_width, \n"
|
||||
" __global float* z_buffer, \n"
|
||||
" __global int* nfaces, \n"
|
||||
" __global float* uv_buffer, \n"
|
||||
" __global int* map_size, \n"
|
||||
" __global float* light_dir, \n"
|
||||
" __global float* norms_buff, \n"
|
||||
" __global uchar* diffuse_map \n"
|
||||
") { \n"
|
||||
" int GROUP_ID = get_group_id(0); \n"
|
||||
" int GROUP_SIZE = get_local_size(0); \n"
|
||||
" int LOCAL_ID = get_local_id(0); \n"
|
||||
" \n"
|
||||
" bool out = true; \n"
|
||||
" float3 vertices3[3]; \n"
|
||||
" float2 uv_coords[3]; \n"
|
||||
" float3 norms[3]; \n"
|
||||
" \n"
|
||||
" for(int i = 0; i < 3; i++) { \n"
|
||||
" float4 vertex;// = vertices[faces[GROUP_ID * 3 + i ][0]]; \n"
|
||||
" for(int j = 0; j < 4; j ++) { \n"
|
||||
" vertex[j] = vertices[4 * faces[GROUP_ID*3 + i][0] + j]; \n"
|
||||
" } \n"
|
||||
" \n"
|
||||
" for(int j = 0; j < 2; j++) { \n"
|
||||
" uv_coords[i][j] = uv_buffer [2 * faces[GROUP_ID * 3 + i][1] + j]; \n"
|
||||
" } \n"
|
||||
" \n"
|
||||
" for( int j = 0; j < 3; j++ ) { \n"
|
||||
" vertices3[j][i] = (vertex[j]/vertex[3]); \n"
|
||||
" norms[j][i] = norms_buff[3 * (faces[GROUP_ID * 3 + i][2]) + j]; \n"
|
||||
" } \n"
|
||||
" \n"
|
||||
" if ( vertices3[0][i] > 0 && vertices3[0][i] < *screen_width \n"
|
||||
" && vertices3[1][i] > 0 && vertices3[1][i] < *screen_width ) \n"
|
||||
" out = false; \n"
|
||||
" } \n"
|
||||
" \n"
|
||||
" if(out) return; \n"
|
||||
" \n"
|
||||
" //if(vertices3[1][0] == vertices3[1][1] && vertices3[1][2] == vertices3[1][1]) return; \n"
|
||||
" \n"
|
||||
" int2 bounding_box_min = (int2) { *screen_width - 1, *screen_width - 1 }; \n"
|
||||
" int2 bounding_box_max = (int2) { 0, 0 }; \n"
|
||||
" int2 clamper = (int2) { *screen_width - 1, *screen_width - 1 }; \n"
|
||||
" \n"
|
||||
" for(int i = 0; i < 3; i++) { \n"
|
||||
" for(int j = 0; j < 2; j++) { \n"
|
||||
" bounding_box_min[j] = max(0, min(bounding_box_min[j], (int)vertices3[j][i])); \n"
|
||||
" bounding_box_max[j] = min(clamper[j], max(bounding_box_max[j], (int)vertices3[j][i])); \n"
|
||||
" } \n"
|
||||
" } \n"
|
||||
" \n"
|
||||
" if(bounding_box_min[0] > *screen_width || bounding_box_max[0] < 0 || bounding_box_min[1] > *screen_width || bounding_box_max[1] < 0) return; \n"
|
||||
" \n"
|
||||
" \n"
|
||||
" int X_PER_ITEM = (int)(ceil((float)(bounding_box_max[0] - bounding_box_min[0]) / 16.f)); \n"
|
||||
" int Y_PER_ITEM = (int)(ceil((float)(bounding_box_max[1] - bounding_box_min[1]) / 16.f)); \n"
|
||||
" float STARTING_X = bounding_box_min[0] + X_PER_ITEM * (LOCAL_ID % 16); \n"
|
||||
" float ENDING_X = STARTING_X + X_PER_ITEM; \n"
|
||||
" float STARTING_Y = bounding_box_min[1] + Y_PER_ITEM * (LOCAL_ID / 16); \n"
|
||||
" float ENDING_Y = STARTING_Y + Y_PER_ITEM; \n"
|
||||
" \n"
|
||||
" \n"
|
||||
" float3 point; \n"
|
||||
" for(point[0] = STARTING_X; point[0] <= ENDING_X; point[0]++) { \n"
|
||||
" for(point[1] = STARTING_Y; point[1] <= ENDING_Y; point[1]++) { \n"
|
||||
" if(point[1] >= *screen_width || point[1] >= *screen_width) break; \n"
|
||||
" float3 bc_coord = barycentric(vertices3, point); \n"
|
||||
" if (bc_coord[0] < 0 || bc_coord[1] < 0 || bc_coord[2] < 0) continue; \n"
|
||||
" \n"
|
||||
" float2 uv_vec = (float2){0, 0}; \n"
|
||||
" float3 normal = (float3){0, 0, 0}; \n"
|
||||
" \n"
|
||||
" point[2] = dot(vertices3[2], bc_coord); \n"
|
||||
" if (z_buffer[(int)(point[0] + point[1] * *screen_width)] > point[2]) { \n"
|
||||
" continue; \n"
|
||||
" } \n"
|
||||
" \n"
|
||||
" \n"
|
||||
" for (int i = 0; i < 3; i++) { \n"
|
||||
" uv_vec[0] += uv_coords[i][0] * bc_coord[i]; \n"
|
||||
" uv_vec[1] += uv_coords[i][1] * bc_coord[i]; \n"
|
||||
" normal[i] = dot(norms[i], bc_coord); \n"
|
||||
" } \n"
|
||||
" int2 uv_point = (int2) { (int)(uv_vec[0] * map_size[0]), (int)(uv_vec[1] * map_size[1]) }; \n"
|
||||
" \n"
|
||||
" \n"
|
||||
" int col_index = 3 * (uv_point[0] + uv_point[1] * map_size[0]); \n"
|
||||
" \n"
|
||||
" \n"
|
||||
" \n"
|
||||
" \n"
|
||||
" float3 normalized_norm = normalize(normal); \n"
|
||||
" \n"
|
||||
" \n"
|
||||
" float intensity = clamp((dot(normalized_norm , (float3){light_dir[0], light_dir[1], light_dir[2]})), 0.f, 1.f) + 0.2; \n"
|
||||
" \n"
|
||||
" int color = 0; \n"
|
||||
" color |= ((int)fmin((float)(diffuse_map[col_index + 0]) * intensity, (float) 0xff)) << 16; \n"
|
||||
" color |= ((int)fmin((float)(diffuse_map[col_index + 1]) * intensity, (float) 0xff)) << 8; \n"
|
||||
" color |= ((int)fmin((float)(diffuse_map[col_index + 2]) * intensity, (float) 0xff)) << 0; \n"
|
||||
" \n"
|
||||
" //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 16; \n"
|
||||
" //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 8; \n"
|
||||
" //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 0; \n"
|
||||
" \n"
|
||||
" \n"
|
||||
" z_buffer[(int) (point[0] + point[1] * *screen_width)] = point[2]; \n"
|
||||
" pixels [(int) (point[0] + point[1] * *screen_width)] = color;// & 0x00ffffff; \n"
|
||||
" } \n"
|
||||
" } \n"
|
||||
"} \n";
|
||||
//#include "kernels.h"
|
||||
//
|
||||
//const char* vertex_shader_kernel_source =
|
||||
//"__kernel \n"
|
||||
//"void vertex_shader( __global float* m, \n"
|
||||
//" __global float* VertexBuffer, \n"
|
||||
//" __global float* NewVertexBuffer) \n"
|
||||
//"{ \n"
|
||||
//" int local_index = get_local_id(0); \n"
|
||||
//" int global_index = get_group_id(0); \n"
|
||||
//" NewVertexBuffer[4*global_index+local_index] = \n"
|
||||
//" m[local_index*4]*VertexBuffer[3*global_index] \n"
|
||||
//" + m[local_index*4 + 1]*VertexBuffer[3*global_index+1] \n"
|
||||
//" + m[local_index*4 + 2]*VertexBuffer[3*global_index+2] \n"
|
||||
//" + m[local_index*4 + 3]; \n"
|
||||
//"} \n";
|
||||
//
|
||||
//const char* fragment_shader_kernel_source =
|
||||
//"float3 barycentric(float3* pts, float3 P) \n"
|
||||
//"{ \n"
|
||||
//" float3 u = cross( \n"
|
||||
//" (float3){pts[0][2] - pts[0][0], pts[0][1] - pts[0][0], pts[0][0] - P[0]}, // AC_x, AB_x, distance_x \n"
|
||||
//" (float3){pts[1][2] - pts[1][0], pts[1][1] - pts[1][0], pts[1][0] - P[1]} // AC_y, AB_y, distance_y \n"
|
||||
//" ); \n"
|
||||
//" if (fabs(u[2]) < 1) return (float3){-1, 1, 1}; \n"
|
||||
//" return (float3){1.f - (u[0] + u[1]) / u[2], u[1] / u[2], u[0] / u[2]}; \n"
|
||||
//"} \n"
|
||||
//" \n"
|
||||
//"__kernel void fragment_shader ( \n"
|
||||
//" __global int3* faces, \n"
|
||||
//" __global float* vertices, \n"
|
||||
//" __global int* pixels, \n"
|
||||
//" __global int* screen_width, \n"
|
||||
//" __global float* z_buffer, \n"
|
||||
//" __global int* nfaces, \n"
|
||||
//" __global float* uv_buffer, \n"
|
||||
//" __global int* map_size, \n"
|
||||
//" __global float* light_dir, \n"
|
||||
//" __global float* norms_buff, \n"
|
||||
//" __global uchar* diffuse_map \n"
|
||||
//") { \n"
|
||||
//" int GROUP_ID = get_group_id(0); \n"
|
||||
//" int GROUP_SIZE = get_local_size(0); \n"
|
||||
//" int LOCAL_ID = get_local_id(0); \n"
|
||||
//" \n"
|
||||
//" bool out = true; \n"
|
||||
//" float3 vertices3[3]; \n"
|
||||
//" float2 uv_coords[3]; \n"
|
||||
//" float3 norms[3]; \n"
|
||||
//" \n"
|
||||
//" for(int i = 0; i < 3; i++) { \n"
|
||||
//" float4 vertex;// = vertices[faces[GROUP_ID * 3 + i ][0]]; \n"
|
||||
//" for(int j = 0; j < 4; j ++) { \n"
|
||||
//" vertex[j] = vertices[4 * faces[GROUP_ID*3 + i][0] + j]; \n"
|
||||
//" } \n"
|
||||
//" \n"
|
||||
//" for(int j = 0; j < 2; j++) { \n"
|
||||
//" uv_coords[i][j] = uv_buffer [2 * faces[GROUP_ID * 3 + i][1] + j]; \n"
|
||||
//" } \n"
|
||||
//" \n"
|
||||
//" for( int j = 0; j < 3; j++ ) { \n"
|
||||
//" vertices3[j][i] = (vertex[j]/vertex[3]); \n"
|
||||
//" norms[j][i] = norms_buff[3 * (faces[GROUP_ID * 3 + i][2]) + j]; \n"
|
||||
//" } \n"
|
||||
//" \n"
|
||||
//" if ( vertices3[0][i] > 0 && vertices3[0][i] < *screen_width \n"
|
||||
//" && vertices3[1][i] > 0 && vertices3[1][i] < *screen_width ) \n"
|
||||
//" out = false; \n"
|
||||
//" } \n"
|
||||
//" \n"
|
||||
//" if(out) return; \n"
|
||||
//" \n"
|
||||
//" //if(vertices3[1][0] == vertices3[1][1] && vertices3[1][2] == vertices3[1][1]) return; \n"
|
||||
//" \n"
|
||||
//" int2 bounding_box_min = (int2) { *screen_width - 1, *screen_width - 1 }; \n"
|
||||
//" int2 bounding_box_max = (int2) { 0, 0 }; \n"
|
||||
//" int2 clamper = (int2) { *screen_width - 1, *screen_width - 1 }; \n"
|
||||
//" \n"
|
||||
//" for(int i = 0; i < 3; i++) { \n"
|
||||
//" for(int j = 0; j < 2; j++) { \n"
|
||||
//" bounding_box_min[j] = max(0, min(bounding_box_min[j], (int)vertices3[j][i])); \n"
|
||||
//" bounding_box_max[j] = min(clamper[j], max(bounding_box_max[j], (int)vertices3[j][i])); \n"
|
||||
//" } \n"
|
||||
//" } \n"
|
||||
//" \n"
|
||||
//" if(bounding_box_min[0] > *screen_width || bounding_box_max[0] < 0 || bounding_box_min[1] > *screen_width || bounding_box_max[1] < 0) return; \n"
|
||||
//" \n"
|
||||
//" \n"
|
||||
//" int X_PER_ITEM = (int)(ceil((float)(bounding_box_max[0] - bounding_box_min[0]) / 16.f)); \n"
|
||||
//" int Y_PER_ITEM = (int)(ceil((float)(bounding_box_max[1] - bounding_box_min[1]) / 16.f)); \n"
|
||||
//" float STARTING_X = bounding_box_min[0] + X_PER_ITEM * (LOCAL_ID % 16); \n"
|
||||
//" float ENDING_X = STARTING_X + X_PER_ITEM; \n"
|
||||
//" float STARTING_Y = bounding_box_min[1] + Y_PER_ITEM * (LOCAL_ID / 16); \n"
|
||||
//" float ENDING_Y = STARTING_Y + Y_PER_ITEM; \n"
|
||||
//" \n"
|
||||
//" \n"
|
||||
//" float3 point; \n"
|
||||
//" for(point[0] = STARTING_X; point[0] <= ENDING_X; point[0]++) { \n"
|
||||
//" for(point[1] = STARTING_Y; point[1] <= ENDING_Y; point[1]++) { \n"
|
||||
//" if(point[1] >= *screen_width || point[1] >= *screen_width) break; \n"
|
||||
//" float3 bc_coord = barycentric(vertices3, point); \n"
|
||||
//" if (bc_coord[0] < 0 || bc_coord[1] < 0 || bc_coord[2] < 0) continue; \n"
|
||||
//" \n"
|
||||
//" float2 uv_vec = (float2){0, 0}; \n"
|
||||
//" float3 normal = (float3){0, 0, 0}; \n"
|
||||
//" \n"
|
||||
//" point[2] = dot(vertices3[2], bc_coord); \n"
|
||||
//" if (z_buffer[(int)(point[0] + point[1] * *screen_width)] > point[2]) { \n"
|
||||
//" continue; \n"
|
||||
//" } \n"
|
||||
//" \n"
|
||||
//" \n"
|
||||
//" for (int i = 0; i < 3; i++) { \n"
|
||||
//" uv_vec[0] += uv_coords[i][0] * bc_coord[i]; \n"
|
||||
//" uv_vec[1] += uv_coords[i][1] * bc_coord[i]; \n"
|
||||
//" normal[i] = dot(norms[i], bc_coord); \n"
|
||||
//" } \n"
|
||||
//" int2 uv_point = (int2) { (int)(uv_vec[0] * map_size[0]), (int)(uv_vec[1] * map_size[1]) }; \n"
|
||||
//" \n"
|
||||
//" \n"
|
||||
//" int col_index = 3 * (uv_point[0] + uv_point[1] * map_size[0]); \n"
|
||||
//" \n"
|
||||
//" \n"
|
||||
//" \n"
|
||||
//" \n"
|
||||
//" float3 normalized_norm = normalize(normal); \n"
|
||||
//" \n"
|
||||
//" \n"
|
||||
//" float intensity = clamp((dot(normalized_norm , (float3){light_dir[0], light_dir[1], light_dir[2]})), 0.f, 1.f) + 0.2; \n"
|
||||
//" \n"
|
||||
//" int color = 0; \n"
|
||||
//" color |= ((int)fmin((float)(diffuse_map[col_index + 0]) * intensity, (float) 0xff)) << 16; \n"
|
||||
//" color |= ((int)fmin((float)(diffuse_map[col_index + 1]) * intensity, (float) 0xff)) << 8; \n"
|
||||
//" color |= ((int)fmin((float)(diffuse_map[col_index + 2]) * intensity, (float) 0xff)) << 0; \n"
|
||||
//" \n"
|
||||
//" //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 16; \n"
|
||||
//" //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 8; \n"
|
||||
//" //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 0; \n"
|
||||
//" \n"
|
||||
//" \n"
|
||||
//" z_buffer[(int) (point[0] + point[1] * *screen_width)] = point[2]; \n"
|
||||
//" pixels [(int) (point[0] + point[1] * *screen_width)] = color;// & 0x00ffffff; \n"
|
||||
//" } \n"
|
||||
//" } \n"
|
||||
//"} \n";
|
||||
|
||||
Reference in New Issue
Block a user