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();
}
