From 9294e7d9075239b803e93e9ec5821e7dfeac14cc Mon Sep 17 00:00:00 2001
From: Nora Abi Akar <nora.abiakar@gmail.com>
Date: Mon, 13 Dec 2021 12:47:54 +0100
Subject: [PATCH] Modcc: Forward declare procedure kernels in generated GPU
code (#1787)
---
modcc/printer/gpuprinter.cpp | 14 ++++++++++++--
1 file changed, 12 insertions(+), 2 deletions(-)
diff --git a/modcc/printer/gpuprinter.cpp b/modcc/printer/gpuprinter.cpp
index 3b431592..97154abe 100644
--- a/modcc/printer/gpuprinter.cpp
+++ b/modcc/printer/gpuprinter.cpp
@@ -186,17 +186,27 @@ std::string emit_gpu_cu_source(const Module& module_, const printer_options& opt
<< "using ::arb::gpu::max;\n\n";
// Procedures as __device__ functions.
- auto emit_procedure_kernel = [&] (ProcedureExpression* e) {
+ auto emit_procedure_proto = [&] (ProcedureExpression* e) {
out << fmt::format("__device__\n"
"void {}(arb_mechanism_ppack params_, int tid_",
e->name());
for(auto& arg: e->args()) out << ", arb_value_type " << arg->is_argument()->name();
- out << ") {\n" << indent
+ out << ")";
+ };
+
+ auto emit_procedure_kernel = [&] (ProcedureExpression* e) {
+ emit_procedure_proto(e);
+ out << " {\n" << indent
<< "PPACK_IFACE_BLOCK;\n"
<< cuprint(e->body())
<< popindent << "}\n\n";
};
+ for (auto& p: module_normal_procedures(module_)) {
+ emit_procedure_proto(p);
+ out << ";\n\n";
+ }
+
for (auto& p: module_normal_procedures(module_)) {
emit_procedure_kernel(p);
}
--
GitLab