未验证 提交 9f5a685a 编写于 作者: Y Yanzhan Yang 提交者: GitHub

accelerate opencl compilation time test=develop (#1961)

* accelerate opencl compilation time test=develop

* fix style test=develop
上级 d60b8d61
...@@ -29,7 +29,7 @@ limitations under the License. */ ...@@ -29,7 +29,7 @@ limitations under the License. */
namespace paddle_mobile { namespace paddle_mobile {
extern const std::map<std::string, std::vector<unsigned char>> opencl_kernels; extern const std::map<std::string, std::vector<unsigned char>> opencl_kernels;
extern const std::vector<std::string> need_conv_header_kernels; extern const std::map<std::string, std::vector<unsigned char>> opencl_headers;
namespace framework { namespace framework {
...@@ -48,7 +48,7 @@ class CLScope { ...@@ -48,7 +48,7 @@ class CLScope {
const std::string &kernel_name, const std::string &file_name, const std::string &kernel_name, const std::string &file_name,
const std::string &options) { const std::string &options) {
DLOG << " to get program " << file_name; DLOG << " to get program " << file_name;
auto program = Program(file_name, options); auto program = Program(file_name, kernel_name, options);
DLOG << " end get program ~ "; DLOG << " end get program ~ ";
DLOG << " to create kernel: " << kernel_name; DLOG << " to create kernel: " << kernel_name;
std::unique_ptr<_cl_kernel, CLKernelDeleter> kernel( std::unique_ptr<_cl_kernel, CLKernelDeleter> kernel(
...@@ -60,8 +60,12 @@ class CLScope { ...@@ -60,8 +60,12 @@ class CLScope {
cl_context Context() { return context_; } cl_context Context() { return context_; }
cl_program Program(const std::string &file_name, const std::string &options) { cl_program Program(const std::string &file_name,
std::string program_key = file_name; const std::string &kernel_name,
const std::string &options) {
if (opencl_kernels.find(kernel_name) != opencl_kernels.end() &&
opencl_headers.find(file_name) != opencl_headers.end()) {
std::string program_key = file_name + kernel_name;
if (!options.empty()) { if (!options.empty()) {
program_key += options; program_key += options;
} }
...@@ -69,17 +73,11 @@ class CLScope { ...@@ -69,17 +73,11 @@ class CLScope {
if (it != programs_.end()) { if (it != programs_.end()) {
return it->second.get(); return it->second.get();
} }
auto src_it = opencl_kernels.find(kernel_name);
if (opencl_kernels.find(file_name) != opencl_kernels.end()) { std::string source(src_it->second.begin(), src_it->second.end());
auto it = opencl_kernels.find(file_name); auto header_it = opencl_headers.find(file_name);
std::string source(it->second.begin(), it->second.end()); std::string header(header_it->second.begin(), header_it->second.end());
if (std::find(need_conv_header_kernels.begin(), source = header + "\n" + source;
need_conv_header_kernels.end(),
file_name) != need_conv_header_kernels.end()) {
auto it = opencl_kernels.find("conv_kernel.inc.cl");
std::string header(it->second.begin(), it->second.end());
source = header + source;
}
auto program = CLEngine::Instance()->CreateProgramWithSource( auto program = CLEngine::Instance()->CreateProgramWithSource(
context_, source.c_str()); context_, source.c_str());
...@@ -88,7 +86,16 @@ class CLScope { ...@@ -88,7 +86,16 @@ class CLScope {
DLOG << " --- end build program -> " << program_key << " --- "; DLOG << " --- end build program -> " << program_key << " --- ";
programs_[program_key] = std::move(program); programs_[program_key] = std::move(program);
return programs_[program_key].get();
} else { } else {
std::string program_key = file_name;
if (!options.empty()) {
program_key += options;
}
auto it = programs_.find(program_key);
if (it != programs_.end()) {
return it->second.get();
}
auto program = CLEngine::Instance()->CreateProgramWith( auto program = CLEngine::Instance()->CreateProgramWith(
context_, context_,
CLEngine::Instance()->GetCLPath() + "/cl_kernel/" + file_name); CLEngine::Instance()->GetCLPath() + "/cl_kernel/" + file_name);
...@@ -98,10 +105,9 @@ class CLScope { ...@@ -98,10 +105,9 @@ class CLScope {
DLOG << " --- end build program -> " << program_key << " --- "; DLOG << " --- end build program -> " << program_key << " --- ";
programs_[program_key] = std::move(program); programs_[program_key] = std::move(program);
}
return programs_[program_key].get(); return programs_[program_key].get();
} }
}
CLLocalWorkSizeInfo LocalWorkSizeInfo() { return localWorkSizeInfo_; } CLLocalWorkSizeInfo LocalWorkSizeInfo() { return localWorkSizeInfo_; }
......
...@@ -24,10 +24,12 @@ def gen_opencl_kernels(): ...@@ -24,10 +24,12 @@ def gen_opencl_kernels():
#include <string> #include <string>
#include <vector> #include <vector>
namespace paddle_mobile { namespace paddle_mobile {
// func name => source
extern const std::map<std::string, std::vector<unsigned char>> opencl_kernels = { extern const std::map<std::string, std::vector<unsigned char>> opencl_kernels = {
%s %s
}; };
extern const std::vector<std::string> need_conv_header_kernels = { // file name => header
extern const std::map<std::string, std::vector<unsigned char>> opencl_headers = {
%s %s
}; };
} }
...@@ -41,66 +43,143 @@ def gen_opencl_kernels(): ...@@ -41,66 +43,143 @@ def gen_opencl_kernels():
hex_list.append(hex_) hex_list.append(hex_)
return hex_list return hex_list
infile = open("cl_kernel/cl_common.h", "r") def clean_source(content):
common_content = infile.read() new_content = re.sub(r"/\*([^*]|[\r\n]|(\*+([^*/]|[\r\n])))*\*+/", "", content, flags=re.DOTALL)
infile.close() lines = new_content.split("\n")
common_content = re.sub(r"/\*[^*]*\*/", "", common_content, flags=re.DOTALL)
lines = common_content.split("\n")
new_lines = [] new_lines = []
for i in range(len(lines)): for i in range(len(lines)):
line = lines[i] line = lines[i]
line = re.sub(r"//.*$", "", line)
line = line.strip() line = line.strip()
if line == "": if line == "":
continue continue
if line.startswith("//"):
continue
line = re.sub(r"//.*$", "", line)
new_lines.append(line) new_lines.append(line)
common_content = "\n".join(new_lines) new_content = "\n".join(new_lines)
return new_content
need_conv_header_kernels = [] infile = open("cl_kernel/cl_common.h", "r")
common_content = infile.read()
infile.close()
common_content = clean_source(common_content)
infile = open("cl_kernel/conv_kernel.inc.cl", "r")
inc_content = infile.read()
infile.close()
inc_content = clean_source(inc_content)
def get_header_raw(content):
lines = content.split("\n")
new_lines = []
for line in lines:
if "__kernel void" in line:
break
new_lines.append(line)
header = "\n".join(new_lines)
return header
common_header = get_header_raw(common_content)
inc_header = get_header_raw(inc_content)
def get_header(content):
lines = content.split("\n")
new_lines = []
for line in lines:
if "__kernel void" in line:
break
new_lines.append(line)
for i in range(len(new_lines)):
if "#include \"conv_kernel.inc.cl\"" in new_lines[i]:
new_lines[i] = inc_header
header = "\n".join(new_lines)
new_lines = header.split("\n")
for i in range(len(new_lines)):
if "#include \"cl_common.h\"" in new_lines[i]:
new_lines[i] = common_header
header = "\n".join(new_lines)
return header
def get_funcs(content):
funcs = {}
lines = content.split("\n")
first_kernel_idx = None
for i in range(len(lines)):
if "__kernel void" in lines[i]:
first_kernel_idx = i
break
if first_kernel_idx is None:
return funcs
lines = lines[first_kernel_idx:]
func = []
name = ""
for line in lines:
if "__kernel void" in line:
if name != "":
funcs[name] = "\n".join(func)
name = ""
func = []
pattern = re.compile("__kernel void ([^(]+)\(")
match = pattern.search(line)
name = match.group(1)
func.append(line)
if name != "":
funcs[name] = "\n".join(func)
name = ""
func = []
return funcs
cores = ""
filenames = os.listdir("cl_kernel") filenames = os.listdir("cl_kernel")
file_count = len(filenames) file_count = len(filenames)
headers = {}
funcs = {}
for i in range(file_count): for i in range(file_count):
filename = filenames[i] filename = filenames[i]
infile = open("cl_kernel/" + filename, "r") infile = open("cl_kernel/" + filename, "r")
new_lines = []
content = infile.read() content = infile.read()
content = re.sub(r"/\*[^*]*\*/", "", content, flags=re.DOTALL)
infile.close() infile.close()
lines = content.split("\n") content = clean_source(content)
for i in range(len(lines)): header = get_header(content)
line = lines[i] headers[filename] = header
line = line.strip() funcs_temp = get_funcs(content)
if line == "": for key in funcs_temp:
continue funcs[key] = funcs_temp[key]
if line.startswith("//"):
continue core1 = ""
line = re.sub(r"//.*$", "", line) core2 = ""
if "cl_common.h" in line:
line = common_content for i in range(len(funcs)):
elif "conv_kernel.inc.cl" in line: func_name = list(funcs.keys())[i]
need_conv_header_kernels.append("\"%s\"" % filename) content = funcs[func_name]
continue
new_lines.append(line)
content = "\n".join(new_lines)
if content == "": if content == "":
content = " " content = " "
hexes = [] hexes = []
for char in content: for char in content:
hexes.append(hex(ord(char))) hexes.append(hex(ord(char)))
core = " {\"%s\", {" % filename core = " {\"%s\", {" % func_name
for item in hexes: for item in hexes:
core += str(item) + ", " core += str(item) + ", "
core = core[: -2] core = core[: -2]
core += "}}" core += "}}"
if i != file_count - 1: if i != len(funcs) - 1:
core += ",\n" core += ",\n"
cores += core core1 += core
source = source % (cores, ",".join(need_conv_header_kernels)) for i in range(len(headers)):
file_name = list(headers.keys())[i]
content = headers[file_name]
if content == "":
content = " "
hexes = []
for char in content:
hexes.append(hex(ord(char)))
core = " {\"%s\", {" % file_name
for item in hexes:
core += str(item) + ", "
core = core[: -2]
core += "}}"
if i != len(headers) - 1:
core += ",\n"
core2 += core
source = source % (core1, core2)
print(source) print(source)
def gen_empty_opencl_kernels(): def gen_empty_opencl_kernels():
...@@ -111,9 +190,11 @@ def gen_empty_opencl_kernels(): ...@@ -111,9 +190,11 @@ def gen_empty_opencl_kernels():
#include <string> #include <string>
#include <vector> #include <vector>
namespace paddle_mobile { namespace paddle_mobile {
// func name => source
extern const std::map<std::string, std::vector<unsigned char>> opencl_kernels = { extern const std::map<std::string, std::vector<unsigned char>> opencl_kernels = {
}; };
extern const std::vector<std::string> need_conv_header_kernels = { // file name => header
extern const std::map<std::string, std::vector<unsigned char>> opencl_headers = {
}; };
} }
#endif #endif
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册