static bool_t modify_urls (blockmail_t *blockmail, receiver_t *rec, block_t *block) /*{{{*/ { int n; int len; const xmlChar *cont; int lstore; int state; char initial; char ch; int start, end; int mask; int clen; bool_t changed; xmlBufferEmpty (block -> out); len = xmlBufferLength (block -> in); cont = xmlBufferContent (block -> in); lstore = 0; state = ST_INITIAL; if (block -> nr == 1) initial = 'h'; else if (block -> nr == 2) initial = '<'; else initial = '\0'; start = -1; end = -1; mask = 1 << (block -> nr - 1); changed = false; for (n = 0; n <= len; ) { if (n < len) { clen = xmlCharLength (cont[n]); if ((clen == 1) && isascii ((char) cont[n])) { ch = (char) cont[n]; switch (state) { case ST_INITIAL: if (tolower (ch) == initial) { if (block -> nr == 1) { state = 1; start = n; } else if (block -> nr == 2) state = 100; } break; # define CHK(ccc) do { if ((ccc) == ch) ++state; else state = ST_INITIAL; } while (0) # define CCHK(ccc) do { if ((ccc) == tolower (ch)) ++state; else state = ST_INITIAL; } while (0) case 1: CCHK ('t'); break; case 2: CCHK ('t'); break; case 3: CCHK ('p'); break; case 4: ++state; if (tolower (ch) == 's') break; /* Fall through . . . */ case 5: CHK (':'); break; case 6: CHK ('/'); break; case 7: if (ch == '/') state = ST_START_FOUND; else state = ST_INITIAL; break; case 100: CCHK ('a'); break; # define HCHK(ccc) do { if ((ccc) == tolower (ch)) ++state; else if ('>' == ch) state = ST_INITIAL; else state = 101; } while (0) case 101: HCHK ('h'); break; case 102: HCHK ('r'); break; case 103: HCHK ('e'); break; case 104: HCHK ('f'); break; # undef HCHK case 105: CHK ('='); break; case 106: ++state; if (ch == '"') break; /* Fall through . . */ case 107: if (tolower (ch) == 'h') { state = 1; start = n; } else state = ST_INITIAL; break; case ST_START_FOUND: if (isspace ((int) ((unsigned char) ch)) || (ch == '"') || (ch == '>')) { end = n; state = ST_END_FOUND; } break; } # undef CHK # undef CCHK } else { if (state == ST_START_FOUND) { end = n; state = ST_END_FOUND; } else state = ST_INITIAL; } n += clen; } else { if (state == ST_START_FOUND) { end = n; state = ST_END_FOUND; } ++n; } if (state == ST_END_FOUND) { int ulen, m, o; ulen = end - start; for (m = 0; m < blockmail -> url_count; ++m) if ((blockmail -> url[m] -> usage & mask) && (blockmail -> url[m] -> dlen == ulen) && (! xmlStrncmp (blockmail -> url[m] -> dptr, cont + start, ulen))) break; if (m < blockmail -> url_count) { if (lstore < start) xmlBufferAdd (block -> out, cont + lstore, start - lstore); for (o = 0; o < rec -> url_count; ++o) if (rec -> url[o] -> uid == blockmail -> url[m] -> uid) break; if (o < rec -> url_count) xmlBufferAdd (block -> out, rec -> url[o] -> dptr, rec -> url[o] -> dlen); lstore = end; changed = true; } state = ST_INITIAL; } } if (changed) { if (lstore < len) xmlBufferAdd (block -> out, cont + lstore, len - lstore); SWAP (block); } return true; }/*}}}*/
static void run_test(int image_width, int image_height, int image_pitch, int chan_order, int chan_type) { unsigned int num_platforms; int err, i; #define NAME(x) [x & 0xf] = #x static const char *channel_orders[] = { NAME(CL_R), NAME(CL_A), NAME(CL_RG), NAME(CL_RA), NAME(CL_RGB), NAME(CL_RGBA), NAME(CL_BGRA), NAME(CL_ARGB), NAME(CL_INTENSITY), NAME(CL_LUMINANCE), NAME(CL_Rx), NAME(CL_RGx), NAME(CL_RGBx), }; static const char *channel_types[] = { NAME(CL_SNORM_INT8), NAME(CL_SNORM_INT16), NAME(CL_UNORM_INT8), NAME(CL_UNORM_INT16), NAME(CL_UNORM_SHORT_565), NAME(CL_UNORM_SHORT_555), NAME(CL_UNORM_INT_101010), NAME(CL_SIGNED_INT8), NAME(CL_SIGNED_INT16), NAME(CL_SIGNED_INT32), NAME(CL_UNSIGNED_INT8), NAME(CL_UNSIGNED_INT16), NAME(CL_UNSIGNED_INT32), NAME(CL_HALF_FLOAT), NAME(CL_FLOAT), }; size_t global; // global domain size for our calculation size_t local; // local domain size for our calculation size_t bin_size; size_t len; unsigned char *bin; cl_platform_id platform; cl_device_id device_id; // compute device id cl_context context; // compute context cl_command_queue commands; // compute command queue cl_program program; // compute program cl_kernel kernel; // compute kernel cl_mem input1, input2; // device memory used for the input array cl_mem output; // device memory used for the output array RD_START("image", "width=%d, height=%d, pitch=%d, order=%s, type=%s", image_width, image_height, image_pitch, channel_orders[chan_order & 0xf], channel_types[chan_type & 0xf]); CCHK(clGetPlatformIDs(1, &platform, &num_platforms)); CCHK(clGetPlatformInfo(platform, CL_PLATFORM_EXTENSIONS, sizeof(buffer), buffer, &len)); DEBUG_MSG("extensions=%s\n", buffer); CCHK(clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL)); CCHK(clGetDeviceInfo(device_id, CL_DEVICE_VERSION, sizeof(buffer), buffer, &len)); DEBUG_MSG("version=%s\n", buffer); CCHK(clGetDeviceInfo(device_id, CL_DEVICE_EXTENSIONS, sizeof(buffer), buffer, &len)); DEBUG_MSG("device extensions=%s\n", buffer); context = clCreateContext(0, 1, &device_id, NULL, NULL, NULL); commands = clCreateCommandQueue(context, device_id, 0, NULL); program = clCreateProgramWithSource(context, 1, (const char **) &KernelSource, NULL, NULL); CCHK(clBuildProgram(program, 0, NULL, "-cl-mad-enable -DFILTER_SIZE=1 " "-DSAMP_MODE=CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_CLAMP_TO_EDGE|CLK_FILTER_NEAREST", NULL, NULL)); kernel = clCreateKernel(program, "test_kernel", NULL); /* fill buffer with dummy pattern: */ for (i = 0; i < 256; i++) buffer[i] = i; const cl_image_format format = { chan_order, chan_type }; input1 = clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, &format, image_width, image_height, image_pitch, buffer, &err); CCHK(err); input2 = clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, &format, image_width+1, image_height+1, image_pitch, buffer + 4, &err); CCHK(err); output = clCreateImage2D(context, CL_MEM_WRITE_ONLY, &format, image_width+2, image_height+2, 0, NULL, &err); CCHK(err); CCHK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &input1)); CCHK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &input2)); CCHK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &output)); CCHK(clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL)); global = 1024; CCHK(clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global, &local, 0, NULL, NULL)); CCHK(clFinish(commands)); // CCHK(clEnqueueReadBuffer( commands, output, CL_TRUE, 0, sizeof(float) * count, results, 0, NULL, NULL)); clReleaseMemObject(input1); clReleaseMemObject(input2); clReleaseMemObject(output); clReleaseProgram(program); clReleaseKernel(kernel); clReleaseCommandQueue(commands); clReleaseContext(context); RD_END(); }