用于字符串连接的 OpenCL 内核

Question

我没有找到太多使用 GPU 对字符串执行操作的文献或示例。具体来说，我有 2 个字符串数组，我需要将第二个数组的元素连接到第一个数组的相应元素。我不知道如何为此编写内核。

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 字符复制到输出。

改进空间：

• 尝试不同的全局工作尺寸，以找到适合您的设备和输入尺寸的最佳值
• 全局内存访问应该不错，如果您想尝试单个工作组并使用本地内存，这可能会有所帮助，但您会受到全局读取速度的限制。