diff --git a/modcc/printer/gpuprinter.cpp b/modcc/printer/gpuprinter.cpp
index 3b431592744eb5828812e6ef99f755266de5e3b3..97154abefd9fdb0dcb42b56135f996020608bac4 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);
     }