Company

Resources

Company

Resources

Technology

OpenCL, 병렬 컴퓨팅을 위한 개방형 표준

OpenCL, 병렬 컴퓨팅을 위한 개방형 표준

OpenCL, 병렬 컴퓨팅을 위한 개방형 표준

이번에는 OpenCL을 이용하여 어떻게 GPU를 활용하는지에 대해 알아보고자 합니다. OpenCL이란 Khronos Group에서 제작 및 관리하는 개방형 병렬 컴퓨팅 프레임워크로써 Graphics Library 위에 구현된 Vulkan Compute Shader와는 달리 GPGPU를 위해 설계된 GPGPU Library입니다.

Jinwhan Shin

2024년 8월 2일

안녕하세요? ENERZAi에서 Runtime을 개발하고 있는 신진환이라고 합니다. 지난 글에서는 Vulkan Compute Shader를 통해 어떻게 GPU를 활용하여 연산을 하는지에 대해 알아보았습니다. 이번에는 OpenCL을 이용하여 어떻게 GPU를 활용하는지에 대해 알아보고자 합니다. OpenCL이란 Khronos Group에서 제작 및 관리하는 개방형 병렬 컴퓨팅 프레임워크로써 Graphics Library 위에 구현된 Vulkan Compute Shader와는 달리 GPGPU를 위해 설계된 GPGPU Library입니다. NVIDIA의 CUDA와 동일한 역할을 수행하지만 GPU만을 지원하는 CUDA와 달리 CPU 또는 NPU와 같은 가속기에 대해서도 프로그래밍을 할 수 있도록 설계가 된 것이 차이점이며, 공개된 기술 표준이기 때문에 지원하고자 하는 Vendor 누구나 지원할 수 있다는 장점이 있습니다.

지난번 Vulkan Shader에 관한 글과 마찬가지로 OpenCL을 이해하기 위해 필요한 용어와 개념들을 먼저 소개드리고 간략한 코드를 통해 어떻게 실제로 사용할 수 있는지 설명 드리겠습니다.

OpenCL Runtime

OpenCL Runtime은 위와 같은 Flow를 통해 컴파일된 코드가 GPU에서 실행될 수 있도록 지원합니다. 세부적인 기능들을 아래 주요 키워드별로 정리하였으니 참고 부탁드립니다.

Platform

Platform은 OpenCL에서 계산을 수행하는 유닛인 Device와 Device를 관리하는 하나의 Host를 묶어 표현하는 단위입니다. 간단하게 표현하면 하나의 OpenCL 구현체를 Platform이라고도 볼 수 있습니다.

따라서 만약 한 컴퓨터에 Intel CPU와 내장 그래픽 카드, 2장의 NVIDIA 외장 그래픽 카드를 설치했다면, 아래와 같이 3가지 Platform이 존재한다고 볼 수 있습니다. (일반적으로 Intel CPU Platform과 GPU Platform은 별개의 Platform으로 봅니다)

  • Intel CPU Platform: Intel CPU Device

  • Intel GPU Platform: Intel GPU Device

  • NVIDIA GPU Platform: NVIDIA GPU Device & NVIDIA GPU Device

Device

Device는 OpenCL에서 계산을 수행하는 유닛입니다. 이 유닛은 구현 방식에 따라 CPU, GPU, DSP 등 다양한 형태로 구현될 수 있습니다. OpenCL에서는 아래와 같은 Device 유형 상수를 통해 Device 유형을 정의하고 있습니다.

  • CL_DEVICE_TYPE_CPU : CPU

  • CL_DEVICE_TYPE_GPU : GPU

  • CL_DEVICE_TYPE_ACCELERATOR : Hexagon DSP, GNA, TPU 등의 AI Accelerator

  • CL_DEVICE_TYPE_CUSTOM : 기타 나머지 Device

Device는 CPU 코어처럼 연산을 수행하는 여러 개의 Compute Unit으로 구성되어 있으며, Compute Unit 또한 여러 개의 Processing Element로 이루어져 있습니다. 예를 들어,

  • Intel i5–12400은 12개의 Compute Unit으로 구성되어 있으며, 각 Compute Unit은 1개의 Processing Element로 이루어져 있습니다.

  • NVIDIA RTX 3060은 28개의 Compute Unit으로 구성되어 있으며, 각 Compute Unit은 128개의 Processing Element로 이루어져 있습니다.

Work Group & Work Item

Work Item은 하나의 Processing Element가 수행 하는 일의 단위 입니다. 쉽게 하나의 코어가 하나의 함수를 실행한다 라고 보셔도 됩니다. 여러 개의 Work Item을 그룹 지어 하나의 Work Group이 되고, Work Group의 집합이 총 수행하는 연산이 됩니다.

한 번에 실행 가능한 Work Group 크기에 제한이 존재합니다. 이 내용은 clGetDeviceInfo 함수로 Query할 수 있는데, 사용하는 Tag는 다음과 같습니다.

  • CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : Device에서 사용 가능한 최대 Work Items의 차원 수를 나타냅니다.

  • CL_DEVICE_MAX_WORK_ITEM_SIZES : Device에서 사용 가능한 차원 별 최대 Work Items의 수를 나타냅니다.

  • CL_DEVICE_MAX_WORK_GROUP_SIZE : Device에서 사용 가능한 최대 Work Group의 크기를 나타냅니다. Work Group의 모든 차원 수의 곱이 이 값을 초과할 수 없습니다.

Context

Context는 OpenCL에서 Buffer, Program, Command Queue와 같은 자원을 관리하는 객체입니다. Context는 clCreateContext로 생성하거나 clCreateContextFromType으로 Platform, Device 선택부터 Context 생성까지 한 번에 진행할 수도 있습니다.

Command Queue

Command Queue는 Buffer에 읽고 쓰거나, Kernel을 실행하는 등 Device와 상호작용하는 모든 명령을 입력하는 창구입니다. clEnqueue~ 계열 함수들이 Command Queue에 명령을 삽입하는 함수들입니다(clCreateCommandQueueWithProperties으로 Command Queue를 생성할 수 있습니다).

OpenCL에서는 기본값으로 한 번에 한 명령만 수행합니다(Sequential). 따라서 필요한 경우 CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE을 통해(Parallel) 더 나은 성능을 추구할 수 있습니다.

동기화 관련 내용은 후술할 Events 내용을 참고하시면 되겠습니다.

Buffer

Buffer는 OpenCL에서 메모리를 나타내는 객체입니다. 이 안에 어떠한 데이터가 들어갈지에 따라 아래와 같이 나누어집니다.

  • Buffer: 1차원 배열로 나타낼 수 있는 데이터. 아무거나 들어갈 수 있습니다. (int, float같은 스칼라 타입부터 float4 같은 벡터 타입, 구조체 같은 것까지요!) clCreateBuffer 으로 생성합니다. clCreateSubBuffer 를 통해 Buffer의 일부분을 사용할 수도 있습니다.

  • Image: 1/2/3차원 이미지를 나타내는 메모리. 텍스쳐에 대한 연산을 진행할 때 보통 사용합니다. clCreateImage으로 생성합니다.

  • Pipe: 작은 크기의 데이터를 FIFO로 저장할 때 사용하는 메모리입니다. 호스트에서 직접 접근할 수 없고 디바이스에서 read_pipewrite_pipe를 통해서 접근할 수 있습니다. clCreatePipe으로 생성합니다.

또한, 메모리가 어디에 존재하는지에 따라 아래와 같이 나누어집니다.

1. Host Memory

  • Host 장치의 메모리 입니다.

2. Device Memory

  • Global Memory: Device 장치의 메모리 입니다. 쉽게 GPU Memory라고 생각해도 됩니다. Host에 Mapping되어 Host가 읽고 쓸 수 있습니다. 또한 모든 Work Group 간에 공유 되는 메모리 입니다.

  • Local Memory: Device 장치의 메모리이나 Host가 접근할 수 없습니다. 같은 Work Group 안에서 공유되는 메모리 입니다. 다른 Work Group 간에는 공유되지 않습니다.

  • Private Memory: Device 장치의 메모리이나 Host가 접근할 수 없습니다. Processing Element 고유 메모리 입니다. 각각의 Processing Element 자신만 접근 가능합니다.

Buffer를 생성할 때, Buffer의 용도에 따라 다양한 옵션을 줄 수 있습니다.

  • CL_MEM_READ_WRITE: Device가 읽고 쓸 수 있는 Buffer를 생성합니다.

  • CL_MEM_READ_ONLY: Device가 읽기만 가능한 Buffer를 생성합니다.

  • CL_MEM_WRITE_ONLY: Device가 쓰기만 가능한 Buffer를 생성합니다.

  • CL_MEM_USE_HOST_PTR: Device가 접근 가능한 Buffer를 생성할 때 주어진 Host Memory를 Mapping하여 사용합니다.

  • CL_MEM_ALLOC_HOST_PTR: Device가 접근 가능한 Buffer를 생성할 때 Host Memory를 할당하고 그 메모리를 Mapping하여 사용합니다.

  • CL_MEM_COPY_HOST_PTR: Device가 접근 가능한 Buffer를 생성하고 필요할 때 Device Memory의 내용을 주어진 Host Memory에 복사합니다.

  • CL_MEM_HOST_WRITE_ONLY: Host가 쓰기만 가능한 Buffer를 생성합니다.

  • CL_MEM_HOST_READ_ONLY: Host가 읽기만 가능한 Buffer를 생성합니다.

  • CL_MEM_HOST_NO_ACCESS: Host가 접근 불가능한 Buffer를 생성합니다.

Note

CL_MEM_HOST_NO_ACCESS 옵션이 주어진 경우 그 Buffer는 Host의 영향을 받지 않으므로, 성능상 이점이 있다고 합니다. 따라서 Intermediate Buffer를 할당할 경우 필요한 경우가 아니라면 CL_MEM_HOST_NO_ACCESS Flags을 주는 것이 좋다고 합니다.

CL_MEM_USE_HOST_PTR옵션이 항상 Zero-copy(Device-mapped-memory)가 아닐 수 있습니다. Host와 Device 양쪽의 Alignment가 충족되지 않는 경우 Copy가 발생할 수 있습니다. 또한 이 옵션은 장치에 따라 권장되기도(Intel), 지양되기도(Arm) 합니다. 타겟 장치에 따라 적절히 사용 해야 할 필요성이 있습니다.

Buffer의 최대 크기 제한이 존재합니다. 제약 사항은 clGetDeviceInfo 함수로 Query 할 수 있는데, CL_DEVICE_MAX_MEM_ALLOC_SIZE Tag를 이용하여 알 수 있습니다.

Program

Program은 OpenCL에서 Kernel들과 함수들이 들어있는 하나의 집합입니다. 쉽게 설명하면 하나의 소스 파일이 컴파일된 결과라고 봐도 무방합니다. 또한, 하나의 소스 파일이 컴파일된 결과가 Program이기에 이들 여러 개를 하나로 묶어 커다란 Program을 생성할 수 있습니다(clLinkProgram).

Program은 크게 3개로 분류됩니다.

  • Source: OpenCL C 또는 OpenCL C++ 으로 작성된 소스. clCreateProgramWithSource으로 생성합니다.

  • IL: Intermeidate Representation Binary로 작성된 소스. clCreateProgramWithIL으로 생성합니다. 이것은 cl_khr_il_program Extension이 지원될 때 사용 가능합니다. SPIR-V는 cl_khr_spir Extension이 지원될 때 사용 가능하며, 옵션으로 -x spir 을 명시해야 합니다.

  • Binary: Target-device Dependent한 Binary으로 작성된 소스. clCreateProgramWithBinary으로 생성합니다.

이후 Source 또는 IL으로 생성된 프로그램은 clCompileProgram으로 컴파일 된 후 clLinkProgram을 통해 하나의 큰 Program으로 만들어 지게 됩니다.

Kernel

Kernel은 하나의 Processing Element가 수행하는 명령의 집합입니다. 함수와의 다른 점은 아래와 같습니다.

  • 엔트리 포인트의 역할

  • 함수에 kernel 또는 __kernel 키워드가 있어야 한다는 점

  • 반드시 반환 형이 void 이어야 한다는 점

Kernel Argument에 제약이 존재합니다. 제약 사항은 clGetDeviceInfo 함수로 Query할 수 있는데, 사용하는 Tag는 다음과 같습니다.

  • CL_DEVICE_MAX_PARAMETER_SIZE : 하나의 Parameter당 최대 크기를 나타냅니다. Buffer를 제외하고 Value로 전달되는 값의 크기를 나타냅니다.

  • CL_DEVICE_GLOBAL_MEM_SIZE : Global Memory의 크기를 나타냅니다.

  • CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE : Constant Buffer(constant 또는 __constant 가 붙은 pointers)의 최대 크기를 나타냅니다.

  • CL_DEVICE_MAX_CONSTANT_ARGS : 사용할 수 있는 최대 Constant Buffer의 개수를 나타냅니다.

  • CL_DEVICE_LOCAL_MEM_SIZE : Local Memory의 크기를 나타냅니다.

Events

Event는 Enqueue한 명령 사이의 의존성을 설명하는 동기화 객체입니다. clEnqueue~ 계열 함수 끝에 wait_list를 입력하고 Output Event를 출력할 수 있는데, wait_list는 대기할 Event의 목록, Event는 이 명령이 끝날 때 Signal되는 Event 객체입니다. 이를 통해 명령 간의 동기화를 수행할 수 있습니다.

사용자 정의 Event를 생성할 때에는 clCreateUserEvent를 사용하여 생성하고, clSetUserEventStatus를 통해 Signal할 수 있습니다.

이 외에도 Marker와 Barrier의 2가지 동기화 수단이 추가로 존재합니다. 두 가지 수단의 같은 점은 wait_list의 Event가 모두 Signal 될 때 Output Event가 Signal 된다는 점이며, 다른 점은 Marker는 Marker 이후에 있는 명령이 wait_list의 Event가 Signal 되지 않아도 실행 가능(signal event when events in wait_list are signaled)하지만 Barrier는 Barrier 이후에 있는 명령이 wait_list의 Event가 모두 Signal될 때 까지 Block된다는 점(block execution until events in wait_list are signaled)입니다.

Marker는 clEnqueueMarkerWithWaitList, Barrier는 clEnqueueBarrierWithWaitList를 통해 사용할 수 있습니다.

Extension

Extension은 기본 OpenCL Functionality 외의 기능을 Khronos Group이나 Device 제조사에서 추가하는 기능 입니다. 사용할 수 있는 함수 뿐만 아니라 타입도 Extension 지원 유무에 따라 사용 유무가 갈립니다. 따라서 기능 및 타입을 사용하기 전에 Device에서 Extension을 지원하는지 먼저 Query를 진행해야 합니다.

Extension이라 사용하지 못할 수도 있는 기능 몇 가지만 추려봤습니다.

  • cl_khr_fp16 : float16 연산 지원

  • cl_khr_fp64 : float64 연산 지원

  • cl_khr_il_program, cl_khr_spir : SPIR-V IR 지원

  • cl_khr_global_int32_base_atomics, cl_khr_global_int32_extended_atomics: Global Memory에 대한 int32 atomic 연산 지원

  • cl_khr_local_int32_base_atomics, cl_khr_local_int32_extended_atomics: Local Memory에 대한 int32 atomic 연산 지원

  • cl_khr_global_int64_base_atomics, cl_khr_global_int64_extended_atomics: Global Memory에 대한 int64 atomic 연산 지원

  • cl_khr_local_int64_base_atomics, cl_khr_local_int64_extended_atomics: Local Memory에 대한 int64 atomic 연산 지원

Execute Kernel with OpenCL

  1. OpenCL에서 Kernel을 실행하기 위해서는 Kernel을 실행할 Device 먼저 선택해야 합니다. 이를 위해, Platform과 Device를 Enumerate하여 실행할 Target Device를 선택합니다.

    // Get platform IDs
    std::vector<cl_platform_id> platforms;
    cl_uint num_platforms = 0;
    
    cl_device device;
    
    clGetPlatformIDs(0, nullptr, &num_platforms);
    
    platforms.resize(num_platforms);
    clGetPlatformIDs(num_platforms, platforms.data(), &num_platforms);
    
    // Get device IDs per platform
    for (auto platform : platforms) {
      std::vector<cl_device_id> devices;
      cl_uint num_devices = 0;
    
      clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &num_devices);
      
      devices.resize(num_devices);
      clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, num_devices, devices.data(), &num_devices);
    
      // select proper device
    }


  2. 사용할 Device를 선택했다면, Context와 Command Queue를 생성합니다.

    cl_context_properties context_properties[] = {
      CL_CONTEXT_PLATFORM, platform,
      0
    };
    cl_queue_properties queue_properties[] = {
      0
    };
    
    cl_int errcode;
    cl_context context = clCreateContext(
      context_properties, 1, &device, nullptr, nullptr, &errcode
    );
    cl_command_queue queue = clCreateCommandQueueWithProperties(
      context, device, queue_properties, &errcode
    );


  3. Context와 Command Queue를 생성했다면 Program을 로드합니다.

    std::string source = R"(...)";
    const char* c_str = source.c_str();
    cl_int errcode = CL_SUCCESS;
    
    // Create program with OpenCL C
    cl_program c_program = clCreateProgramWithSource(context, 1, &c_str, source.length(), &errcode);
    
    // Create program with SPIR-V IR
    std::vector<uint8_t> ir;
    cl_program ir_program = clCreateProgramWithIL(context, reinterpret_cast<const void*>(ir.data()), ir.size(), &errcode);
    
    // Create program with device-specific binary
    std::vector<uint8_t> binary;
    size_t length = binary.size();
    const uint8_t* binary_ptr = binary.data();
    cl_int status = 0;
    cl_program binary_program = clCreateProgramWithBinary(context, 1, &device, &length, &binary_ptr, &status, &errcode);


  4. 로드한 Program을 Compile하고 Link합니다.

    // Compile program
    clCompileProgram(c_program, 1, &device, "-cl-std=CL2.0", 0, nullptr, nullptr, nullptr, nullptr);
    clCompileProgram(ir_program, 1, &device, nullptr, 0, nullptr, nullptr, nullptr, nullptr);
    // no need to compile binary program; it is already compiled.
    
    // Link program
    cl_program programs[] = { c_program, ir_program, binary_program };
    cl_program linked_program = clLinkProgram(context, 1, &device, nullptr, 3, programs, nullptr, nullptr, nullptr);


  5. Linked Program에서 Kernel을 생성합니다.

    cl_int errcode;
    cl_kernel kernel = clCreateKernel(linked_program, “conv2d”, &errcode);
    
    // can gather all kernels in program at once.
    std::vector<cl_kernel> kernels;
    cl_uint num_kernels;
    
    clCreateKernelsInProgram(linked_program, 0, nullptr, &num_kernels);
    
    kernels.resize(num_kernels);
    clCreateKernelsInProgram(linked_program, num_kernels, kernels.data(), &num_kernels);


  6. Kernel의 Input and Output Buffer를 생성합니다.

    cl_int errcode;
    cl_mem input_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY, INPUT_SIZE, nullptr, &errcode);
    cl_mem output_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY|CL_MEM_HOST_READ_ONLY, OUTPUT_SIZE, nullptr, &errcode);


  7. Kernel의 Argument를 작성합니다.

    // __kernel void conv2d(__global const float* input, __global float* output);
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &input_buffer);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &output_buffer);
    
    // __kernel void DoSomeMath(__global const float4* input, __global float4* output, float4 alpha, float beta);
    cl_float4 alpha = { 1.0f, 2.0f, 3.0f, 4.0f };
    cl_float beta = -1.0f;
    
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &input_buffer);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &output_buffer);
    clSetKernelArg(kernel, 2, sizeof(cl_float4), &alpha);
    clSetKernelArg(kernel, 3, sizeof(cl_float), &beta);


  8. 필요한 경우 Buffer에 데이터를 Write합니다.

    // using clEnqueueWriteBuffer
    std::vector<uint8_t> data;
    clEnqueueWriteBuffer(queue, input_buffer, CL_TRUE, 0, INPUT_SIZE, data.data(), 0, nullptr, nullptr);
    
    // can be asynchronous, but unsupported devices may present(AMD, etc.)
    cl_event event;
    clEnqueueWriteBuffer(queue, input_buffer, CL_FALSE, 0, INPUT_SIZE, data.data(), N, barrier, &event);
    clWaitForEvents(1, &event);
    
    // using clEnqueueMapBuffer
    cl_int errcode;
    void* buffer = clEnqueueMapBuffer(queue, input_buffer, CL_TRUE, CL_MAP_WRITE, 0, INPUT_SIZE, 0, nullptr, nullptr, &errcode);
    memcpy(buffer, data.data(), data.size() * sizeof(uint8_t));
    
    // can be asychronous, but unsupported devices may present(AMD, etc.)
    void* buffer = clEnqueueMapBuffer(queue, input_buffer, CL_FALSE, CL_MAP_WRITE, 0, INPUT_SIZE, N, barrier, &event, &errcode);
    clWaitForEvents(1, &event);
    memcpy(buffer, data.data(), data.size() * sizeof(uint8_t));
    
    // Must be unmapped after map buffer.
    clEnqueueUnmapMemObject(queue, input_buffer, buffer, 0, nullptr, &event);
    clWaitForEvents(1, &event);


  9. Kernel을 실행합니다.

    cl_event event;
    cl_uint work_dim = 1;
    size_t global_worksize[1] = { 1024 };
    size_t local_worksize[1] = { 64 };
    clEnqueueNDRangeKernel(queue, kernel, work_dim, nullptr, global_worksize, local_worksize, 0, nullptr, &event);
    clWaitForEvents(1, &event);


  10. 필요한 경우 Buffer에서 데이터를 Read합니다.

    // using clEnqueueReadBuffer
    std::vector<uint8_t> data;
    clEnqueueWriteBuffer(queue, input_buffer, CL_TRUE, 0, OUTPUT_SIZE, data.data(), 0, nullptr, nullptr);
    
    // can be asynchronous, but unsupported devices may present(AMD, etc.)
    cl_event event;
    clEnqueueWriteBuffer(queue, input_buffer, CL_FALSE, 0, OUTPUT_SIZE, data.data(), N, barrier, &event);
    clWaitForEvents(1, &event);
    
    // using clEnqueueMapBuffer
    cl_int errcode;
    void* buffer = clEnqueueMapBuffer(queue, input_buffer, CL_TRUE, CL_MAP_READ, 0, OUTPUT_SIZE, 0, nullptr, nullptr, &errcode);
    memcpy(data.data(), buffer, data.size() * sizeof(uint8_t));
    
    // can be asychronous, but unsupported devices may present(AMD, etc.)
    void* buffer = clEnqueueMapBuffer(queue, input_buffer, CL_FALSE, CL_MAP_READ, 0, OUTPUT_SIZE, N, barrier, &event, &errcode);
    clWaitForEvents(1, &event);
    memcpy(data.data(), buffer, data.size() * sizeof(uint8_t));
    
    // Must be unmapped after map buffer.
    clEnqueueUnmapMemObject(queue, input_buffer, buffer, 0, nullptr, &event);
    clWaitForEvents(1, &event);


  11. 필요한 경우 8 ~ 10을 반복합니다.

  12. 사용한 Resources를 전부 Release합니다.

    clReleaseKernel(kernel);
    clReleaseProgram(linked_program);
    clReleaseProgram(binary_program);
    clReleaseProgram(ir_program);
    clReleaseProgram(c_program);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);

지금까지 OpenCL의 개념과 사용 방법에 대해 알아봤습니다. 혹시 글을 읽고 이해는 되었지만 병렬 프로그래밍과 프로그래밍 최적화를 직접 수행하기는 어려울 것 같고 부담스러우신가요? 걱정 마세요! 이러한 분들을 위해 저희가 열심히 Optimium을 개발하고 있습니다.

ENERZAi의 AI 추론 최적화 엔진인 Optimium은 자동으로 최적화를 수행하고, 현재 AMD64, Arm CPU에서 Single-thread, Multi-thread 모두 지원하고 있습니다. 또한 AI 모델 뿐만 아니라 전처리 과정 역시 가속화가 가능하며, GPU 또한 빠른 시일 내에 준비가 될 예정입니다! Optimium에 관심 있으신 분이나 저희 AI 추론 최적화 기술 관련 문의 사항이 있으신 분은 주저 없이 contact@enerzai.com으로 연락주시거나 저희 LinkedIn(www.linkedin.com/company/enerzai)에 방문해주세요! 그럼 다음 게시글에서 다시 뵙겠습니다 🙂

Optimium

Optimium

Solutions

Resources

ENERZAi

Copyright ⓒ ENERZAi Inc. All Rights Reserved

사업자등록번호: 246-86-01405

이메일: contact@enerzai.com

연락처: +82 (2) 883 1231

주소: 대한민국 서울특별시 강남구 테헤란로27길 27

Optimium

Optimium

Solutions

Resources

ENERZAi

Copyright ⓒ ENERZAi Inc. All Rights Reserved

사업자등록번호: 246-86-01405

이메일: contact@enerzai.com

연락처: +82 (2) 883 1231

주소: 대한민국 서울특별시 강남구 테헤란로27길 27

Optimium

Optimium

Solutions

Resources

ENERZAi

Copyright ⓒ ENERZAi Inc. All Rights Reserved

사업자등록번호: 246-86-01405

이메일: contact@enerzai.com

연락처: +82 (2) 883 1231

주소: 대한민국 서울특별시 강남구 테헤란로27길 27