Ядро OpenCL для конкатенации строк

Question

Я не нашел много литературы или примеров выполнения операций с строками с графическим процессором. В частности, у меня есть 2 массива строк, и мне нужно объединить элементы 2-го массива с соответствующими элементами 1-го массива. Я не могу понять, как написать ядро ​​для этого.

Пример конкатенации в C будет:

#include <stdio.h> 

void concatenate_string(char*, char*, char*); 

int main() 
{ 
    char original[100], add[100], result[100]; 
    printf("Enter source string\n"); 
    scanf("%s", original); 
    printf("Enter string to concatenate\n"); 
    scanf("%s", add); 
    concatenate_string(original, add, result); 

    printf("String after concatenation is \"%s\"\n", result); 

    return 0; 
} 

void concatenate_string(char *original, char *add, char *result) 
{ 
    while(*original) 
    { 
     *result = *original; 
     original++; 
     result++; 
    } 
    while(*add) 
    { 
     *result = *add; 
     add++; 
     result++; 
    } 
    *result = '\0'; 
} 

Ниже мой OpenCL код-хозяин, содержащая ядро. Ядро следует за тем же потоком, что и функция concatenate_string выше. Программа выполняется успешно, но не дает результата.

#include <stdio.h> 
#include <stdlib.h> 
#include <string.h> 
#ifdef __APPLE__ 
#include <OpenCL/cl.h> 
#else 
#include <CL/cl.h> 
#endif 
#include <ocl_macros.h> 
#include <iostream> 
#include <string> 
//Common defines 
#define VENDOR_NAME "AMD" 
#define DEVICE_TYPE CL_DEVICE_TYPE_GPU 
#define VECTOR_SIZE 1024 
using namespace std; 

//OpenCL kernel which is run for every work item created. 
//The below const char string is compiled by the runtime complier 
//when a program object is created with clCreateProgramWithSource 
//and built with clBuildProgram. 
const char *concat_kernel = 
"__kernel          \n" 
"void concat_kernel(       \n" 
"     __global uchar *D,   \n" 
"     __global uchar *E,   \n" 
"     __global uchar *F)   \n" 
"{            \n" 
" //Get the index of the work-item   \n" 
" int index = get_global_id(0);    \n" 
" while(D[index])       \n" 
" {           \n" 
"  *F[index] = *D[index];     \n" 
"  D[index]++;       \n" 
"  F[index]++;       \n" 
" }           \n" 
" while(E[index])       \n" 
" {           \n" 
"  *F[index] = *E[index];     \n" 
"  E[index]++;       \n" 
"  F[index]++;       \n" 
" }           \n" 
" *F[index] = '\0';       \n" 
"}            \n"; 

int main(void) { 

    cl_int clStatus; //Keeps track of the error values returned. 

    // Get platform and device information 
    cl_platform_id * platforms = NULL; 

    // Set up the Platform. Take a look at the MACROs used in this file. 
    // These are defined in common/ocl_macros.h 
    OCL_CREATE_PLATFORMS(platforms); 

    // Get the devices list and choose the type of device you want to run on 
    cl_device_id *device_list = NULL; 
    OCL_CREATE_DEVICE(platforms[0], DEVICE_TYPE, device_list); 

    // Create OpenCL context for devices in device_list 
    cl_context context; 
    cl_context_properties props[3] = 
    { 
     CL_CONTEXT_PLATFORM, 
     (cl_context_properties)platforms[0], 
     0 
    }; 
    // An OpenCL context can be associated to multiple devices, either CPU or GPU 
    // based on the value of DEVICE_TYPE defined above. 
    context = clCreateContext(NULL, num_devices, device_list, NULL, NULL, &clStatus); 
    LOG_OCL_ERROR(clStatus, "clCreateContext Failed..."); 

    // Create a command queue for the first device in device_list 
    cl_command_queue command_queue = clCreateCommandQueue(context, device_list[0], 0, &clStatus); 
    LOG_OCL_ERROR(clStatus, "clCreateCommandQueue Failed..."); 

    // Allocate space for vectors D, E, and F 
    string *D = (string*)malloc(sizeof(string)*VECTOR_SIZE); 
    string *E = (string*)malloc(sizeof(string)*VECTOR_SIZE); 
    string *F = (string*)malloc(sizeof(string)*VECTOR_SIZE); 
    for(int i = 0; i < VECTOR_SIZE; i++) 
    { 
     D[i] = ".25_numstring"; 
    } 
    for(int i = 0; i < VECTOR_SIZE; i++) 
    { 
     E[i] = "string_2"; 
     F[i] = "0"; 
    } 
    // Create memory buffers on the device for each vector 
    cl_mem D_clmem = clCreateBuffer(context, CL_MEM_READ_ONLY, 
      VECTOR_SIZE * sizeof(string), NULL, &clStatus); 
    cl_mem E_clmem = clCreateBuffer(context, CL_MEM_READ_ONLY, 
      VECTOR_SIZE * sizeof(string), NULL, &clStatus); 
    cl_mem F_clmem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 
      VECTOR_SIZE * sizeof(string), NULL, &clStatus); 

    // Copy the Buffer D and E to the device. We do a blocking write to the device buffer. 
    clStatus = clEnqueueWriteBuffer(command_queue, D_clmem, CL_TRUE, 0, 
      VECTOR_SIZE * sizeof(string), D, 0, NULL, NULL); 
    LOG_OCL_ERROR(clStatus, "clEnqueueWriteBuffer Failed..."); 
    clStatus = clEnqueueWriteBuffer(command_queue, E_clmem, CL_TRUE, 0, 
      VECTOR_SIZE * sizeof(string), E, 0, NULL, NULL); 
    LOG_OCL_ERROR(clStatus, "clEnqueueWriteBuffer Failed..."); 

    // Create a program from the kernel source 
    cl_program program = clCreateProgramWithSource(context, 1, 
      (const char **)&concat_kernel, NULL, &clStatus); 
    LOG_OCL_ERROR(clStatus, "clCreateProgramWithSource Failed..."); 

    // Build the program 
    clStatus = clBuildProgram(program, 1, device_list, NULL, NULL, NULL); 
    if(clStatus != CL_SUCCESS) 
     LOG_OCL_COMPILER_ERROR(program, device_list[0]); 

    // Create the OpenCL kernel 
    cl_kernel kernel = clCreateKernel(program, "concat_kernel", &clStatus); 

    // Set the arguments of the kernel. Take a look at the kernel definition in concat_kernel 
    // variable. First parameter is a constant and the other three are buffers. 
    clStatus |= clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&D_clmem); 
    clStatus |= clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&E_clmem); 
    clStatus |= clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&F_clmem); 
    LOG_OCL_ERROR(clStatus, "clSetKernelArg Failed..."); 

    // Execute the OpenCL kernel on the list 
    size_t global_size = VECTOR_SIZE; // Process one vector element in each work item 
    size_t local_size = 64;   // Process in work groups of size 64. 
    cl_event concat_event; 
    clStatus = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, 
      &global_size, &local_size, 0, NULL, &concat_event); 
    LOG_OCL_ERROR(clStatus, "clEnqueueNDRangeKernel Failed..."); 

    // Read the memory buffer F_clmem on the device to the host allocated buffer C 
    // This task is invoked only after the completion of the event concat_event 
    clStatus = clEnqueueReadBuffer(command_queue, F_clmem, CL_TRUE, 0, 
      VECTOR_SIZE * sizeof(string), F, 1, &concat_event, NULL); 
    LOG_OCL_ERROR(clStatus, "clEnqueueReadBuffer Failed..."); 

    // Clean up and wait for all the comands to complete. 
    clStatus = clFinish(command_queue); 

    // Display the result to the screen 
    for(int i = 0; i < VECTOR_SIZE; i++) 
     printf("%s + %s = %s\n", D[i].c_str(), E[i].c_str(), F[i].c_str()); 

    // Finally release all OpenCL objects and release the host buffers. 
    clStatus = clReleaseKernel(kernel); 
    clStatus = clReleaseProgram(program); 
    clStatus = clReleaseMemObject(D_clmem); 
    clStatus = clReleaseMemObject(E_clmem); 
    clStatus = clReleaseMemObject(F_clmem); 
    clStatus = clReleaseCommandQueue(command_queue); 
    clStatus = clReleaseContext(context); 
    free(D); 
    free(E); 
    free(F); 
    free(platforms); 
    free(device_list); 

    return 0; 
} 

Answer 1

Я не думаю, что вы будете видеть большую часть выручки, перенося операцию Concat к GPU, но вот как я бы это сделать:

__kernel void concat_kernel(__global uchar *D,__global uchar *E,__global uchar *F, const int dSize, const int eSize) 
{ 
    int gid = get_global_id(0); 
    int globalSize = get_global_size(0); 

    int i; 
    for(i=gid; i< dSize; i+= globalSize){ 
     F[i] = D[i]; 
    } 

    for(i=gid; i< eSize; i+= globalSize){ 
     F[i+dSize] = E[i]; 
    } 

    if(gid == globalSize-1){ 
     //using the last work item here because it will be 
     //idle when (dSize+eSize) % globalSize != 0 
     F[dSize + eSize -1] = '\0'; 
    } 
} 

Вам необходимо пройти в размеры строк, которые вы хотите объединить, вместо поиска нулевого значения. Это ядро ​​будет работать с любым количеством рабочих элементов и с различными входами D и E. Как обычно, F должен быть достаточно большим, чтобы удерживать dSize + eSise + 1 символ.

Каждый рабочий элемент будет копировать символы (dSize + eSize)/globalSize на выход.

Комната для усовершенствования:

• попробовать различные размеры глобальной работы, чтобы найти оптимальное значение для вашего устройства и ввода размера
• глобальный доступ к памяти должен быть очень хорошо, если вы хотите попробовать один работу группе и использовать локальную память, это может помочь, но вы будете связаны глобальной скоростью чтения.