提交 7b7e0d90 编写于 作者: Y Yanzhan Yang 提交者: GitHub

accelerate opencl compilation time test=develop (#1961)

* accelerate opencl compilation time test=develop

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