Реших да го нахвърлям набързо да видим за колко ще се изпълни върху GPU-то (1080ti). 100 хиляди сравнения, низовете са случайни с дължина до 200 байта, kernel-а е тотално неоптимизиран, кажи-речи копи-пейст на примерна имплементация. Обаче естествено има измама, в паметта, низовете са оформени като "паскалски" такива, с първите 4 байта дължината на низа и нататък, това спестява strlen-а имплементиран в kernel-а, което не че е сложно, ама като нагази в глобалната памет да го смята, ще стане лееееееекинко по-тегаво.
main.c:
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <CL/cl.h>
#define MAX_SOURCE_SIZE (0x100000)
#define SET_SIZE 100032
#define STRING_MAX 200
#define MY_STRING "boiko borisov pederast pederast! boiko borisov pederast pederast!"
int main(void)
{
// Create the inputs
int i,j,sz;
char *A = (char*)malloc(SET_SIZE*(STRING_MAX + sizeof(int)));
char *B = (char*)malloc(STRING_MAX + sizeof(int));
char *results = (char*)malloc(SET_SIZE * sizeof(int));
char temp[STRING_MAX + sizeof(int)];
double exec_time = 0.0;
memset(A,SET_SIZE*(STRING_MAX + sizeof(int)), 0);
memset(B,STRING_MAX + sizeof(int), 0);
// Create the set A, SET_SIZE pascal-type strings up to STRING_MAX length
for(i = 0; i < SET_SIZE; i++)
{
sz = 1 + rand() % (STRING_MAX-1);
for (j = 0; j < sz; j++)
{
A[i * (STRING_MAX + sizeof(int)) + sizeof(int) + j] = (char)(rand() % 57 + 65);
}
memcpy((char*)(A + i * (STRING_MAX + sizeof(int))), &sz, sizeof(int));
}
// Create the set B, single pascal-type string
memcpy((char*)(B + sizeof(int)), MY_STRING, strlen(MY_STRING));
sz = strlen(MY_STRING);
memcpy(B, &sz, 4);
// Load the kernel
FILE *fp;
char *source_str;
size_t source_size;
fp = fopen("kernel.cl", "r");
if (!fp) {
fprintf(stderr, "Failed to load kernel.\n");
exit(1);
}
source_str = (char*)malloc(MAX_SOURCE_SIZE);
source_size = fread( source_str, 1, MAX_SOURCE_SIZE, fp);
fclose( fp );
// OpenCL stuff initialization
cl_platform_id platform_id = NULL;
cl_device_id device_id = NULL;
cl_uint ret_num_devices;
cl_uint ret_num_platforms;
cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
ret = clGetDeviceIDs( platform_id, CL_DEVICE_TYPE_DEFAULT, 1,
&device_id, &ret_num_devices);
cl_context context = clCreateContext( NULL, 1, &device_id, NULL, NULL, &ret);
cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
// Create the buffers and copy to device
cl_mem set_a = clCreateBuffer(context, CL_MEM_READ_ONLY, SET_SIZE*(STRING_MAX + sizeof(int)), NULL, &ret);
cl_mem set_b = clCreateBuffer(context, CL_MEM_READ_ONLY, STRING_MAX + sizeof(int), NULL, &ret);
cl_mem results_set = clCreateBuffer(context, CL_MEM_WRITE_ONLY, SET_SIZE * sizeof(int), NULL, &ret);
// Copy the lists A and B to their respective memory buffers
ret = clEnqueueWriteBuffer(command_queue, set_a, CL_TRUE, 0, SET_SIZE*(STRING_MAX + sizeof(int)), A, 0, NULL, NULL);
ret = clEnqueueWriteBuffer(command_queue, set_b, CL_TRUE, 0, STRING_MAX + sizeof(int), B, 0, NULL, NULL);
// Compile
cl_program program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);
ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
if (ret != 0)
{
size_t log_size;
clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
char *log = (char *) malloc(log_size);
clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
printf("%s\n", log);
}
cl_kernel kernel = clCreateKernel(program, "levenshtein", &ret);
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&set_a);
ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&set_b);
ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&results_set);
// Execute
size_t global_item_size = SET_SIZE;
size_t local_item_size = 64;
clock_t start = clock();
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_item_size, &local_item_size, 0, NULL, NULL);
printf("execution retstatus=%d\n",ret);
clock_t end = clock();
exec_time += (double)(end - start) / CLOCKS_PER_SEC;
printf("execution time = %f seconds\n", exec_time);
unsigned int *C = (int*)malloc(sizeof(int)*SET_SIZE);
ret = clEnqueueReadBuffer(command_queue, results_set, CL_TRUE, 0, SET_SIZE * sizeof(int), C, 0, NULL, NULL);
// Print distances
//for(i = 0; i < SET_SIZE; i++)
// printf("%d = %u\n", i, C[i]);
ret = clFlush(command_queue);
ret = clFinish(command_queue);
ret = clReleaseKernel(kernel);
ret = clReleaseProgram(program);
ret = clReleaseMemObject(set_a);
ret = clReleaseMemObject(set_b);
ret = clReleaseMemObject(results_set);
ret = clReleaseCommandQueue(command_queue);
ret = clReleaseContext(context);
free(A);
free(B);
free(C);
return 0;
}
kernel.cl:
#define MIN3(a, b, c) ((a) < (b) ? ((a) < (c) ? (a) : (c)) : ((b) < (c) ? (b) : (c)))
#define STRING_MAX 200
__kernel void levenshtein(__global const uchar *A, __global const char *B, __global uint *C)
{
uint s1len, s2len, x, y, lastdiag, olddiag;
uint temp;
__local uint column[STRING_MAX];
int i = get_global_id(0);
temp = i * (STRING_MAX+sizeof(uint));
s1len = (A[temp+3] << 24) | (A[temp+2] << 16) | (A[temp+1] << 8) | A[temp];
s2len = (B[3] << 24) | (B[2] << 16) | (B[1] << 8) | B[0];
for (y = 1; y <= s1len; y++)
column[y] = y;
for (x = 1; x <= s2len; x++)
{
column[0] = x;
for (y = 1, lastdiag = x-1; y <= s1len; y++)
{
olddiag = column[y];
column[y] = MIN3(column[y] + 1, column[y-1] + 1, lastdiag + (A[temp + sizeof(int) + y-1] == B[sizeof(int) + x - 1] ? 0 : 1));
lastdiag = olddiag;
}
}
C[i] = column[s1len];
}
Резултат от изпълнението:
root@debian:/home/gat3way/leventest# gcc main.c -o levtest -l OpenCL -I/usr/local/cuda-9.0/include/
root@debian:/home/gat3way/leventest# ./levtest
execution retstatus=0
execution time = 0.000041 seconds