문자열 연결을위한 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

나는 당신이 GPU에 CONCAT 작업을 오프로드하여 이득을 많이 볼 수 있다고 생각하지만, 여기에 내가 어떻게 할 것입니다하지 않습니다

__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'; 
    } 
} 

당신은에 전달해야 null 값을 검색하는 대신 연결하려는. 자 열의 크기. 이 커널은 다양한 크기의 D 및 E 입력을 사용하여 작업 항목 수에 관계없이 작동합니다. 늘 그렇듯이 F는 dSize + eSise + 1 문자를 보유 할만큼 충분히 커야합니다.

각 작업 항목은 약 (dSize + eSize)/globalSize 문자를 출력으로 복사합니다. 개선을위한

룸 :

• 는 장치와 입력 크기에 대한 최적의 값을 찾기 위해 다른 글로벌 작업 크기를 시도
• 단일 작업을 시도 할 경우 글로벌 메모리 액세스가 꽤 잘되어야합니다 그룹을 만들고 로컬 메모리를 사용하면 도움이 될 수 있지만 전역 읽기 속도에 구속을받습니다.