/****************** Listing 1 / Koordinatentransformation **********/ // Caption: "Die Koordinatentransformationen können überraschend effizient mit // Hilfe eingebauter OpenCL-Funktionen in einem Kernel erfolgen. Die transformierten // Koordinaten werden dann in einem entsprechenden Buffer auf der GPU für die weitere // Verwendung gespeichert." /** * A struct holding the coordinates, colors and opacity of a single triangle, which is * defined via a surrounding circle */ typedef struct triangle_circle_struct { float middleX; float middleY; float radius; float angle1; float angle2; float angle3; ushort4 rgba_c; // == 8 byte; color definition } t_circle; /** * A struct holding the coordinates, colors and opacity of a single triangle in condensed form */ typedef struct triangle_condensed_struct { float2 tr_one; float2 tr_two; float2 tr_three; ushort4 rgba_c; // == 8 byte; color definition } t_condensed; /** * This kernel calculates cartesic coordinates for the triangle edges. It then encodes the * triangles so that they fit into a 128 bit wide struct, holding 6 half floats for the * coordinates plus 4 uchar values for the colors and opacity) */ __kernel void monalisa_triangle_transcode ( __global const t_circle *triangles_h // The triangle specifications, as provided by the host , __global t_condensed *triangles_c // The condensed triangle specifications, as passed to the picture evaluation ) { // Retrieve our global id int gid = get_global_id(0); // Copy the data of "our" triangle to private memory float middleX = triangles_h[gid].middleX; float middleY = triangles_h[gid].middleY; float radius = triangles_h[gid].radius; float angle1 = triangles_h[gid].angle1; float angle2 = triangles_h[gid].angle2; float angle3 = triangles_h[gid].angle3; ushort4 rgba_c = triangles_h[gid].rgba_c; // and store them in the triangle struct triangles_c[gid].tr_one = (float2)( mad(radius, native_cos(angle1*2.0f*M_PI_F), middleX)" // similar to middleX+radius*cos(angle1*2.0f*M_PI_F); , mad(radius, native_sin(angle1*2.0f*M_PI_F), middleY) // similar to middleY+radius*sin(angle1*2.0f*M_PI_F); ); triangles_c[gid].tr_two = (float2)( mad(radius, native_cos(angle2*2.0f*M_PI_F), middleX) // similar to middleX+radius*cos(angle2*2.0f*M_PI_F); , mad(radius, native_sin(angle2*2.0f*M_PI_F), middleY) // similar to middleY+radius*sin(angle2*2.0f*M_PI_F); ); triangles_c[gid].tr_three = (float2)( mad(radius, native_cos(angle3*2.0f*M_PI_F), middleX) // similar to middleX+radius*cos(angle3*2.0f*M_PI_F); , mad(radius, native_sin(angle3*2.0f*M_PI_F), middleY) // similar to middleY+radius*sin(angle3*2.0f*M_PI_F); ); triangles_c[gid].rgba_c = rgba_c; } /*******************************************************************/ /****************** Listing 2 / Zeichnen der Dreiecke **********/ // Caption: "Der Check auf Containment des Pixels in einem Dreieck // erfolgt mit der Methode der Baryzentrischen Koordinaten. Hierbei werden // zur Performance-Steigerung viele eingebaute Funktionen von OpenCL // eingesetzt, die jeweils durch den Compiler soweit wie möglich auf die verfügbare // Hardware abgebildet werden." /** * Checks whether a given pixel is inside a given triangle and if so, adds the triangles' * colors to the pixel */ inline void checkAndAdd(__global t_condensed *t, float4 *c_pixel, float2 *pos_f) { float dot11 = dot(t->tr_three - t->tr_one, t->tr_three - t->tr_one); float dot12 = dot(t->tr_three - t->tr_one, t->tr_two - t->tr_one); float dot22 = dot(t->tr_two - t->tr_one, t->tr_two - t->tr_one); float dot1p = dot(t->tr_three - t->tr_one, *pos_f - t->tr_one); float dot2p = dot(t->tr_two - t->tr_one, *pos_f - t->tr_one); float denom = fmax(dot11*dot22 - dot12*dot12, 0.0000001f); float u = native_divide((dot22*dot1p - dot12*dot2p), denom); float v = native_divide((dot11*dot2p - dot12*dot1p), denom); if((u >= 0.f) && (v >= 0.f) && (u+v < 1.f)) { float4 tr_cols = convert_float4(t->rgba_c)*native_recip((float)(MAXCOL)); float4 alphaVec = (float4)(tr_cols.s3); *c_pixel = mix(*c_pixel, tr_cols, alphaVec); } } /** * This kernel checks for a given pixel whether it is contained in a set of * triangles, ordered according to their alpha channel. If the pixel is contained in a triangle, * it adds its colors to the pixel. */ __kernel void monalisa_candidate_creator ( __global t_condensed *triangles_c , __write_only image2d_t candidate ){ // Initialize the candidate pixel with black float4 c_pixel = (float4)(1.f); float2 pos_f = (float2)((float)(get_global_id(0)+1)*XDIMINV, (float)(get_global_id(1)+1)*YDIMINV); for(size_t i_t=0; i_t