**OpenCL(Open Computing Language)**是一种用于异构平台(包括CPU、GPU、FPGA、DSP等)上的并行计算框架和编程标准。它由Khronos Group制定,旨在提供一种跨平台、统一的编程接口,使开发者可以利用不同硬件设备进行高性能并行计算。
OpenCL的核心概念
- 平台(Platform):代表不同厂商的计算平台,如NVIDIA CUDA、AMD GPU、Intel CPU等。
- 设备(Device):平台下的具体计算硬件,比如GPU、CPU核。
- 上下文(Context):管理设备和资源的环境。
- 命令队列(Command Queue):提交任务(如内核执行、内存操作)到设备。
- 内核(Kernel):在设备上运行的计算函数。
- 程序对象(Program):包含编译后的内核代码。
OpenCL的使用流程
- 获取平台和设备
- 创建上下文和命令队列
- 编译内核程序(Kernel)代码
- 创建缓冲区(Buffer)
- 设置内核参数
- 运行(Enqueue)内核
- 读取结果
- 清理释放资源
OpenCL的代码结构通常由两部分组成:主程序代码(Host Program) 和 内核代码(Kernel Code)。它们可以用C或C++编写,取决于你的编译环境和API的使用。
- OpenCL API本身是基于C的,所以在“主程序”中,通常用C或C++都可以调用。
- 可以用C++,比如用STL容器、更复杂的封装,也可以用
cl.hpp(C++封装版的OpenCL头文件),提供更面向对象的接口。 - 内核代码通常用C语言风格,代码在
.cl文件中。
1. 文件后缀
-
Host程序(主代码):一般用
.c(纯C)或.cpp(C++)文件编写。例如:main.cpphost.c
-
内核代码(Device端代码):用特殊的源文件,常用后缀包括:
.cl(主要文件扩展名)- 也可以用
.cpp或其他扩展,不过标准和习惯是用.cl
总结:
.cl文件:存放OpenCL内核程序(GPU、CPU上的设备程序)- 主程序文件(C或C++):调用OpenCL API,负责加载、编译内核、管理数据等
2. 代码结构
例:典型的OpenCL程序结构(包括两个文件)
a. kernel.cl (内核代码)(在GPU上运行)
// kernel function 向量加法
__kernel void vector_add(__global const float* A, __global const float* B, __global float* C, int N) {
int i = get_global_id(0);
if (i < N) {
C[i] = A[i] + B[i];
}
}
b. 主程序用C++调用OpenCL API(在CPU上执行)】
- 包含OpenCL API调用:
- 选择平台和设备
- 创建上下文和命令队列
- 加载内核代码(读入
kernel.cl文件内容) - 编译程序
- 设置参数、分配缓冲区
- 启动核函数
- 读出和处理结果
方法1:
#include <iostream>
#include <vector>
#include <fstream>
#include <streambuf>
#include <CL/cl.h>
const char* kernel_file = "kernel.cl";
int main() {
// 1. 读取内核源码文件
std::ifstream kernel_stream(kernel_file);
if (!kernel_stream.is_open()) {
std::cerr << "Failed to open kernel file." << std::endl;
return -1;
}
std::string kernel_code((std::istreambuf_iterator<char>(kernel_stream)), std::istreambuf_iterator<char>());
const char* kernel_source = kernel_code.c_str();
// 2. 获取平台
cl_platform_id platform;
clGetPlatformIDs(1, &platform, NULL);
// 3. 获取设备(GPU或CPU)
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, NULL);
// 4. 创建上下文
cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
// 5. 创建命令队列
cl_command_queue queue = clCreateCommandQueue(context, device, 0, NULL);
// 6. 编译内核程序
const size_t source_size = kernel_code.size();
cl_program program = clCreateProgramWithSource(context, 1, &kernel_source, &source_size, NULL);
if (clBuildProgram(program, 1, &device, NULL, NULL, NULL) != CL_SUCCESS) {
// 输出编译错误信息
size_t log_size;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
std::vector<char> build_log(log_size);
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, build_log.data(), NULL);
std::cerr << "Error in kernel:\n" << build_log.data() << std::endl;
clReleaseProgram(program);
clReleaseContext(context);
return -1;
}
// 7. 创建内核
cl_kernel kernel = clCreateKernel(program, "vector_add", NULL);
// 8. 创建向量数据
const int N = 1024;
std::vector<float> A(N, 1.0f);
std::vector<float> B(N, 2.0f);
std::vector<float> C(N, 0);
// 9. 创建缓冲区
cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * N, A.data(), NULL);
cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * N, B.data(), NULL);
cl_mem bufC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * N, NULL, NULL);
// 10. 设置内核参数
clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufB);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufC);
clSetKernelArg(kernel, 3, sizeof(int), &N);
// 11. 定义全局与本地工作项数
size_t global_size = ((N + 255) / 256) * 256; // 以256为块大小的倍数
size_t local_size = 256;
// 12. 执行内核
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, &local_size, 0, NULL, NULL);
// 13. 读取结果
clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, sizeof(float) * N, C.data(), 0, NULL, NULL);
// 14. 输出前几个结果验证
std::cout << "C[0] = " << C;
// 释放资源
clReleaseMemObject(bufA);
clReleaseMemObject(bufB);
clReleaseMemObject(bufC);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);
}
方法2:
#include <CL/cl.h>
#include <iostream>
#include <vector>
const char* kernel_source = R"(
__kernel void vector_add(__global const float* A, __global const float* B, __global float* C, int N) {
int i = get_global_id(0);
if (i < N) {
C[i] = A[i] + B[i];
}
}
)";
int main() {
// 1. 获取平台和设备
cl_platform_id platform;
cl_device_id device;
clGetPlatformIDs(1, &platform, NULL);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
// 2. 创建上下文和命令队列
cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
cl_command_queue queue = clCreateCommandQueue(context, device, 0, NULL);
int N = 1024;
std::vector<float> A(N, 1.0f), B(N, 2.0f), C(N);
// 3. 创建缓冲区
cl_mem bufferA = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * N, A.data(), NULL);
cl_mem bufferB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * N, B.data(), NULL);
cl_mem bufferC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * N, NULL, NULL);
// 4. 编译程序
cl_program program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, NULL);
clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
cl_kernel kernel = clCreateKernel(program, "vector_add", NULL);
// 5. 设置内核参数
clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufferA);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufferB);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufferC);
clSetKernelArg(kernel, 3, sizeof(int), &N);
// 6. 运行内核
size_t global_size = N;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL);
// 7. 读取结果
clEnqueueReadBuffer(queue, bufferC, CL_TRUE, 0, sizeof(float) * N, C.data(), 0, NULL, NULL);
// 8. 输出和清理
std::cout << "C[0] = " << C[0] << std::endl;
clReleaseMemObject(bufferA);
clReleaseMemObject(bufferB);
clReleaseMemObject(bufferC);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
cmake 编译opencl代码
用CMake管理OpenCL项目,包括编译.cl内核文件和保护源代码,主要涉及以下几个方面:
1. 在CMake中如何处理.cl文件
方案一:将.cl文件作为资源文件(非源代码)打包到项目中
- 将内核代码存放在项目目录内,比如
kernels/vector_add.cl。 - 使用
configure_file()或file()命令,将内核文件复制到输出目录。 - 在运行时,主程序用文件IO加载
vector_add.cl的内容。
示例:CMakeLists.txt
cmake_minimum_required(VERSION 3.14)
project(OpenCLExample)
# 添加可执行文件
add_executable(opencl_example main.cpp)
# 包含OpenCL头文件路径(根据你的环境调整)
target_include_directories(opencl_example PRIVATE /path/to/OpenCL/headers)
# 安装内核源码文件
install(FILES kernels/vector_add.cl DESTINATION ${CMAKE_BINARY_DIR}/kernels)
# 在配置阶段复制内核文件到输出目录
configure_file(kernels/vector_add.cl ${CMAKE_BINARY_DIR}/kernels/vector_add.cl COPYONLY)
在代码中加载内核
// 在运行时加载内核文件内容
std::ifstream kernel_file("path/to/kernels/vector_add.cl");
std::string kernel_source((std::istreambuf_iterator<char>(kernel_file)), std::istreambuf_iterator<char>());
或使用与configure_file()相配合的路径。
2. 编译内核文件到二进制(.bin)以保护源代码
- 可以在CMake中调用
clCreateProgramWithBinary(),预编译内核成二进制(.bin文件),避免暴露源码。
步骤:
- 先用
clBuildProgram()生成内核二进制(用OpenCL工具或程序) - 将
.bin文件存放在项目中 - 使用
clCreateProgramWithBinary()加载预编译的二进制
示例(伪代码):
// 读取内核的二进制文件
std::ifstream bin_file("vector_add.bin", std::ios::binary);
std::vector<unsigned char> binary((std::istreambuf_iterator<char>(bin_file)), std::istreambuf_iterator<char>());
// 传入clCreateProgramWithBinary
cl_program program = clCreateProgramWithBinary(context, 1, &device, &binary.size(), (const unsigned char**)&binary[0], &binary_status, &err);
3. 如何保护代码不泄露?
- 预编译成二进制(pocl,SPIR-V、OpenCL二进制):
- 编译内核为二进制文件,这样源代码不会被直接暴露。
- 代码混淆:
- 不常用,难以实现,效果有限。
- 硬件保护方案(如GPU的安全特性):
- 一些GPU厂商提供程序保护,但难以完全防止逆向。
4. 总结
| 操作 | 说明 |
|---|---|
.cl源文件处理 | 在CMake中复制到可访问路径,运行时加载,保护较弱 |
| 预编译成二进制文件 | 用OpenCL工具或程序编译成.bin文件,运行时用clCreateProgramWithBinary()加载,增强保护 |
| 保护源码 | 最好只发行二进制,或者使用硬件/软件级别的保护方案 |
5.示例
一、项目结构示范
OpenCLProject/
├── CMakeLists.txt
├── src/
│ └── main.cpp
└── kernels/
└── vector_add.cl
二、操作流程和示例
1. 准备OpenCL内核源代码
**kernels/vector_add.cl**内容:
__kernel void vector_add(__global const float* A, __global const float* B, __global float* C, int N) {
int i = get_global_id(0);
if (i < N) {
C[i] = A[i] + B[i];
}
}
2. 预编译内核,生成二进制(二进制保护)
方法:
- 使用官方OpenCL SDK提供的工具(如
clBuildProgram后用clGetProgramInfo()的CL_PROGRAM_BINARY_SIZES和CL_PROGRAM_BINARIES)在程序运行时生成 - 或者用OpenCL API在代码运行时将源编译为二进制文件(
clCompileProgram)简便方案:
- 用OpenCL程序在第一次运行时,将
clBuildProgram的二进制保存到文件
示例:在main.cpp中添加代码(仅做提示,实际在项目中编写)
// 省略初始化(平台、设备、上下文)...
cl_program program = clCreateProgramWithSource(context, 1, &kernel_source, &source_size, &err);
clBuildProgram(program, 1, &device, NULL, NULL, NULL);
// 获取二进制
size_t binary_size;
clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binary_size, NULL);
unsigned char* binary = new unsigned char[binary_size];
unsigned char* binaries[] = {binary};
clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(unsigned char*), &binaries, NULL);
// 保存二进制到文件
std::ofstream bin_file("vector_add.bin", std::ios::binary);
bin_file.write((char*)binary, binary_size);
bin_file.close();
delete[] binary;
clReleaseProgram(program);
此文件vector_add.bin可以在部署时分发,用于替换源代码,避免泄露。
3. 在正式运行时用二进制程序
用二进制加载(示例代码片段):
// 加载二进制文件
std::ifstream bin_file("vector_add.bin", std::ios::binary);
size_t bin_size;
bin_file.seekg(0, std::ios::end);
bin_size = bin_file.tellg();
bin_file.seekg(0, std::ios::beg);
unsigned char* binary_data = new unsigned char[bin_size];
bin_file.read((char*)binary_data, bin_size);
bin_file.close();
cl_int binary_status;
cl_program bin_program = clCreateProgramWithBinary(context, 1, &device, &bin_size, (const unsigned char**)&binary_data, &binary_status, &err);
delete[] binary_data;
clBuildProgram(bin_program, 0, NULL, NULL, NULL, NULL);
cl_kernel kernel_bin = clCreateKernel(bin_program, "vector_add", NULL);
完整main.cpp
支持两种方式:一次性加载内核源码执行,以及加载预编译的二进制文件(保护源码)。
#include <CL/cl.h>
#include <iostream>
#include <vector>
#include <fstream>
#include <string>
// 选择使用源码或二进制
const bool USE_BINARY = true; // 设置为true加载二进制,为false加载源码
const std::string kernel_source_file = "kernels/vector_add.cl";
const std::string kernel_binary_file = "vector_add.bin";
int main() {
cl_int err;
// 1. 获取平台
cl_platform_id platform;
clGetPlatformIDs(1, &platform, NULL);
// 2. 获取设备
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
// 3. 创建上下文
cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
// 4. 创建命令队列
cl_command_queue queue = clCreateCommandQueue(context, device, 0, &err);
// 5. 加载内核程序
cl_program program;
if (USE_BINARY) {
// 载入二进制
std::ifstream bin_file(kernel_binary_file, std::ios::binary);
if (!bin_file.is_open()) {
std::cerr << "Failed to open kernel binary: " << kernel_binary_file << std::endl;
return -1;
}
size_t bin_size;
bin_file.seekg(0, std::ios::end);
bin_size = bin_file.tellg();
bin_file.seekg(0, std::ios::beg);
std::vector<unsigned char> binary_data(bin_size);
bin_file.read((char*)binary_data.data(), bin_size);
bin_file.close();
const unsigned char* binaries[] = { binary_data.data() };
program = clCreateProgramWithBinary(context, 1, &device, &bin_size, binaries, NULL, &err);
if (err != CL_SUCCESS) {
std::cerr << "Failed to create program with binary" << std::endl;
return -1;
}
// 编译二进制(某些平台可能不需要)
err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
if (err != CL_SUCCESS) {
size_t log_size;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
std::vector<char> log(log_size);
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, log.data(), NULL);
std::cerr << "Build log:\n" << log.data() << std::endl;
return -1;
}
} else {
// 载入源码
std::ifstream kernel_file(kernel_source_file);
if (!kernel_file.is_open()) {
std::cerr << "Failed to open kernel file" << std::endl;
return -1;
}
std::string kernel_code((std::istreambuf_iterator<char>(kernel_file)), std::istreambuf_iterator<char>());
const char* kernel_source = kernel_code.c_str();
program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, &err);
if (err != CL_SUCCESS) {
std::cerr << "Failed to create program with source" << std::endl;
return -1;
}
// 编译程序
if (clBuildProgram(program, 1, &device, NULL, NULL, NULL) != CL_SUCCESS) {
size_t log_size;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
std::vector<char> log(log_size);
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, log.data(), NULL);
std::cerr << "Build log:\n" << log.data() << std::endl;
return -1;
}
}
// 6. 创建核函数
cl_kernel kernel = clCreateKernel(program, "vector_add", &err);
if (err != CL_SUCCESS) {
std::cerr << "Failed to create kernel" << std::endl;
return -1;
}
// 7. 准备数据
const int N = 1024;
std::vector<float> A(N, 1.0f);
三、CMakeLists.txt配置示例
cmake_minimum_required(VERSION 3.14)
project(OpenCLExample)
# 查找OpenCL
find_package(OpenCL REQUIRED)
# 添加可执行文件
add_executable(opencl_example src/main.cpp)
# 添加源码路径(要根据实际路径调整)
target_include_directories(opencl_example PRIVATE ${OPENCL_INCLUDE_DIRS})
target_link_libraries(opencl_example PRIVATE ${OPENCL_LIBRARIES})
# 复制kernel文件到构建目录(可选)
configure_file(kernels/vector_add.cl ${CMAKE_BINARY_DIR}/kernels/vector_add.cl COPYONLY)
四、总结
- 开发时:用
.cl源码工程方便调试。 - 发布时:用OpenCL API在程序中生成二进制文件,存为
.bin,并加载二进制,避免源码泄露。 - CMake主要负责布局和资源管理,不涉及二进制生成的细节,但可以利用
configure_file()复制资源。
OPENCL编译模式
在编译OpenCL程序时,特别是在使用CMake或其他构建工具,对OpenCL内核进行预编译或处理时,通常会涉及一些标志或参数。这些缩写(IL、BC、CL、CLS)代表不同的概念,主要与OpenCL内核的中间表示和二进制格式有关:
1. IL(Intermediate Language,中间语言)
- 定义:OpenCL的内核可以被编译成一种中间表示(Intermediate Language),类似于LLVM IR或SPIR-V,方便在不同硬件和驱动之间进行移植和优化。
- 作用:IL是一种平台无关的“中间码”,可以被存储和传输,然后再在目标设备上被编译成硬件特定的机器码。
- 使用场景:通常在预编译、特定平台支持或多通用硬件环境中使用。
2. BC(Binary Code,二进制代码)
- 定义:这是OpenCL内核被编译后生成的二进制格式,比如为特定GPU/CPU生成的专用机器代码。
- 区别:比中间语言更接近设备底层,可以直接加载到设备上执行。
- 用途:减少运行时编译时间,保护内核源码。
3. CL(OpenCL C代码)
- 定义:这是OpenCL标准定义的C语言风格的内核源代码(
.cl文件)。 - 处理方式:需要在运行时由OpenCL驱动或API(如
clBuildProgram())编译成设备可执行二进制或中间格式。 - 扩展:当你用“CL”作为参数或标志,通常指操作的目标是OpenCL的源代码或相关操作。
4. CLS(OpenCL C Source)
- 定义:即OpenCL的C源代码,也就是
.cl文件,纯文本的源代码形式。 - 作用:这是最原始的代码形式,用于开发和调试。
总结区别与用途:
表格
| 缩写 | 全名 | 类型 | 描述 |
|---|---|---|---|
| IL | Intermediate Language | 中间语言(中间表示) | 跨平台、中立的中间码,方便后续优化和转换 |
| BC | Binary Code | 二进制代码 | 设备专用、编译后生成的二进制形式,直接加载执行 |
| CL | OpenCL C Source | 源代码(.cl) | 开发时的源代码,需要在运行时编译成可执行形式 |
| CLS | OpenCL C Source | 与CL类似,纯文本源代码 | 通常指源代码文件,也可以理解为直接的好理解的源码文本 |
怎么在CMake或者编译流程中使用?
- 源代码(
.cl,CLS):用clCreateProgramWithSource加载,然后运行clBuildProgram编译成二进制。 - 预编译二进制(
.bc,.bin, IL):用clCreateProgramWithBinary,直接加载预先编译好的二进制(推荐保护代码,提升加载速度) - 生成IL或BC:可以在不同平台用工具链(如XLA、SPIR-V工具)或在运行时由API自动生成。
- 如果你的目标是保护源码,建议预编译成二进制(BC/IL),在加载时用
clCreateProgramWithBinary()。
RK系列芯片上怎么使用opencl
在RK(Rockchip)系列芯片上使用OpenCL,通常需要满足以下几个条件和步骤:
一、确认支持情况
- 硬件支持:
RK系列芯片(如RK3399、RK3566等)多配备ARM Mali GPU或Mali兼容的OpenCL驱动支持。 - 软件支持:
需要安装支持的OpenCL平台驱动(通常是Mali的OpenCL SDK或Rockchip的专用驱动包)。
二、获取和安装OpenCL驱动
- 获取官方驱动:
- 你可以从Rockchip官网或者芯片供应商那里获取专用的OpenCL SDK或驱动包,通常包含动态库(libOpenCL.so)、头文件和示例。
- 实现或集成驱动:
- 使用来自Rockchip的预编译OpenCL库,或者自己编译Mali的OpenCL实现(如果你拥有源代码和编译环境)。
- 配置环境:
- 将驱动库(
libOpenCL.so)放在系统库路径(如/usr/lib/或/usr/local/lib/)。 - 可能需要设置环境变量,比如
LD_LIBRARY_PATH=/usr/local/lib。
- 将驱动库(
三、使用OpenCL开发
-
编写应用程序:
- 你可以用标准OpenCL API(和之前讲解的一样)进行开发。
-
运行时环境:
- 确保运行环境中有匹配的
libOpenCL.so库,应用程序在启动时会识别到GPU的OpenCL平台。
- 确保运行环境中有匹配的
-
验证支持:
- 运行简单的OpenCL程序(如列出平台和设备)确认支持。
// 简单示例:列出平台和设备
cl_platform_id platform;
clGetPlatformIDs(1, &platform, NULL);
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
// 打印设备信息
char buffer[256];
clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(buffer), buffer, NULL);
printf("Device: %s\n", buffer);
四、特定驱动可能的注意事项
-
驱动版本兼容:
需要确保OpenCL驱动版本与你的硬件和软件环境匹配,否则可能无法正确加载或运行。 -
硬件优化:
利用Mali GPU的特性优化内核代码,以获得更好的性能。 -
调试和调优:
使用厂商提供的工具或OpenCL调试工具监测性能。
五、总结
表格
| 步骤 | 说明 |
|---|---|
| 获取驱动 | 从Rockchip或Mali提供商下载官方OpenCL SDK或驱动包 |
| 安装驱动 | 将库文件放到系统路径,配置好环境变量 |
| 开发应用 | 使用OpenCL API(如前述代码示例)编写程序 |
| 运行测试 | 运行示例程序确认支持,开始GPU加速 |

















![[蓝桥杯]图形排版](https://i-blog.csdnimg.cn/direct/b88570d063f646e7bdc5011519d44a5f.png)

