diff --git a/AdaptiveCpp b/AdaptiveCpp
new file mode 160000
index 0000000000..be5d2d8117
--- /dev/null
+++ b/AdaptiveCpp
@@ -0,0 +1 @@
+Subproject commit be5d2d8117d21f19aaf54a9b72de11771669339f
diff --git a/CMakeLists.txt b/CMakeLists.txt
index efad28e458..73d60793ef 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -27,6 +27,34 @@ set(CMAKE_LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}"
 option(BUILD_SHARED_LIBS    "Build shared libraries." ON)
 option(MIM_BUILD_DOCS       "If ON, build the documentation (requires Doxygen)." OFF)
 option(MIM_BUILD_EXAMPLES   "If ON, build examples." OFF)
+option(MIM_WITH_ADAPTIVECPP "If ON, locate AdaptiveCpp (acpp-rt + acpp-common) so the pCUDA backend can shell to llvm-link with sscp_stubs.ll, and so users can link MimIR-emitted host LL with the SSCP runtime." ON)
+
+# AdaptiveCpp discovery for the pCUDA backend in the nvptx plugin. The plugin
+# itself does NOT link against acpp-rt at build time (it only shells out at
+# emit time and bakes a stubs path into a configured header), but downstream
+# user binaries built from MimIR-emitted .ll do need the runtime libs.
+if(MIM_WITH_ADAPTIVECPP)
+    # Default hint: the in-tree AdaptiveCpp build install dir, if present.
+    set(_acpp_in_tree "${CMAKE_CURRENT_LIST_DIR}/AdaptiveCpp/build/install")
+    if(EXISTS "${_acpp_in_tree}/lib/cmake/AdaptiveCpp/adaptivecpp-config.cmake")
+        list(APPEND CMAKE_PREFIX_PATH "${_acpp_in_tree}")
+    endif()
+    find_package(AdaptiveCpp QUIET)
+    if(AdaptiveCpp_FOUND)
+        # The find_package config sets PACKAGE_PREFIX_DIR locally; derive a
+        # globally-available install dir from the imported acpp-rt location.
+        get_target_property(_acpp_rt_loc AdaptiveCpp::acpp-rt LOCATION)
+        if(_acpp_rt_loc)
+            get_filename_component(MIM_ACPP_LIB_DIR     "${_acpp_rt_loc}" DIRECTORY)
+            get_filename_component(MIM_ACPP_INSTALL_DIR "${MIM_ACPP_LIB_DIR}/.." ABSOLUTE)
+            message(STATUS "AdaptiveCpp found: ${MIM_ACPP_INSTALL_DIR}")
+        else()
+            message(WARNING "AdaptiveCpp_FOUND but AdaptiveCpp::acpp-rt LOCATION not set; pCUDA backend will fall back to cwd-relative stubs path.")
+        endif()
+    else()
+        message(STATUS "AdaptiveCpp not found; pCUDA backend will use cwd-relative scripts/sscp_stubs.ll.")
+    endif()
+endif()
 
 message(STATUS "Build type: ${CMAKE_BUILD_TYPE}; shared libs: ${BUILD_SHARED_LIBS}")
 
diff --git a/include/mim/plug/pcuda/be/hcf_adapter.h b/include/mim/plug/pcuda/be/hcf_adapter.h
new file mode 100644
index 0000000000..ba6ff8f0cb
--- /dev/null
+++ b/include/mim/plug/pcuda/be/hcf_adapter.h
@@ -0,0 +1,74 @@
+#pragma once
+
+#include <cstddef>
+#include <cstdint>
+#include <string>
+#include <vector>
+
+namespace mim::ll::pcuda {
+
+/// HCF (Heterogeneous Container Format) builder for the pCUDA backend.
+///
+/// Produces the exact text-plus-binary-appendix format that AdaptiveCpp's
+/// runtime parses via `hipsycl::common::hcf_container::parse`, mirroring the
+/// SSCP plugin's `generateHCF` (TargetSeparationPass.cpp:434-516). The output
+/// is embedded into the host LLVM IR as `@__acpp_local_sscp_hcf_content`.
+///
+/// Layout:
+///   root: object-id, generator
+///     images / llvm-ir.global: variant=global-module, format=llvm-ir,
+///                              exported-symbols (list), imported-symbols (list),
+///                              __binary { start, size }
+///     kernels / <name>: image-providers, host-side-parameter-sizes,
+///                       compile-flags, compile-options,
+///                       parameters / <i>: byte-offset, byte-size,
+///                                         original-index, type, annotations
+///   __acpp_hcf_binary_appendix<raw device bitcode bytes>
+
+enum class HCFParamType { Integer, FloatingPoint, Pointer, OtherByValue };
+
+struct HCFParam {
+    std::size_t byte_offset;
+    std::size_t byte_size;
+    std::size_t original_index;
+    HCFParamType type;
+    std::vector<std::string> annotations;
+};
+
+struct HCFKernel {
+    std::string name;
+    std::vector<std::size_t> host_side_parameter_sizes;
+    std::vector<HCFParam> parameters;
+};
+
+class HCFBuilder {
+public:
+    HCFBuilder() = default;
+
+    void set_object_id(std::uint64_t id) { object_id_ = id; }
+    void set_generator(std::string s) { generator_ = std::move(s); }
+
+    /// Raw LLVM bitcode bytes (output of llvm-as on the device .ll).
+    void set_device_bitcode(std::string bytes) { device_bitcode_ = std::move(bytes); }
+
+    void set_exported_symbols(std::vector<std::string> syms) { exported_ = std::move(syms); }
+    void set_imported_symbols(std::vector<std::string> syms) { imported_ = std::move(syms); }
+
+    void add_kernel(HCFKernel k) { kernels_.push_back(std::move(k)); }
+
+    std::uint64_t object_id() const { return object_id_; }
+    const std::vector<HCFKernel>& kernels() const { return kernels_; }
+
+    /// Serialize to the wire format that the AdaptiveCpp runtime accepts.
+    std::string serialize() const;
+
+private:
+    std::uint64_t object_id_ = 0;
+    std::string generator_ = "MimIR pCUDA backend";
+    std::string device_bitcode_;
+    std::vector<std::string> exported_;
+    std::vector<std::string> imported_;
+    std::vector<HCFKernel> kernels_;
+};
+
+} // namespace mim::ll::pcuda
diff --git a/include/mim/plug/pcuda/be/hcf_container.hpp b/include/mim/plug/pcuda/be/hcf_container.hpp
new file mode 100644
index 0000000000..f868605845
--- /dev/null
+++ b/include/mim/plug/pcuda/be/hcf_container.hpp
@@ -0,0 +1,244 @@
+// SPDX-License-Identifier: BSD-2-Clause AND MIT
+//
+// Vendored from AdaptiveCpp (BSD 2-Clause):
+//   AdaptiveCpp/include/hipSYCL/common/hcf_container.hpp
+// Copyright The AdaptiveCpp Contributors.
+//
+// MimIR-local modifications: debug-logging macros stubbed to std::cerr to drop
+// the dependency on AdaptiveCpp's runtime/application.hpp. Format wire-compat
+// with libacpp-rt's hcf_container parser is preserved verbatim — node start/end
+// markers, binary appendix sentinel, key=value lines, all unchanged.
+//
+// AdaptiveCpp BSD 2-Clause license terms apply to this file in addition to
+// MimIR's MIT license. See AdaptiveCpp/LICENSE for the full BSD text.
+
+#pragma once
+
+#include <algorithm>
+#include <cassert>
+#include <cstddef>
+#include <iostream>
+#include <locale>
+#include <sstream>
+#include <stdexcept>
+#include <string>
+#include <utility>
+#include <vector>
+
+namespace mim::ll::pcuda::hcf {
+
+class hcf_container {
+public:
+    struct node {
+        std::vector<std::pair<std::string, std::string>> key_value_pairs;
+        std::vector<node> subnodes;
+        std::string node_id;
+
+        const node* get_subnode(const std::string& name) const {
+            for (size_t i = 0; i < subnodes.size(); ++i)
+                if (subnodes[i].node_id == name) return &subnodes[i];
+            return nullptr;
+        }
+        node* get_subnode(const std::string& name) {
+            for (size_t i = 0; i < subnodes.size(); ++i)
+                if (subnodes[i].node_id == name) return &subnodes[i];
+            return nullptr;
+        }
+        const std::string* get_value(const std::string& key) const {
+            for (size_t i = 0; i < key_value_pairs.size(); ++i)
+                if (key_value_pairs[i].first == key) return &key_value_pairs[i].second;
+            return nullptr;
+        }
+        bool has_key(const std::string& key) const { return get_value(key) != nullptr; }
+        bool has_subnode(const std::string& name) const { return get_subnode(name) != nullptr; }
+        std::vector<std::string> get_subnodes() const {
+            std::vector<std::string> r;
+            for (const auto& s : subnodes) r.push_back(s.node_id);
+            return r;
+        }
+        bool has_binary_data_attached() const { return has_subnode("__binary"); }
+        bool is_binary_content() const { return node_id == "__binary"; }
+
+        node* add_subnode(const std::string& unique_name) {
+            for (size_t i = 0; i < subnodes.size(); ++i) {
+                if (subnodes[i].node_id == unique_name) {
+                    std::cerr << "hcf: Subnode already exists with name " << unique_name << "\n";
+                    return nullptr;
+                }
+            }
+            node n;
+            n.node_id = unique_name;
+            subnodes.push_back(n);
+            return &subnodes.back();
+        }
+        void set(const std::string& key, const std::string& value) {
+            key_value_pairs.emplace_back(key, value);
+        }
+        void set_as_list(const std::string& key, const std::vector<std::string>& list_entries) {
+            auto* N = add_subnode(key);
+            if (N)
+                for (const auto& e : list_entries) N->add_subnode(e);
+        }
+        std::vector<std::string> get_as_list(const std::string& key) const {
+            if (!has_subnode(key)) return {};
+            return get_subnode(key)->get_subnodes();
+        }
+    };
+
+    hcf_container() { _root_node.node_id = "root"; }
+
+    explicit hcf_container(const std::string& container) {
+        std::string appendix_id{_binary_appendix_id};
+        std::size_t appendix_begin = container.find(appendix_id);
+        if (appendix_begin != std::string::npos)
+            _binary_appendix = container.substr(appendix_begin + appendix_id.length());
+        std::string parseable_data = container;
+        if (appendix_begin != std::string::npos) parseable_data.erase(appendix_begin);
+        parse(parseable_data);
+    }
+
+    const node* root_node() const { return &_root_node; }
+    node* root_node() { return &_root_node; }
+
+    bool get_binary_attachment(const node* n, std::string& out) const {
+        if (!n) return false;
+        const node* descriptor_node = nullptr;
+        if (n->is_binary_content()) descriptor_node = n;
+        else if (n->has_binary_data_attached()) descriptor_node = n->get_subnode("__binary");
+        else {
+            std::cerr << "hcf: Node " << n->node_id
+                      << " is not a binary content node, nor does it carry a binary attachment\n";
+            return false;
+        }
+        assert(descriptor_node);
+        const std::string* start_entry = descriptor_node->get_value("start");
+        const std::string* size_entry  = descriptor_node->get_value("size");
+        if (!start_entry || !size_entry) {
+            std::cerr << "hcf: Node missing binary start/size\n";
+            return false;
+        }
+        std::size_t start = std::stoull(*start_entry);
+        std::size_t size  = std::stoull(*size_entry);
+        if (start + size > _binary_appendix.size()) {
+            std::cerr << "hcf: Binary content address is out-of-bounds\n";
+            return false;
+        }
+        out = _binary_appendix.substr(start, size);
+        return true;
+    }
+
+    bool attach_binary_content(node* n, const std::string& binary_content) {
+        node* binary_node = n->add_subnode(_binary_marker);
+        if (!binary_node) return false;
+        std::size_t start  = _binary_appendix.size();
+        std::size_t length = binary_content.size();
+        _binary_appendix += binary_content;
+        binary_node->set("start", std::to_string(start));
+        binary_node->set("size", std::to_string(length));
+        return true;
+    }
+
+    std::string serialize() const {
+        std::stringstream sstr;
+        serialize_node(_root_node, sstr);
+        sstr << _binary_appendix_id;
+        return sstr.str() + _binary_appendix;
+    }
+
+private:
+    void serialize_node(const node& n, std::ostream& out) const {
+        for (const auto& p : n.key_value_pairs) out << p.first << "=" << p.second << "\n";
+        for (const auto& s : n.subnodes) {
+            out << _node_start_id << s.node_id << "\n";
+            serialize_node(s, out);
+            out << _node_end_id << s.node_id << "\n";
+        }
+    }
+
+    static void trim_left(std::string& s) {
+        auto it = std::find_if(s.begin(), s.end(),
+                               [](char ch) { return !std::isspace<char>(ch, std::locale::classic()); });
+        s.erase(s.begin(), it);
+    }
+    static void trim_right(std::string& s) {
+        auto it = std::find_if(s.rbegin(), s.rend(),
+                               [](char ch) { return !std::isspace<char>(ch, std::locale::classic()); });
+        s.erase(it.base(), s.end());
+    }
+
+    bool parse_node_start(const std::string& line, std::string& node_id) const {
+        std::string proc = line;
+        std::string ns{_node_start_id};
+        trim_left(proc);
+        if (proc.find(ns) != 0) {
+            std::cerr << "hcf: Invalid node start: " << proc << "\n";
+            return false;
+        }
+        proc.erase(0, ns.length());
+        trim_right(proc);
+        node_id = proc;
+        return true;
+    }
+
+    bool parse_node_interior(const std::vector<std::string>& lines,
+                             std::size_t node_start_line, std::size_t node_end_line,
+                             node& current) const {
+        if (node_start_line == node_end_line) return true;
+        for (std::size_t i = node_start_line; i < node_end_line; ++i) {
+            assert(i < lines.size());
+            const std::string& cur = lines[i];
+            if (cur.find(_node_start_id) == 0) {
+                node nn;
+                if (!parse_node_start(cur, nn.node_id)) return false;
+                std::size_t num_lines = std::string::npos;
+                std::string end_marker = std::string{_node_end_id} + nn.node_id;
+                for (std::size_t j = i + 1; j < node_end_line; ++j) {
+                    if (lines[j] == end_marker) {
+                        num_lines = j - i;
+                        break;
+                    }
+                }
+                if (num_lines == std::string::npos) {
+                    std::cerr << "hcf: Syntax error: missing node end marker: " << end_marker << "\n";
+                    return false;
+                }
+                if (!parse_node_interior(lines, i + 1, i + num_lines, nn)) return false;
+                current.subnodes.push_back(nn);
+                i += num_lines;
+            } else if (cur.find('=') != std::string::npos) {
+                std::size_t pos = cur.find('=');
+                current.key_value_pairs.emplace_back(cur.substr(0, pos), cur.substr(pos + 1));
+            } else if (cur.find(_node_end_id) == 0) {
+                std::cerr << "hcf: Syntax error: unexpected node end: " << cur << "\n";
+                return false;
+            } else {
+                std::cerr << "hcf: Syntax error: invalid line: " << cur << "\n";
+                return false;
+            }
+        }
+        return true;
+    }
+
+    bool parse(const std::string& data) {
+        std::stringstream ss(data);
+        std::string line;
+        std::vector<std::string> lines;
+        while (std::getline(ss, line)) {
+            trim_left(line);
+            trim_right(line);
+            if (!line.empty()) lines.push_back(line);
+        }
+        _root_node.node_id = "root";
+        return parse_node_interior(lines, 0, lines.size(), _root_node);
+    }
+
+    static constexpr char _binary_appendix_id[] = "__acpp_hcf_binary_appendix";
+    static constexpr char _node_start_id[]      = "{.";
+    static constexpr char _node_end_id[]        = "}.";
+    static constexpr char _binary_marker[]      = "__binary";
+
+    node _root_node;
+    std::string _binary_appendix;
+};
+
+} // namespace mim::ll::pcuda::hcf
diff --git a/include/mim/plug/pcuda/be/ll_pcuda.h b/include/mim/plug/pcuda/be/ll_pcuda.h
new file mode 100644
index 0000000000..305cb78035
--- /dev/null
+++ b/include/mim/plug/pcuda/be/ll_pcuda.h
@@ -0,0 +1,16 @@
+#pragma once
+
+#include "mim/plug/core/be/ll.h"
+
+namespace mim {
+
+class World;
+
+namespace ll::pcuda {
+
+void emit_host(World&, std::ostream&);
+void emit_device(World&, std::ostream&);
+void emit_host_with_embedded_device(World&, std::ostream&);
+
+} // namespace ll::pcuda
+} // namespace mim
diff --git a/include/mim/plug/pcuda/pcuda.h b/include/mim/plug/pcuda/pcuda.h
new file mode 100644
index 0000000000..a602019d90
--- /dev/null
+++ b/include/mim/plug/pcuda/pcuda.h
@@ -0,0 +1,3 @@
+#pragma once
+
+#include "mim/plug/pcuda/autogen.h"
diff --git a/lit/gpu/array_kernel.mim b/lit/gpu/array_kernel.mim
index 8344c6498c..fbc1683a8e 100644
--- a/lit/gpu/array_kernel.mim
+++ b/lit/gpu/array_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/const_mem_global_symbol_kernel.mim b/lit/gpu/const_mem_global_symbol_kernel.mim
index d979ff299a..67cb295e74 100644
--- a/lit/gpu/const_mem_global_symbol_kernel.mim
+++ b/lit/gpu/const_mem_global_symbol_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/const_mem_kernel.mim b/lit/gpu/const_mem_kernel.mim
index b49172e45b..9155692b2e 100644
--- a/lit/gpu/const_mem_kernel.mim
+++ b/lit/gpu/const_mem_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/dyn_smem_kernel.mim b/lit/gpu/dyn_smem_kernel.mim
index 728638f8ad..71cbe5d98b 100644
--- a/lit/gpu/dyn_smem_kernel.mim
+++ b/lit/gpu/dyn_smem_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/generic_kernel.mim b/lit/gpu/generic_kernel.mim
index 14ef5307ab..abb478067f 100644
--- a/lit/gpu/generic_kernel.mim
+++ b/lit/gpu/generic_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/global_symbol_kernel.mim b/lit/gpu/global_symbol_kernel.mim
index 29ce8c8319..48bf6e3efd 100644
--- a/lit/gpu/global_symbol_kernel.mim
+++ b/lit/gpu/global_symbol_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/malloc_kernel.mim b/lit/gpu/malloc_kernel.mim
index fb55c63dad..6edcb6309d 100644
--- a/lit/gpu/malloc_kernel.mim
+++ b/lit/gpu/malloc_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/manual_streams.mim b/lit/gpu/manual_streams.mim
index a6f4cc97de..c967d4a4ff 100644
--- a/lit/gpu/manual_streams.mim
+++ b/lit/gpu/manual_streams.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/manual_sync_kernel.mim b/lit/gpu/manual_sync_kernel.mim
index d1e09c2579..457118f40e 100644
--- a/lit/gpu/manual_sync_kernel.mim
+++ b/lit/gpu/manual_sync_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/mem_arg_kernel.mim b/lit/gpu/mem_arg_kernel.mim
index 7b9d7ae5b5..b87aa8ce07 100644
--- a/lit/gpu/mem_arg_kernel.mim
+++ b/lit/gpu/mem_arg_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s; test $? -ne 0
 
diff --git a/lit/gpu/no_arg_kernel.mim b/lit/gpu/no_arg_kernel.mim
index fb7f952f90..6ffa9ceef0 100644
--- a/lit/gpu/no_arg_kernel.mim
+++ b/lit/gpu/no_arg_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s -O 0
 
diff --git a/lit/gpu/no_gpu_deinit.mim b/lit/gpu/no_gpu_deinit.mim
index eb671e6311..58920e865c 100644
--- a/lit/gpu/no_gpu_deinit.mim
+++ b/lit/gpu/no_gpu_deinit.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s; test $? -ne 0
 
diff --git a/lit/gpu/no_gpu_init.mim b/lit/gpu/no_gpu_init.mim
index bdf0a3f3f0..ed65366037 100644
--- a/lit/gpu/no_gpu_init.mim
+++ b/lit/gpu/no_gpu_init.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s; test $? -ne 0
 
diff --git a/lit/gpu/scope_streams.mim b/lit/gpu/scope_streams.mim
index 9339f8b0d9..c275672fe7 100644
--- a/lit/gpu/scope_streams.mim
+++ b/lit/gpu/scope_streams.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/scope_sync_kernel.mim b/lit/gpu/scope_sync_kernel.mim
index 99cebcafcc..4dc071e207 100644
--- a/lit/gpu/scope_sync_kernel.mim
+++ b/lit/gpu/scope_sync_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/simple_kernel.mim b/lit/gpu/simple_kernel.mim
index aaa2dabd44..b51d7aafd3 100644
--- a/lit/gpu/simple_kernel.mim
+++ b/lit/gpu/simple_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/gpu/two_arg_kernel.mim b/lit/gpu/two_arg_kernel.mim
index dbf474ebcc..8c80b6fd5d 100644
--- a/lit/gpu/two_arg_kernel.mim
+++ b/lit/gpu/two_arg_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/lit.cfg.py b/lit/lit.cfg.py
index 07b8416b81..629228b040 100644
--- a/lit/lit.cfg.py
+++ b/lit/lit.cfg.py
@@ -1,4 +1,5 @@
 import lit.formats
+import shutil
 import os
 
 config.name = 'mim regression'
@@ -11,6 +12,16 @@
 
 config.substitutions.append(('%mim', config.mim))
 
+acpp_lib_dir = getattr(config, 'acpp_lib_dir', '')
+if acpp_lib_dir:
+    config.substitutions.append(
+        ('%acpp_link',
+         '-L{0} -lacpp-rt -lacpp-common -Wl,-rpath={0} -pthread -ldl'.format(acpp_lib_dir)))
+    config.available_features.add("pcuda")
+
+if shutil.which("nvidia-smi") is not None:
+    config.available_features.add("nvptx")
+
 # inhert env vars
 config.environment = os.environ
 
diff --git a/lit/lit.site.cfg.py.in b/lit/lit.site.cfg.py.in
index 0e909371a2..a29e2a780c 100644
--- a/lit/lit.site.cfg.py.in
+++ b/lit/lit.site.cfg.py.in
@@ -3,6 +3,7 @@ import sys
 
 config.my_src_root = r'@CMAKE_SOURCE_DIR@'
 config.my_obj_root = r'@CMAKE_BINARY_DIR@'
+config.acpp_lib_dir = r'@MIM_ACPP_LIB_DIR@'
 if sys.platform == "win32":
     config.mim = r'@CMAKE_BINARY_DIR@/bin/mim.exe'
 else:
diff --git a/lit/nvptx/coalesced_gemm.mim b/lit/nvptx/coalesced_gemm.mim
index b5bd15a313..1fb978b5de 100644
--- a/lit/nvptx/coalesced_gemm.mim
+++ b/lit/nvptx/coalesced_gemm.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/nvptx/softmax_kernel.mim b/lit/nvptx/softmax_kernel.mim
index f2493dd041..25f16d07e7 100644
--- a/lit/nvptx/softmax_kernel.mim
+++ b/lit/nvptx/softmax_kernel.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/nvptx/unoptimized_gemm.mim b/lit/nvptx/unoptimized_gemm.mim
index e13db288af..653f527a14 100644
--- a/lit/nvptx/unoptimized_gemm.mim
+++ b/lit/nvptx/unoptimized_gemm.mim
@@ -1,3 +1,4 @@
+// REQUIRES: nvptx
 // RUN: rm -f %t.ll
 // RUN: %mim %s --output-ll %t.ll --device-target nvptx -o -
 // RUN: rm -f %t
diff --git a/lit/pcuda/array_kernel.mim b/lit/pcuda/array_kernel.mim
new file mode 100644
index 0000000000..d98e066f06
--- /dev/null
+++ b/lit/pcuda/array_kernel.mim
@@ -0,0 +1,53 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 42
+// RUN: %t 1; test $? -eq 43
+// RUN: %t 1 2; test $? -eq 44
+// RUN: %t 1 2 3; test $? -eq 45
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), data: %gpu.GlobalPtr «4; I32»,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let el = %mem.lea (data, item_id);
+    let val = %core.bitcast I32 item_id;
+    let val = %core.wrap.add 0 (val, 42I32);
+    let m1 = %mem.store (m1, el, val);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, arr) = %mem.alloc («4; I32», 0) m0;
+
+    let arr_0 = %mem.lea (arr, 0_4);
+    let m0 = %mem.store (m0, arr_0, argc);
+    let arr_1 = %mem.lea (arr, 1_4);
+    let m0 = %mem.store (m0, arr_1, argc);
+    let arr_2 = %mem.lea (arr, 2_4);
+    let m0 = %mem.store (m0, arr_2, argc);
+    let arr_3 = %mem.lea (arr, 3_4);
+    let m0 = %mem.store (m0, arr_3, argc);
+
+    let (m0, m1, m4, (), ()) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr) = %mem.alloc («4; I32», 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, arr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, arr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let idx = %core.bitcast (Idx 4) (%core.wrap.sub 0 (argc, 1I32));
+    let val_arr = %mem.lea (arr, idx);
+    let (m0, val) = %mem.load (m0, val_arr);
+    let m0 = %mem.free (m0, arr);
+
+    return (m0, val);
diff --git a/lit/pcuda/coalesced_gemm.mim b/lit/pcuda/coalesced_gemm.mim
new file mode 100644
index 0000000000..1a46d8912b
--- /dev/null
+++ b/lit/pcuda/coalesced_gemm.mim
@@ -0,0 +1,154 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+
+import gpu;
+import pcuda;
+import core;
+import libc;
+import math;
+import mem;
+import print;
+
+let n = 32768;
+let size = %core.nat.mul (n, n);
+let n_i32 = %core.bitcast I32 n;
+
+let numBlocks = 1048576;
+let numThreads = 1024;
+
+// GEMM of matrix dimensions:  nxn * nxn -> nxn
+con gemm_kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+                 group_id: Idx numBlocks, item_id: Idx numThreads, (), (),
+                 (out: %gpu.GlobalPtr «size; %math.F32»,
+                 in1: %gpu.GlobalPtr «size; %math.F32», in2: %gpu.GlobalPtr «size; %math.F32»),
+                 return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let rid = %core.bitcast I32 group_id;
+    let cid = %core.bitcast I32 item_id;
+
+    let id = %core.bitcast I32 (
+        %core.wrap.add 0 (%core.wrap.mul 0 (rid, n_i32), cid)
+    );
+
+    let row_oob = %core.icmp.sge (rid, n_i32);
+    let col_oob = %core.icmp.sge (cid, n_i32);
+    let tot_oob = %core.icmp.sge (id, %core.bitcast I32 (size));
+
+    let is_idle_thread = %core.bit2.or_ 0 (%core.bit2.or_ 0 (row_oob, col_oob), tot_oob);
+    (continue, return)#is_idle_thread (m1, m3, m4, m5)
+    where
+        con continue (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM) =
+            let id = %core.bitcast (Idx size) id;
+
+            let fp_zero = %math.conv.s2f %math.f32 0I32;
+            wrapped_loop (m1, 0I32, fp_zero)
+            where
+                con wrapped_loop (m1: %gpu.GlobalM, i: I32, acc: %math.F32) =
+                    let cond = %core.icmp.sl (i, n_i32);
+                    let next_i = %core.wrap.add 0 (i, 1I32);
+                    (write_and_exit, body)#cond m1
+                    where
+                        con body (m1: %gpu.GlobalM) =
+                            let in1_idx = %core.bitcast (Idx size) (
+                                %core.wrap.add 0 (%core.wrap.mul 0 (rid, n_i32), i)
+                            );
+                            let in2_idx = %core.bitcast (Idx size) (
+                                %core.wrap.add 0 (%core.wrap.mul 0 (i, n_i32), cid)
+                            );
+                            let in1_addr = %mem.lea (in1, in1_idx);
+                            let in2_addr = %mem.lea (in2, in2_idx);
+                            let (m1, in1_val) = %mem.load (m1, in1_addr);
+                            let (m1, in2_val) = %mem.load (m1, in2_addr);
+                            let new_acc = %pcuda.fmaf (in1_val, in2_val, acc);
+                            wrapped_loop (m1, next_i, new_acc);
+                        con write_and_exit (m1: %gpu.GlobalM) =
+                            let out_addr = %mem.lea (out, id);
+                            let m1 = %mem.store (m1, out_addr, acc);
+                            return (m1, m3, m4, m5);
+                    end;
+            end;
+    end;
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0)) =
+    let (m0, h_in1) = %mem.alloc («size; %math.F32», 0) m0;
+    let (m0, h_in2) = %mem.alloc («size; %math.F32», 0) m0;
+    let (m0, h_out) = %mem.alloc («size; %math.F32», 0) m0;
+
+    // generate ints in [0, cutoff), convert them to floats in [a, b), and fill h_in1 and h_in2 with these values
+    let seed = 42I32;
+    let m0 = %libc.srand (m0, seed);
+    let cutoff = 1000000I32;
+    let a = %math.minus 0 (%math.conv.s2f %math.f32 50I32);
+    let b = %math.conv.s2f %math.f32 50I32;
+    let div = %math.arith.div 0 (%math.conv.s2f %math.f32 cutoff, %math.arith.sub 0 (b, a));
+
+    fun loop_in1 (m0: %mem.M 0) : %mem.M 0 =
+        wrapped_loop (m0, 0I32, return)
+        where
+            con wrapped_loop (m0: %mem.M 0, i: I32, exit: Cn [%mem.M 0]) =
+                let cond = %core.icmp.sl (i, %core.bitcast I32 size);
+                let next_i = %core.wrap.add 0 (i, 1I32);
+                (exit, body)#cond m0
+                where
+                    con body (m0: %mem.M 0) =
+                        let idx = %core.bitcast (Idx size) i;
+                        let addr = %mem.lea (h_in1, idx);
+                        let (m0, value_i32) = %libc.rand m0;
+                        let (m0, value_i32) = %core.div.srem (m0, (value_i32, cutoff));
+                        let value = %math.conv.s2f %math.f32 value_i32;
+                        let value = %math.arith.div 0 (value, div);
+                        let value = %math.arith.add 0 (value, a);
+                        let m0 = %mem.store (m0, addr, value);
+                        wrapped_loop (m0, next_i, exit);
+                end;
+        end;
+    ret m0 = loop_in1 $ m0;
+
+    fun loop_in2 (m0: %mem.M 0) : %mem.M 0 =
+        wrapped_loop (m0, 0I32, return)
+        where
+            con wrapped_loop (m0: %mem.M 0, i: I32, exit: Cn [%mem.M 0]) =
+                let cond = %core.icmp.sl (i, %core.bitcast I32 size);
+                let next_i = %core.wrap.add 0 (i, 1I32);
+                (exit, body)#cond m0
+                where
+                    con body (m0: %mem.M 0) =
+                        let idx = %core.bitcast (Idx size) i;
+                        let addr = %mem.lea (h_in2, idx);
+                        let (m0, value_i32) = %libc.rand m0;
+                        let (m0, value_i32) = %core.div.srem (m0, (value_i32, cutoff));
+                        let value = %math.conv.s2f %math.f32 value_i32;
+                        let value = %math.arith.div 0 (value, div);
+                        let value = %math.arith.add 0 (value, a);
+                        let m0 = %mem.store (m0, addr, value);
+                        wrapped_loop (m0, next_i, exit);
+                end;
+        end;
+    ret m0 = loop_in2 $ m0;
+
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m0, m1, d_in1) = %gpu.alloc_copy.block (m0, m1, h_in1);
+    let (m0, m1, d_in2) = %gpu.alloc_copy.block (m0, m1, h_in2);
+    let (m1, d_out) = %mem.alloc («size; %math.F32», %gpu.addr_space_global) m1;
+
+    ret m0 = %gpu.launch (numBlocks, numThreads, %gpu.default_stream, ff, 0, (), (), ())
+             gemm_kernel (d_out, d_in1, d_in2) $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_out, h_out);
+    let m1 = %gpu.free.block (m1, d_in1);
+    let m1 = %gpu.free.block (m1, d_in2);
+    let m1 = %gpu.free.block (m1, d_out);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    ret m0 = %print.arr_f32 $ (m0, size, h_out, 10I32);
+
+    let m0 = %mem.free (m0, h_in1);
+    let m0 = %mem.free (m0, h_in2);
+    let m0 = %mem.free (m0, h_out);
+
+    return m0;
diff --git a/lit/pcuda/const_mem_global_symbol_kernel.mim b/lit/pcuda/const_mem_global_symbol_kernel.mim
new file mode 100644
index 0000000000..e4198aca7d
--- /dev/null
+++ b/lit/pcuda/const_mem_global_symbol_kernel.mim
@@ -0,0 +1,54 @@
+// pCUDA backend: global/constant memory symbol pointers not supported by pCUDA
+// XFAIL: pcuda
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 42
+// RUN: %t 1 2 3; test $? -eq 42
+// RUN: %t 1 2 3 4 5; test $? -eq 42
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4,
+            (global_sym: %gpu.GlobalSymPtr (0, I32), const_sym: %gpu.ConstSymPtr (0, I32)), (), dummy: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let (m1, global) = %gpu.symptr2ptr (m1, global_sym);
+    let (m4, const) = %gpu.symptr2ptr (m4, const_sym);
+    let (m4, val) = %mem.load (m4, const);
+    let m1 = %mem.store (m1, global, val);
+    let m1 = %mem.store (m1, dummy, val);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, global_ptr, const_ptr) = %gpu.init (1, 1, m0, I32, I32);
+
+    let (m0, const_val) = %mem.slot (I32, 0) (m0, 0);
+    let m0 = %mem.store (m0, const_val, 42I32);
+    let (m0, m4) = %gpu.symbol_copy_to_device.block (m0, m4, const_val, const_ptr);
+
+    let (m0, m1) = %gpu.symbol_copy_to_device.block (m0, m1, ptr, global_ptr);
+
+    let (m1, dummy_global) = %gpu.alloc.block I32 m1;
+
+    ret m0 = %gpu.launch (
+        4, 4, %gpu.default_stream, ff, 2,
+        ((0, I32, %gpu.addr_space_global), (0, I32, %gpu.addr_space_const)),
+        (global_ptr, const_ptr), ()
+    ) kernel dummy_global $ m0;
+
+    let (m0, m1) = %gpu.symbol_copy_to_host.block (m0, m1, global_ptr, ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/const_mem_kernel.mim b/lit/pcuda/const_mem_kernel.mim
new file mode 100644
index 0000000000..78e5ef1c2c
--- /dev/null
+++ b/lit/pcuda/const_mem_kernel.mim
@@ -0,0 +1,47 @@
+// pCUDA backend: global/constant memory symbol pointers not supported by pCUDA
+// XFAIL: pcuda
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 42
+// RUN: %t 1 2 3; test $? -eq 42
+// RUN: %t 1 2 3 4 5; test $? -eq 42
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, const_sym: %gpu.ConstSymPtr (0, I32), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let (m4, const) = %gpu.symptr2ptr (m4, const_sym);
+    let (m4, val) = %mem.load (m4, const);
+    let m1 = %mem.store (m1, data, val);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, const_ptr) = %gpu.init (0, 1, m0, (), I32);
+
+    let (m0, const_val) = %mem.slot (I32, 0) (m0, 0);
+    let m0 = %mem.store (m0, const_val, 42I32);
+    let (m0, m4) = %gpu.symbol_copy_to_device.block (m0, m4, const_val, const_ptr);
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 1, (0, I32, 4), const_ptr, ()) kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/dyn_smem_kernel.mim b/lit/pcuda/dyn_smem_kernel.mim
new file mode 100644
index 0000000000..b9da086c06
--- /dev/null
+++ b/lit/pcuda/dyn_smem_kernel.mim
@@ -0,0 +1,42 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 42
+// RUN: %t 1 2 3; test $? -eq 42
+// RUN: %t 1 2 3 4 5; test $? -eq 42
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM, group_id: Idx 4, item_id: Idx 4,
+            (), smem: %gpu.SharedPtr «4; I32», data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let el = %mem.lea (smem, item_id);
+    let val = %core.bitcast I32 item_id;
+    let m3 = %mem.store (m3, el, val);
+    let m1 = %mem.store (m1, data, 42I32);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, tt, 0, (), (), «4; I32») kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/generic_kernel.mim b/lit/pcuda/generic_kernel.mim
new file mode 100644
index 0000000000..88d71a42a8
--- /dev/null
+++ b/lit/pcuda/generic_kernel.mim
@@ -0,0 +1,43 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 42
+// RUN: %t 1 2 3; test $? -eq 42
+// RUN: %t 1 2 3 4 5; test $? -eq 42
+
+plugin core;
+plugin gpu;
+
+con generic_kernel (n_groups: Nat, n_items: Nat)
+                   (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+                    group_id: Idx n_groups, item_id: Idx n_items, (), (), data: %gpu.GlobalPtr I32,
+                    return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let m1 = %mem.store (m1, data, 42I32);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    let k4 = generic_kernel (4, 4);
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) k4 d_ptr $ m0;
+    let k16 = generic_kernel (16, 16);
+    ret m0 = %gpu.launch (16, 16, %gpu.default_stream, ff, 0, (), (), ()) k16 d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/global_symbol_kernel.mim b/lit/pcuda/global_symbol_kernel.mim
new file mode 100644
index 0000000000..adb0c02410
--- /dev/null
+++ b/lit/pcuda/global_symbol_kernel.mim
@@ -0,0 +1,47 @@
+// pCUDA backend: global/constant memory symbol pointers not supported by pCUDA
+// XFAIL: pcuda
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 42
+// RUN: %t 1 2 3; test $? -eq 42
+// RUN: %t 1 2 3 4 5; test $? -eq 42
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, global_sym: %gpu.GlobalSymPtr (0, I32), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let (m1, glob) = %gpu.symptr2ptr (m1, global_sym);
+    let (m1, val) = %mem.load (m1, glob);
+    let m1 = %mem.store (m1, data, val);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, global_ptr, _) = %gpu.init (1, 0, m0, I32, ());
+
+    let (m0, global_val) = %mem.slot (I32, 0) (m0, 0);
+    let m0 = %mem.store (m0, global_val, 42I32);
+    let (m0, m1) = %gpu.symbol_copy_to_device.block (m0, m1, global_val, global_ptr);
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 1, (0, I32, 1), global_ptr, ()) kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/malloc_kernel.mim b/lit/pcuda/malloc_kernel.mim
new file mode 100644
index 0000000000..e0af7d5948
--- /dev/null
+++ b/lit/pcuda/malloc_kernel.mim
@@ -0,0 +1,42 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 42
+// RUN: %t 1 2 3; test $? -eq 42
+// RUN: %t 1 2 3 4 5; test $? -eq 42
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let (m1, ptr) = %mem.alloc (I32, 1) m1;
+    let m1 = %mem.store (m1, ptr, 42I32);
+    let (m1, val) = %mem.load (m1, ptr);
+    let m1 = %mem.store (m1, data, val);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/manual_streams.mim b/lit/pcuda/manual_streams.mim
new file mode 100644
index 0000000000..2ebba79ddf
--- /dev/null
+++ b/lit/pcuda/manual_streams.mim
@@ -0,0 +1,60 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 84
+// RUN: %t 1 2 3; test $? -eq 84
+// RUN: %t 1 2 3 4 5; test $? -eq 84
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let m1 = %mem.store (m1, data, 42I32);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr1) = %mem.alloc (I32, 0) m0;
+    let (m0, ptr2) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr1, argc);
+    let m0 = %mem.store (m0, ptr2, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m0, s1_ptr) = %mem.slot (%gpu.Stream, 0) (m0, 0);
+    let (m0, s2_ptr) = %mem.slot (%gpu.Stream, 0) (m0, 1);
+    let (m0, m1) = %gpu.stream_init (m0, m1, s1_ptr);
+    let (m0, m1) = %gpu.stream_init (m0, m1, s2_ptr);
+    let (m0, s1) = %mem.load (m0, s1_ptr);
+    let (m0, s2) = %mem.load (m0, s2_ptr);
+
+    let (m0, m1, d_ptr1) = %gpu.alloc_copy.async (m0, m1, ptr1, s1);
+    let (m0, m1, d_ptr2) = %gpu.alloc_copy.async (m0, m1, ptr2, s2);
+
+    ret m0 = %gpu.launch (4, 4, s1, ff, 0, (), (), ()) kernel d_ptr1 $ m0;
+    ret m0 = %gpu.launch (4, 4, s2, ff, 0, (), (), ()) kernel d_ptr2 $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.async (m0, m1, d_ptr1, ptr1, s1);
+    let (m0, m1) = %gpu.copy_to_host.async (m0, m1, d_ptr2, ptr2, s2);
+    let m1 = %gpu.free.async (m1, d_ptr1, s1);
+    let m1 = %gpu.free.async (m1, d_ptr2, s2);
+
+    let (m0, m1) = %gpu.stream_sync (m0, m1, s1);
+    let (m0, m1) = %gpu.stream_sync (m0, m1, s2);
+
+    let (m0, m1) = %gpu.stream_deinit (m0, m1, s1);
+    let (m0, m1) = %gpu.stream_deinit (m0, m1, s2);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val1) = %mem.load (m0, ptr1);
+    let (m0, val2) = %mem.load (m0, ptr2);
+    let val = %core.wrap.add 0 (val1, val2);
+    let m0 = %mem.free (m0, ptr1);
+    let m0 = %mem.free (m0, ptr2);
+
+    return (m0, val);
diff --git a/lit/pcuda/manual_sync_kernel.mim b/lit/pcuda/manual_sync_kernel.mim
new file mode 100644
index 0000000000..ae895635a7
--- /dev/null
+++ b/lit/pcuda/manual_sync_kernel.mim
@@ -0,0 +1,52 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 2
+// RUN: %t 1 2 3; test $? -eq 2
+// RUN: %t 1 2 3 4 5; test $? -eq 2
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let (m3, shared_ptr) = %mem.slot («4; I32», 3) (m3, 0);
+    let store_loc = %mem.lea (shared_ptr, item_id);
+    let val = %core.bitcast I32 item_id;
+    let m3 = %mem.store (m3, store_loc, val);
+    let (m1, m3) = %gpu.sync_work_items (m1, m3);
+    let load_idx = %core.wrap.sub %core.mode.nusw (3_4, item_id);
+    let load_loc = %mem.lea (shared_ptr, load_idx);
+    let (m3, load_val) = %mem.load (m3, load_loc);
+    let cond = %core.icmp.e (item_id, 1_4);
+    (return, store)#cond (m1, m3, m4, m5)
+    where
+        lam store (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM) =
+            let m1 = %mem.store (m1, data, load_val);
+            return (m1, m3, m4, m5);
+    end;
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/mem_arg_kernel.mim b/lit/pcuda/mem_arg_kernel.mim
new file mode 100644
index 0000000000..7b9d7ae5b5
--- /dev/null
+++ b/lit/pcuda/mem_arg_kernel.mim
@@ -0,0 +1,34 @@
+// RUN: rm -f %t.ll
+// RUN: %mim %s; test $? -ne 0
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), (m0: %mem.M 0, h_ptr: %mem.Ptr0 I32, d_ptr: %gpu.GlobalPtr I32),
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let m0 = %mem.store (m0, h_ptr, 42I32);
+    let m1 = %mem.store (m1, d_ptr, 24I32);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel (m0, ptr, d_ptr) $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/no_arg_kernel.mim b/lit/pcuda/no_arg_kernel.mim
new file mode 100644
index 0000000000..fb7f952f90
--- /dev/null
+++ b/lit/pcuda/no_arg_kernel.mim
@@ -0,0 +1,31 @@
+// RUN: rm -f %t.ll
+// RUN: %mim %s -O 0
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), (),
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m0, m1, d_ptr) = %gpu.alloc_copy.block (m0, m1, ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel () $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %gpu.free.block (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/no_gpu_deinit.mim b/lit/pcuda/no_gpu_deinit.mim
new file mode 100644
index 0000000000..eb671e6311
--- /dev/null
+++ b/lit/pcuda/no_gpu_deinit.mim
@@ -0,0 +1,31 @@
+// RUN: rm -f %t.ll
+// RUN: %mim %s; test $? -ne 0
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let m1 = %mem.store (m1, data, 42I32);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, %gpu.GlobalM, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, m1, val);
diff --git a/lit/pcuda/no_gpu_init.mim b/lit/pcuda/no_gpu_init.mim
new file mode 100644
index 0000000000..bdf0a3f3f0
--- /dev/null
+++ b/lit/pcuda/no_gpu_init.mim
@@ -0,0 +1,31 @@
+// RUN: rm -f %t.ll
+// RUN: %mim %s; test $? -ne 0
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let m1 = %mem.store (m1, data, 42I32);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, m1: %gpu.GlobalM, m4: %gpu.ConstM, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/scope_streams.mim b/lit/pcuda/scope_streams.mim
new file mode 100644
index 0000000000..c3dbe38f8a
--- /dev/null
+++ b/lit/pcuda/scope_streams.mim
@@ -0,0 +1,54 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 84
+// RUN: %t 1 2 3; test $? -eq 84
+// RUN: %t 1 2 3 4 5; test $? -eq 84
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let m1 = %mem.store (m1, data, 42I32);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr1) = %mem.alloc (I32, 0) m0;
+    let (m0, ptr2) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr1, argc);
+    let m0 = %mem.store (m0, ptr2, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    ret (m0, m1) = %gpu.with_streams (2, scope)
+    where
+        fun scope (m0: %mem.M 0, m1: %gpu.GlobalM, (s1 s2: %gpu.Stream)) : [%mem.M 0, %gpu.GlobalM] =
+            let (m0, m1, d_ptr1) = %gpu.alloc_copy.async (m0, m1, ptr1, s1);
+            let (m0, m1, d_ptr2) = %gpu.alloc_copy.async (m0, m1, ptr2, s2);
+
+            ret m0 = %gpu.launch (4, 4, s1, ff, 0, (), (), ()) kernel d_ptr1 $ m0;
+            ret m0 = %gpu.launch (4, 4, s2, ff, 0, (), (), ()) kernel d_ptr2 $ m0;
+
+            let (m0, m1) = %gpu.copy_to_host.async (m0, m1, d_ptr1, ptr1, s1);
+            let (m0, m1) = %gpu.copy_to_host.async (m0, m1, d_ptr2, ptr2, s2);
+            let m1 = %gpu.free.async (m1, d_ptr1, s1);
+            let m1 = %gpu.free.async (m1, d_ptr2, s2);
+
+            return (m0, m1);
+    end
+    $ (m0, m1);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val1) = %mem.load (m0, ptr1);
+    let (m0, val2) = %mem.load (m0, ptr2);
+    let val = %core.wrap.add 0 (val1, val2);
+    let m0 = %mem.free (m0, ptr1);
+    let m0 = %mem.free (m0, ptr2);
+
+    return (m0, val);
diff --git a/lit/pcuda/scope_sync_kernel.mim b/lit/pcuda/scope_sync_kernel.mim
new file mode 100644
index 0000000000..423aed70d7
--- /dev/null
+++ b/lit/pcuda/scope_sync_kernel.mim
@@ -0,0 +1,68 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 2
+// RUN: %t 1 2 3; test $? -eq 2
+// RUN: %t 1 2 3 4 5; test $? -eq 2
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let (m1, m3) = %gpu.sync_work_items (m1, m3);
+    let (m3, shared_ptr) = %mem.slot («4; I32», 3) (m3, 0);
+    let (m1, local_data) = %mem.alloc (I32, 1) m1;
+    let (m1, m3) = %gpu.synced_scope (m1, m3, scope)
+    where
+        lam scope (m1: %gpu.GlobalM, m3: %gpu.SharedM) =
+            let store_loc = %mem.lea (shared_ptr, item_id);
+            let val = %core.bitcast I32 item_id;
+            let m3 = %mem.store (m3, store_loc, val);
+            (m1, m3);
+    end;
+    let (m1, m3) = %gpu.synced_scope (m1, m3, scope)
+    where
+        lam scope (m1: %gpu.GlobalM, m3: %gpu.SharedM) =
+            let load_idx = %core.wrap.sub %core.mode.nusw (3_4, item_id);
+            let load_loc = %mem.lea (shared_ptr, load_idx);
+            let (m3, load_val) = %mem.load (m3, load_loc);
+            let cond = %core.icmp.e (item_id, 1_4);
+            let m1 = %mem.store (m1, local_data, load_val);
+            (m1, m3);
+    end;
+    let (m1, val) = %mem.load (m1, local_data);
+    let m1 = %mem.free (m1, local_data);
+    let cond = %core.icmp.e (item_id, 1_4);
+    (return, store)#cond (m1, m3, m4, m5)
+    where
+        lam store (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM) =
+            let m1 = %mem.store (m1, data, val);
+            return (m1, m3, m4, m5);
+    end;
+
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/simple_kernel.mim b/lit/pcuda/simple_kernel.mim
new file mode 100644
index 0000000000..0577909d01
--- /dev/null
+++ b/lit/pcuda/simple_kernel.mim
@@ -0,0 +1,39 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 42
+// RUN: %t 1 2 3; test $? -eq 42
+// RUN: %t 1 2 3 4 5; test $? -eq 42
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), data: %gpu.GlobalPtr I32,
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let m1 = %mem.store (m1, data, 42I32);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr, d_ptr);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel d_ptr $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr, ptr);
+    let m1 = %mem.free (m1, d_ptr);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val) = %mem.load (m0, ptr);
+    let m0 = %mem.free (m0, ptr);
+
+    return (m0, val);
diff --git a/lit/pcuda/two_arg_kernel.mim b/lit/pcuda/two_arg_kernel.mim
new file mode 100644
index 0000000000..159245033c
--- /dev/null
+++ b/lit/pcuda/two_arg_kernel.mim
@@ -0,0 +1,49 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+// RUN: %t; test $? -eq 242
+// RUN: %t 1 2 3; test $? -eq 242
+// RUN: %t 1 2 3 4 5; test $? -eq 242
+
+plugin core;
+plugin gpu;
+
+con kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+            group_id: Idx 4, item_id: Idx 4, (), (), (data1: %gpu.GlobalPtr I32, data2: %gpu.GlobalPtr I32),
+            return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let m1 = %mem.store (m1, data1, 42I32);
+    let m1 = %mem.store (m1, data2, 200I32);
+    return (m1, m3, m4, m5);
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0))
+              : [%mem.M 0, I32] =
+    let (m0, ptr1) = %mem.alloc (I32, 0) m0;
+    let (m0, ptr2) = %mem.alloc (I32, 0) m0;
+    let m0 = %mem.store (m0, ptr1, argc);
+    let m0 = %mem.store (m0, ptr2, argc);
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m1, d_ptr1) = %mem.alloc (I32, 1) m1;
+    let (m1, d_ptr2) = %mem.alloc (I32, 1) m1;
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr1, d_ptr1);
+    let (m0, m1) = %gpu.copy_to_device.block (m0, m1, ptr2, d_ptr2);
+
+    ret m0 = %gpu.launch (4, 4, %gpu.default_stream, ff, 0, (), (), ()) kernel (d_ptr1, d_ptr2) $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr1, ptr1);
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_ptr2, ptr2);
+    let m1 = %mem.free (m1, d_ptr1);
+    let m1 = %mem.free (m1, d_ptr2);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    let (m0, val1) = %mem.load (m0, ptr1);
+    let (m0, val2) = %mem.load (m0, ptr2);
+    let m0 = %mem.free (m0, ptr1);
+    let m0 = %mem.free (m0, ptr2);
+
+    let val = %core.wrap.add 0 (val1, val2);
+    return (m0, val);
diff --git a/lit/pcuda/unoptimized_gemm.mim b/lit/pcuda/unoptimized_gemm.mim
new file mode 100644
index 0000000000..14e2d970f8
--- /dev/null
+++ b/lit/pcuda/unoptimized_gemm.mim
@@ -0,0 +1,154 @@
+// REQUIRES: pcuda
+// RUN: rm -f %t.ll
+// RUN: %mim %s --output-ll %t.ll --device-target pcuda -o -
+// RUN: rm -f %t
+// RUN: clang++ %t.ll -o %t %acpp_link -Wno-override-module
+
+import gpu;
+import pcuda;
+import core;
+import libc;
+import math;
+import mem;
+import print;
+
+let n = 32768;
+let size = %core.nat.mul (n, n);
+let n_i32 = %core.bitcast I32 n;
+
+let numBlocks = 1048576;
+let numThreads = 1024;
+
+// GEMM of matrix dimensions:  nxn * nxn -> nxn
+con gemm_kernel (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM,
+                 group_id: Idx numBlocks, item_id: Idx numThreads, (), (),
+                 (out: %gpu.GlobalPtr «size; %math.F32»,
+                 in1: %gpu.GlobalPtr «size; %math.F32», in2: %gpu.GlobalPtr «size; %math.F32»),
+                 return: Cn [%gpu.GlobalM, %gpu.SharedM, %gpu.ConstM, %gpu.LocalM]) =
+    let cid = %core.bitcast I32 group_id;
+    let rid = %core.bitcast I32 item_id;
+
+    let id = %core.bitcast I32 (
+        %core.wrap.add 0 (%core.wrap.mul 0 (rid, n_i32), cid)
+    );
+
+    let row_oob = %core.icmp.sge (rid, n_i32);
+    let col_oob = %core.icmp.sge (cid, n_i32);
+    let tot_oob = %core.icmp.sge (id, %core.bitcast I32 (size));
+
+    let is_idle_thread = %core.bit2.or_ 0 (%core.bit2.or_ 0 (row_oob, col_oob), tot_oob);
+    (continue, return)#is_idle_thread (m1, m3, m4, m5)
+    where
+        con continue (m1: %gpu.GlobalM, m3: %gpu.SharedM, m4: %gpu.ConstM, m5: %gpu.LocalM) =
+            let id = %core.bitcast (Idx size) id;
+
+            let fp_zero = %math.conv.s2f %math.f32 0I32;
+            wrapped_loop (m1, 0I32, fp_zero)
+            where
+                con wrapped_loop (m1: %gpu.GlobalM, i: I32, acc: %math.F32) =
+                    let cond = %core.icmp.sl (i, n_i32);
+                    let next_i = %core.wrap.add 0 (i, 1I32);
+                    (write_and_exit, body)#cond m1
+                    where
+                        con body (m1: %gpu.GlobalM) =
+                            let in1_idx = %core.bitcast (Idx size) (
+                                %core.wrap.add 0 (%core.wrap.mul 0 (rid, n_i32), i)
+                            );
+                            let in2_idx = %core.bitcast (Idx size) (
+                                %core.wrap.add 0 (%core.wrap.mul 0 (i, n_i32), cid)
+                            );
+                            let in1_addr = %mem.lea (in1, in1_idx);
+                            let in2_addr = %mem.lea (in2, in2_idx);
+                            let (m1, in1_val) = %mem.load (m1, in1_addr);
+                            let (m1, in2_val) = %mem.load (m1, in2_addr);
+                            let new_acc = %pcuda.fmaf (in1_val, in2_val, acc);
+                            wrapped_loop (m1, next_i, new_acc);
+                        con write_and_exit (m1: %gpu.GlobalM) =
+                            let out_addr = %mem.lea (out, id);
+                            let m1 = %mem.store (m1, out_addr, acc);
+                            return (m1, m3, m4, m5);
+                    end;
+            end;
+    end;
+
+fun extern main (m0: %mem.M 0, argc: I32, argv: %mem.Ptr («⊤:Nat; %mem.Ptr («⊤:Nat; I8», 0)», 0)) =
+    let (m0, h_in1) = %mem.alloc («size; %math.F32», 0) m0;
+    let (m0, h_in2) = %mem.alloc («size; %math.F32», 0) m0;
+    let (m0, h_out) = %mem.alloc («size; %math.F32», 0) m0;
+
+    // generate ints in [0, cutoff), convert them to floats in [a, b), and fill h_in1 and h_in2 with these values
+    let seed = 42I32;
+    let m0 = %libc.srand (m0, seed);
+    let cutoff = 1000000I32;
+    let a = %math.minus 0 (%math.conv.s2f %math.f32 50I32);
+    let b = %math.conv.s2f %math.f32 50I32;
+    let div = %math.arith.div 0 (%math.conv.s2f %math.f32 cutoff, %math.arith.sub 0 (b, a));
+
+    fun loop_in1 (m0: %mem.M 0) : %mem.M 0 =
+        wrapped_loop (m0, 0I32, return)
+        where
+            con wrapped_loop (m0: %mem.M 0, i: I32, exit: Cn [%mem.M 0]) =
+                let cond = %core.icmp.sl (i, %core.bitcast I32 size);
+                let next_i = %core.wrap.add 0 (i, 1I32);
+                (exit, body)#cond m0
+                where
+                    con body (m0: %mem.M 0) =
+                        let idx = %core.bitcast (Idx size) i;
+                        let addr = %mem.lea (h_in1, idx);
+                        let (m0, value_i32) = %libc.rand m0;
+                        let (m0, value_i32) = %core.div.srem (m0, (value_i32, cutoff));
+                        let value = %math.conv.s2f %math.f32 value_i32;
+                        let value = %math.arith.div 0 (value, div);
+                        let value = %math.arith.add 0 (value, a);
+                        let m0 = %mem.store (m0, addr, value);
+                        wrapped_loop (m0, next_i, exit);
+                end;
+        end;
+    ret m0 = loop_in1 $ m0;
+
+    fun loop_in2 (m0: %mem.M 0) : %mem.M 0 =
+        wrapped_loop (m0, 0I32, return)
+        where
+            con wrapped_loop (m0: %mem.M 0, i: I32, exit: Cn [%mem.M 0]) =
+                let cond = %core.icmp.sl (i, %core.bitcast I32 size);
+                let next_i = %core.wrap.add 0 (i, 1I32);
+                (exit, body)#cond m0
+                where
+                    con body (m0: %mem.M 0) =
+                        let idx = %core.bitcast (Idx size) i;
+                        let addr = %mem.lea (h_in2, idx);
+                        let (m0, value_i32) = %libc.rand m0;
+                        let (m0, value_i32) = %core.div.srem (m0, (value_i32, cutoff));
+                        let value = %math.conv.s2f %math.f32 value_i32;
+                        let value = %math.arith.div 0 (value, div);
+                        let value = %math.arith.add 0 (value, a);
+                        let m0 = %mem.store (m0, addr, value);
+                        wrapped_loop (m0, next_i, exit);
+                end;
+        end;
+    ret m0 = loop_in2 $ m0;
+
+
+    let (m0, m1, m4, _, _) = %gpu.init (0, 0, m0, (), ());
+
+    let (m0, m1, d_in1) = %gpu.alloc_copy.block (m0, m1, h_in1);
+    let (m0, m1, d_in2) = %gpu.alloc_copy.block (m0, m1, h_in2);
+    let (m1, d_out) = %mem.alloc («size; %math.F32», %gpu.addr_space_global) m1;
+
+    ret m0 = %gpu.launch (numBlocks, numThreads, %gpu.default_stream, ff, 0, (), (), ())
+             gemm_kernel (d_out, d_in1, d_in2) $ m0;
+
+    let (m0, m1) = %gpu.copy_to_host.block (m0, m1, d_out, h_out);
+    let m1 = %gpu.free.block (m1, d_in1);
+    let m1 = %gpu.free.block (m1, d_in2);
+    let m1 = %gpu.free.block (m1, d_out);
+
+    let m0 = %gpu.deinit (m0, m1, m4);
+
+    ret m0 = %print.arr_f32 $ (m0, size, h_out, 10I32);
+
+    let m0 = %mem.free (m0, h_in1);
+    let m0 = %mem.free (m0, h_in2);
+    let m0 = %mem.free (m0, h_out);
+
+    return m0;
diff --git a/src/mim/cli/main.cpp b/src/mim/cli/main.cpp
index 9223e3feef..9b6585cf10 100644
--- a/src/mim/cli/main.cpp
+++ b/src/mim/cli/main.cpp
@@ -21,11 +21,12 @@ using namespace std::literals;
 
 int main(int argc, char** argv) {
     enum Backends { AST, Dot, H, LL, Md, Mim, Nest, Num_Backends };
-    enum DeviceTargets { None, NVPTX, Num_DeviceTargets };
+    enum DeviceTargets { None, NVPTX, PCUDA, Num_DeviceTargets };
 
     absl::btree_map<std::string, DeviceTargets> device_target_names = {
         { "none"s,  DeviceTargets::None},
         {"nvptx"s, DeviceTargets::NVPTX},
+        {"pcuda"s, DeviceTargets::PCUDA},
     };
 
     try {
@@ -71,7 +72,7 @@ int main(int argc, char** argv) {
             | lyra::opt(output[Md ],  "file"               )      ["--output-md"            ]("Emits the input formatted as Markdown.")
             | lyra::opt(output[Mim],  "file"               )["-o"]["--output-mim"           ]("Emits the Mim program again.")
             | lyra::opt(output[Nest], "file"               )      ["--output-nest"          ]("Emits program nesting tree as Dot.")
-            | lyra::opt(device_target_name, "target"       )      ["--device-target"        ]("Target for device code ('none' or 'nvptx'). Default: 'none'")
+            | lyra::opt(device_target_name, "target"       )      ["--device-target"        ]("Target for device code ('none', 'nvptx', or 'pcuda'). Default: 'none'")
               .choices([&device_target_names](std::string value) { return device_target_names.contains(value); })
             | lyra::opt(host_only                          )      ["--ll-host-only"         ]("Emit LLVM only for the host code (Default: compile both and embed device binary in host LLVM)")
             | lyra::opt(device_only                        )      ["--ll-device-only"       ]("Emit LLVM only for the device code (Default: compile both and embed device binary in host LLVM)")
@@ -164,6 +165,7 @@ int main(int argc, char** argv) {
             switch (device_target) {
                 case None: break;
                 case NVPTX: plugins.emplace_back("nvptx"s); break;
+                case PCUDA: plugins.emplace_back("pcuda"s); break;
                 case Num_DeviceTargets: fe::unreachable();
             }
 
@@ -225,6 +227,15 @@ int main(int argc, char** argv) {
                                 backend_name = "ll-host-nvptx-embed-dev";
                             plugin_name = "nvptx";
                             break;
+                        case PCUDA:
+                            if (host_only)
+                                backend_name = "ll-host-pcuda";
+                            else if (device_only)
+                                backend_name = "ll-dev-pcuda";
+                            else
+                                backend_name = "ll-host-pcuda-embed-dev";
+                            plugin_name = "pcuda";
+                            break;
                         case Num_DeviceTargets: fe::unreachable();
                     }
                     if (auto backend = driver.backend(backend_name))
diff --git a/src/mim/plug/CMakeLists.txt b/src/mim/plug/CMakeLists.txt
index 227f8605b8..02cb23e33b 100644
--- a/src/mim/plug/CMakeLists.txt
+++ b/src/mim/plug/CMakeLists.txt
@@ -12,6 +12,7 @@ set(MIM_PLUGINS
     matrix
     mem
     nvptx
+    pcuda
     opt
     ord
     print
diff --git a/src/mim/plug/nvptx/CMakeLists.txt b/src/mim/plug/nvptx/CMakeLists.txt
index cd3c71c0f4..0bb974e5d9 100644
--- a/src/mim/plug/nvptx/CMakeLists.txt
+++ b/src/mim/plug/nvptx/CMakeLists.txt
@@ -5,3 +5,9 @@ add_mim_plugin(nvptx
         ../core/be/ll.cpp # HACK: interim solution to link against LL emitter
     INSTALL
 )
+
+# Make the generated header visible to the plugin's sources.
+target_include_directories(mim_nvptx
+    PRIVATE
+        "${CMAKE_BINARY_DIR}/include"
+)
diff --git a/src/mim/plug/nvptx/be/ll_nvptx.cpp b/src/mim/plug/nvptx/be/ll_nvptx.cpp
index 64b98fa9ef..5f89c3f7e7 100644
--- a/src/mim/plug/nvptx/be/ll_nvptx.cpp
+++ b/src/mim/plug/nvptx/be/ll_nvptx.cpp
@@ -863,7 +863,7 @@ static std::optional<std::string> get_compute_capability() {
     // out should now have form "7.5" referencing the compute capability "sm_75"
 
     auto dot_pos = out.find('.');
-    assert(dot_pos < out.size());
+    if (dot_pos >= out.size()) return std::nullopt; // nvidia-smi unavailable or failed to parse
 
     for (size_t i = 0; i < out.size(); ++i)
         if (i != dot_pos && !std::isdigit(out[i])) return std::nullopt;
diff --git a/src/mim/plug/nvptx/nvptx.cpp b/src/mim/plug/nvptx/nvptx.cpp
index 287e47d3d9..05c5017d58 100644
--- a/src/mim/plug/nvptx/nvptx.cpp
+++ b/src/mim/plug/nvptx/nvptx.cpp
@@ -20,6 +20,7 @@ void reg_stages(Flags2Stages& stages) {
 
 extern "C" MIM_EXPORT Plugin mim_get_plugin() {
     return {"nvptx", nullptr, reg_stages, [](Backends& backends) {
+                // NVPTX (CUDA Driver API) backends
                 backends["ll-host-nvptx"]           = &ll::nvptx::emit_host;
                 backends["ll-dev-nvptx"]            = &ll::nvptx::emit_device;
                 backends["ll-host-nvptx-embed-dev"] = &ll::nvptx::emit_host_with_embedded_device;
diff --git a/src/mim/plug/opt/opt.mim b/src/mim/plug/opt/opt.mim
index 34f8136f1c..074610f12a 100644
--- a/src/mim/plug/opt/opt.mim
+++ b/src/mim/plug/opt/opt.mim
@@ -17,6 +17,7 @@ import direct;
 import gpu;
 import matrix;
 import nvptx;
+import pcuda;
 import refly;
 import regex;
 ///
@@ -75,6 +76,7 @@ lam extern _default_compile (): %compile.Phase =
                 %mem.alloc2malloc_repl,
                 cond_repl "gpu" %gpu.general_repls,
                 cond_repl "nvptx" %nvptx.general_repls,
+                cond_repl "pcuda" %pcuda.general_repls,
                 cond_repl "refly" %refly.remove_dbg_repl,
             )
         ),
diff --git a/src/mim/plug/pcuda/CMakeLists.txt b/src/mim/plug/pcuda/CMakeLists.txt
new file mode 100644
index 0000000000..fb451e1de0
--- /dev/null
+++ b/src/mim/plug/pcuda/CMakeLists.txt
@@ -0,0 +1,35 @@
+add_mim_plugin(pcuda
+    SOURCES
+        pcuda.cpp
+        be/ll_pcuda.cpp
+        be/hcf_adapter.cpp
+        ../core/be/ll.cpp # HACK: interim solution to link against LL emitter
+    INSTALL
+)
+
+# pCUDA backend: bake AdaptiveCpp paths + scripts/sscp_stubs.ll location into
+# a configured header so the orchestrator finds them regardless of cwd.
+if(AdaptiveCpp_FOUND AND MIM_ACPP_INSTALL_DIR)
+    set(MIM_ACPP_FOUND 1)
+    message("Using AdaptiveCpp from ${MIM_ACPP_INSTALL_DIR}")
+else()
+    set(MIM_ACPP_FOUND 0)
+    set(MIM_ACPP_INSTALL_DIR "")
+    set(MIM_ACPP_LIB_DIR "")
+endif()
+set(MIM_SSCP_STUBS_PATH "${CMAKE_BINARY_DIR}/lib/mim/sscp_stubs.ll")
+configure_file(
+    "${CMAKE_CURRENT_SOURCE_DIR}/sscp_stubs.ll"
+    "${MIM_SSCP_STUBS_PATH}"
+    COPYONLY)
+
+configure_file(
+    "${CMAKE_CURRENT_LIST_DIR}/be/pcuda_config.h.in"
+    "${CMAKE_BINARY_DIR}/include/mim/plug/pcuda/be/pcuda_config.h"
+    @ONLY
+)
+# Make the generated header visible to the plugin's sources.
+target_include_directories(mim_pcuda
+    PRIVATE
+        "${CMAKE_BINARY_DIR}/include"
+)
diff --git a/src/mim/plug/pcuda/be/hcf_adapter.cpp b/src/mim/plug/pcuda/be/hcf_adapter.cpp
new file mode 100644
index 0000000000..356399446f
--- /dev/null
+++ b/src/mim/plug/pcuda/be/hcf_adapter.cpp
@@ -0,0 +1,67 @@
+#include "mim/plug/pcuda/be/hcf_adapter.h"
+#include "mim/plug/pcuda/be/hcf_container.hpp"
+
+namespace mim::ll::pcuda {
+
+namespace {
+
+const char* param_type_name(HCFParamType t) {
+    switch (t) {
+        case HCFParamType::Integer:       return "integer";
+        case HCFParamType::FloatingPoint: return "floating-point";
+        case HCFParamType::Pointer:       return "pointer";
+        case HCFParamType::OtherByValue:  return "other-by-value";
+    }
+    return "other-by-value";
+}
+
+} // namespace
+
+std::string HCFBuilder::serialize() const {
+    hcf::hcf_container hcf;
+    auto* root = hcf.root_node();
+
+    root->set("object-id", std::to_string(object_id_));
+    root->set("generator", generator_);
+
+    // images / llvm-ir.global
+    auto* images   = root->add_subnode("images");
+    auto* llvm_img = images->add_subnode("llvm-ir.global");
+    llvm_img->set("variant", "global-module");
+    llvm_img->set("format", "llvm-ir");
+    llvm_img->set_as_list("exported-symbols", exported_);
+    llvm_img->set_as_list("imported-symbols", imported_);
+    hcf.attach_binary_content(llvm_img, device_bitcode_);
+
+    // kernels
+    auto* kernels_node = root->add_subnode("kernels");
+    for (const auto& k : kernels_) {
+        auto* kn = kernels_node->add_subnode(k.name);
+
+        kn->set_as_list("image-providers", {std::string{"llvm-ir.global"}});
+
+        auto* hsps = kn->add_subnode("host-side-parameter-sizes");
+        for (std::size_t i = 0; i < k.host_side_parameter_sizes.size(); ++i)
+            hsps->set(std::to_string(i), std::to_string(k.host_side_parameter_sizes[i]));
+
+        // Empty but present — the runtime parser expects these subnodes.
+        kn->add_subnode("compile-flags");
+        kn->add_subnode("compile-options");
+
+        auto* params = kn->add_subnode("parameters");
+        for (std::size_t i = 0; i < k.parameters.size(); ++i) {
+            const auto& p = k.parameters[i];
+            auto* pn = params->add_subnode(std::to_string(i));
+            pn->set("byte-offset", std::to_string(p.byte_offset));
+            pn->set("byte-size", std::to_string(p.byte_size));
+            pn->set("original-index", std::to_string(p.original_index));
+            pn->set("type", param_type_name(p.type));
+            auto* ann = pn->add_subnode("annotations");
+            for (const auto& a : p.annotations) ann->set(a, "1");
+        }
+    }
+
+    return hcf.serialize();
+}
+
+} // namespace mim::ll::pcuda
diff --git a/src/mim/plug/pcuda/be/ll_pcuda.cpp b/src/mim/plug/pcuda/be/ll_pcuda.cpp
new file mode 100644
index 0000000000..f1bdaca0db
--- /dev/null
+++ b/src/mim/plug/pcuda/be/ll_pcuda.cpp
@@ -0,0 +1,883 @@
+#include "mim/plug/pcuda/be/ll_pcuda.h"
+#include "mim/plug/pcuda/be/hcf_adapter.h"
+#include "mim/plug/pcuda/be/pcuda_config.h"
+
+#include <iomanip>
+#include <sstream>
+
+#include <mim/driver.h>
+#include <mim/util/sys.h>
+
+#include <mim/plug/core/core.h>
+#include <mim/plug/gpu/gpu.h>
+#include <mim/plug/mem/mem.h>
+#include <mim/plug/pcuda/pcuda.h>
+
+using namespace std::string_literals;
+
+namespace mim::ll::pcuda {
+
+namespace core  = mim::plug::core;
+namespace math  = mim::plug::math;
+namespace mem   = mim::plug::mem;
+namespace gpu   = mim::plug::gpu;
+namespace pcuda = mim::plug::pcuda;
+
+// ============================================================================
+// PCUDAHostEmitter Class Definition
+// ============================================================================
+
+class PCUDAHostEmitter : public mim::ll::Emitter {
+public:
+    using Super = mim::ll::Emitter;
+
+    PCUDAHostEmitter(World& world, std::ostream& ostream)
+        : Super(world, "ll_pcuda_host_emitter", ostream) {}
+
+    /// Provide the HCF bytes (output of HCFBuilder::serialize()) and the
+    /// matching object-id. When set (non-empty hcf_blob), start() emits the
+    /// stage-1 IR constants, AdaptiveCpp runtime externs, an llvm.global_ctors
+    /// entry that registers the HCF at program start, and per-kernel name +
+    /// storage globals used by the gpu::launch lowering.
+    void set_hcf_embed(std::string hcf_blob, std::uint64_t object_id) {
+        hcf_blob_  = std::move(hcf_blob);
+        object_id_ = object_id;
+    }
+
+    void start() final;
+    void find_kernels(const Def*);
+    void emit_epilogue(Lam*) final;
+
+    std::optional<std::string> isa_targetspecific_intrinsic(BB&, const Def*) final;
+
+protected:
+    std::string convert(const Def*) override;
+
+private:
+    void emit_hcf_embedding();
+
+    LamMap<int> kernel_ids_;
+    DefSet analyzed_;
+    std::string hcf_blob_;        // serialized HCF bytes; empty = no embedding
+    std::uint64_t object_id_ = 0; // matches @__acpp_local_sscp_hcf_object_id
+};
+
+// ============================================================================
+// PCUDADeviceEmitter Class Definition
+// ============================================================================
+
+class PCUDADeviceEmitter : public mim::ll::Emitter {
+public:
+    using Super = mim::ll::Emitter;
+
+    PCUDADeviceEmitter(World& world, std::ostream& ostream)
+        : Super(world, "ll_pcuda_device_emitter", ostream) {}
+
+    void start() final;
+    std::string prepare() override;
+    std::optional<std::string> isa_targetspecific_intrinsic(BB&, const Def*) final;
+
+protected:
+    std::string convert(const Def*) override;
+
+private:
+    LamSet kernels_;
+    absl::btree_map<std::string, int> symbols_;
+};
+
+// ============================================================================
+// PCUDAHostEmitter Implementation
+// ============================================================================
+
+void PCUDAHostEmitter::start() {
+    for (auto def : world().annexes())
+        find_kernels(def);
+    for (auto def : world().externals().muts())
+        find_kernels(def);
+
+    if (!hcf_blob_.empty()) emit_hcf_embedding();
+
+    Super::start();
+}
+
+namespace {
+
+/// Build an LLVM byte-string literal `c"\XX\XX..."` for arbitrary bytes.
+std::string llvm_byte_literal(std::string_view bytes) {
+    std::ostringstream s;
+    s << "c\"";
+    s << std::hex << std::setfill('0');
+    for (unsigned char c : bytes) s << "\\" << std::setw(2) << static_cast<int>(c);
+    s << "\"";
+    return s.str();
+}
+
+/// Build a NUL-terminated LLVM C-string literal for a kernel name.
+std::string llvm_cstring_literal(std::string_view s) {
+    std::ostringstream out;
+    out << "c\"";
+    out << std::hex << std::setfill('0');
+    for (unsigned char c : s) {
+        if (c >= 0x20 && c < 0x7F && c != '"' && c != '\\')
+            out << static_cast<char>(c);
+        else
+            out << "\\" << std::setw(2) << static_cast<int>(c);
+    }
+    out << "\\00\"";
+    return out.str();
+}
+
+} // namespace
+
+void PCUDAHostEmitter::emit_hcf_embedding() {
+    auto N = hcf_blob_.size();
+
+    // Stage-1 IR constants (resolved by AdaptiveCpp's static_hcf_registration template).
+    print(vars_decls_, "@__acpp_local_sscp_hcf_content = dso_local constant [{} x i8] {}\n", N,
+          llvm_byte_literal(hcf_blob_));
+    print(vars_decls_, "@__acpp_local_sscp_hcf_object_size = dso_local global i64 {}\n", N);
+    print(vars_decls_, "@__acpp_local_sscp_hcf_object_id = dso_local global i64 {}\n", object_id_);
+
+    // Per-kernel name + storage globals — referenced by gpu::launch lowering.
+    for (auto [lam, kid] : kernel_ids_) {
+        auto kname     = std::string{lam->sym().str()};
+        auto kname_len = kname.size() + 1; // +1 for NUL
+        print(vars_decls_, "@.kname_{} = private unnamed_addr constant [{} x i8] {}\n", kid, kname_len,
+              llvm_cstring_literal(kname));
+        print(vars_decls_, "@.kstorage_{} = internal global ptr null\n", kid);
+    }
+
+    // Runtime externs from libacpp-rt.so / libacpp-common.so.
+    declare("void @__acpp_register_hcf(ptr, i64)");
+    declare("void @__acpp_unregister_hcf(i64)");
+    declare("void @__pcudaPushCallConfiguration(i64, i32, i64, i32, i64, ptr)");
+    declare("i32 @__pcudaKernelCall(ptr, ptr, i64, ptr)");
+
+    // Global constructor that registers the HCF blob at program startup.
+    print(func_impls_,
+          "define internal void @__mimir_acpp_register_hcf_ctor() {{\n"
+          "  call void @__acpp_register_hcf(ptr @__acpp_local_sscp_hcf_content, i64 {})\n"
+          "  ret void\n"
+          "}}\n",
+          N);
+    print(vars_decls_,
+          "@llvm.global_ctors = appending global [1 x {{ i32, ptr, ptr }}] "
+          "[{{ i32, ptr, ptr }} {{ i32 65535, ptr @__mimir_acpp_register_hcf_ctor, ptr null }}]\n");
+}
+
+void PCUDAHostEmitter::find_kernels(const Def* def) {
+    if (auto [_, ins] = analyzed_.emplace(def); !ins) return;
+
+    for (auto d : def->deps())
+        find_kernels(d);
+
+    if (auto launch = Axm::isa<gpu::launch>(def)) {
+        auto kernel     = launch->decurry()->decurry()->arg();
+        auto kernel_lam = kernel->isa_mut<Lam>();
+        assert(kernel_lam && "Expect kernel passed to %gpu.launch to be a mutable lambda");
+        if (kernel_ids_.contains(kernel_lam)) return;
+        auto kid                = kernel_ids_.size();
+        kernel_ids_[kernel_lam] = kid;
+    }
+}
+
+std::string PCUDAHostEmitter::convert(const Def* type) {
+    if (auto ptr = Axm::isa<mem::Ptr>(type)) {
+        auto [_, addr_space] = ptr->args<2>();
+        auto lit             = Lit::isa(addr_space);
+        if (lit.value_or(0L) != 0) {
+            return "ptr";
+        }
+    } else if (auto symptr = Axm::isa<gpu::SymPtr>(type)) {
+        auto [_, T, a]      = symptr->args<3>();
+        auto ptr_equivalent = world().call<mem::Ptr>(Defs{T, a});
+        return convert(ptr_equivalent);
+    }
+    return Super::convert(type);
+}
+
+void PCUDAHostEmitter::emit_epilogue(Lam* lam) {
+    auto& bb = lam2bb_[lam];
+    auto app = lam->body()->as<App>();
+    if (auto ret = isa_targetspecific_intrinsic(bb, app)) {
+        assert(ret.has_value());
+        if (app->callee() == root()->ret_var())
+            assert(false && "Return not implemented in pCUDA backend");
+        else if (auto dispatch = Dispatch(app))
+            assert(false && "Dispatch not implemented in pCUDA backend");
+        else if (app->callee()->isa<Bot>())
+            assert(false && "Bot not implemented in pCUDA backend");
+        else if (auto _ = Lam::isa_mut_basicblock(app->callee()))
+            assert(false && "Ordinary Jump not implemented in pCUDA backend");
+        else if (Pi::isa_returning(app->callee_type()))
+            bb.tail("br label {}", ret.value());
+        else
+            assert(false && "Unexpected return case in pCUDA backend");
+    } else {
+        Super::emit_epilogue(lam);
+    }
+}
+
+// MimIR emits LLVM IR directly, so it must reference the *runtime symbols*
+// exported by libacpp-rt.so, not the header-only template wrappers like
+// `pcudaMalloc`/`pcudaMallocHost` (those are defined in
+// AdaptiveCpp/include/hipSYCL/pcuda/pcuda_runtime.hpp and forward to the
+// `pcudaAllocate*` runtime symbols below).
+constexpr auto PCUDA_MALLOC          = "pcudaAllocateDevice";
+constexpr auto PCUDA_FREE            = "pcudaFree";
+constexpr auto PCUDA_MALLOC_HOST     = "pcudaAllocateHost";
+constexpr auto PCUDA_FREE_HOST       = "pcudaFreeHost";
+constexpr auto PCUDA_MEMCPY          = "pcudaMemcpy";
+constexpr auto PCUDA_STREAM_CREATE   = "pcudaStreamCreate";
+constexpr auto PCUDA_STREAM_DESTROY  = "pcudaStreamDestroy";
+constexpr auto PCUDA_STREAM_SYNC     = "pcudaStreamSynchronize";
+constexpr auto PCUDA_GET_DEVICE_PROP = "pcudaGetDeviceProperties";
+constexpr auto PCUDA_SET_DEVICE      = "pcudaSetDevice";
+constexpr auto PCUDA_GET_DEVICE      = "pcudaGetDevice";
+constexpr auto PCUDA_GET_DEVICE_COUNT = "pcudaGetDeviceCount";
+
+std::optional<std::string> PCUDAHostEmitter::isa_targetspecific_intrinsic(BB& bb, const Def* def) {
+    auto name = id(def);
+
+    if (auto default_stream = Axm::isa<gpu::default_stream>(def)) {
+        return "null";
+    } else if (auto init = Axm::isa<gpu::init>(def)) {
+        declare("i32 @{}(ptr)", PCUDA_GET_DEVICE_COUNT);
+        auto dev_count_ptr = bb.assign(name + "_devcount_ptr", "alloca i32");
+        auto devcount_res = bb.assign(name + "_devcount_res",
+                                      "call i32 @{}(ptr {})",
+                                      PCUDA_GET_DEVICE_COUNT, dev_count_ptr);
+
+        declare("i32 @{}(i32)", PCUDA_SET_DEVICE);
+        auto setdev_res = bb.assign(name + "_setdev_res",
+                                    "call i32 @{}(i32 0)",
+                                    PCUDA_SET_DEVICE);
+
+        auto [n, m, mem, global_syms_def, const_syms_def] = init->args<5>();
+        return emit_unsafe(mem);
+    } else if (auto deinit = Axm::isa<gpu::deinit>(def)) {
+        emit_unsafe(deinit->arg(0));
+        return emit_unsafe(deinit->arg(1));
+    } else if (auto stream_init = Axm::isa<gpu::stream_init>(def)) {
+        declare("i32 @{}(ptr)", PCUDA_STREAM_CREATE);
+        emit_unsafe(stream_init->arg(0));
+        emit_unsafe(stream_init->arg(1));
+        auto stream_ptr = emit(stream_init->arg(2));
+        auto res = bb.assign(name, "call i32 @{}(ptr {})",
+                            PCUDA_STREAM_CREATE, stream_ptr);
+        return res;
+    } else if (auto stream_deinit = Axm::isa<gpu::stream_deinit>(def)) {
+        // pcudaStream_t = pcuda::stream* — runtime ABI is a ptr, not i32.
+        declare("i32 @{}(ptr)", PCUDA_STREAM_DESTROY);
+        emit_unsafe(stream_deinit->arg(0));
+        emit_unsafe(stream_deinit->arg(1));
+        auto stream = emit(stream_deinit->arg(2));
+        auto res = bb.assign(name, "call i32 @{}(ptr {})",
+                            PCUDA_STREAM_DESTROY, stream);
+        return res;
+    } else if (auto stream_sync = Axm::isa<gpu::stream_sync>(def)) {
+        declare("i32 @{}(ptr)", PCUDA_STREAM_SYNC);
+        emit_unsafe(stream_sync->arg(0));
+        emit_unsafe(stream_sync->arg(1));
+        auto stream = emit(stream_sync->arg(2));
+        auto res = bb.assign(name, "call i32 @{}(ptr {})",
+                            PCUDA_STREAM_SYNC, stream);
+        return res;
+    } else if (auto alloc = Axm::isa<gpu::alloc>(def)) {
+        declare("i32 @{}(ptr, i64)", PCUDA_MALLOC);
+        emit_unsafe(alloc->arg(0));
+        auto alloc_t    = alloc->decurry()->arg();
+        World& w        = alloc_t->world();
+        auto type_size  = w.call(core::trait::size, alloc_t);
+        auto alloc_size = emit(type_size);
+        auto ptr_t = convert(Axm::as<mem::Ptr>(def->proj(1)->type()));
+        auto alloc_ptr = bb.assign(name + "ptr", "alloca {}", ptr_t);
+        auto alloc_res = bb.assign(name + "res", "call i32 @{}(ptr {}, i64 {})",
+                                   PCUDA_MALLOC, alloc_ptr, alloc_size);
+        return bb.assign(name, "load {}, ptr {}", ptr_t, alloc_ptr);
+    } else if (auto free = Axm::isa<gpu::free>(def)) {
+        declare("i32 @{}(ptr)", PCUDA_FREE);
+        emit_unsafe(free->arg(0));
+        auto ptr = emit(free->arg(1));
+        auto free_res = bb.assign(name + "res", "call i32 @{}(ptr {})",
+                                 PCUDA_FREE, ptr);
+        return free_res;
+    } else if (auto copy_to_device = Axm::isa<gpu::copy_to_device>(def)) {
+        declare("i32 @{}(ptr, ptr, i64, i32)", PCUDA_MEMCPY);
+        auto type      = copy_to_device->decurry()->arg();
+        World& w       = type->world();
+        auto type_size = w.call(core::trait::size, type);
+        emit_unsafe(copy_to_device->arg(0));
+        emit_unsafe(copy_to_device->arg(1));
+        auto host_ptr = emit(copy_to_device->arg(2));
+        auto dev_ptr  = emit(copy_to_device->arg(3));
+        auto size     = emit(w.lit_nat(Lit::as(type_size)));
+        auto copy_res = bb.assign(name + "res", "call i32 @{}(ptr {}, ptr {}, i64 {}, i32 1)",
+                                 PCUDA_MEMCPY, dev_ptr, host_ptr, size);
+        return copy_res;
+    } else if (auto copy_to_host = Axm::isa<gpu::copy_to_host>(def)) {
+        declare("i32 @{}(ptr, ptr, i64, i32)", PCUDA_MEMCPY);
+        auto [type]    = copy_to_host->decurry()->args<1>();
+        World& w       = type->world();
+        auto type_size = w.call(core::trait::size, type);
+        emit_unsafe(copy_to_host->arg(0));
+        emit_unsafe(copy_to_host->arg(1));
+        auto dev_ptr  = emit(copy_to_host->arg(2));
+        auto host_ptr = emit(copy_to_host->arg(3));
+        auto size     = emit(w.lit_nat(Lit::as(type_size)));
+        auto copy_res = bb.assign(name + "res", "call i32 @{}(ptr {}, ptr {}, i64 {}, i32 2)",
+                                 PCUDA_MEMCPY, host_ptr, dev_ptr, size);
+        return copy_res;
+    } else if (auto launch = Axm::isa<gpu::launch>(def)) {
+        auto [implicits, launch_config, kernel_def, arg_def, func_args] = launch->uncurry_args<5>();
+        auto [n_groups_def, n_items_def, stream_def, m, _, __, ___, MT] = launch_config->projs<8>();
+        auto [mem, ret_lam_def] = func_args->projs<2>();
+
+        Lam* lam = kernel_def->isa_mut<Lam>();
+        if (!lam) error("kernel is not a lambda {}", kernel_def);
+        if (!kernel_ids_.contains(lam)) error("unknown kernel {}", lam);
+
+        // Host-only debug mode (no HCF blob set): consume operands but skip the
+        // real launch ABI — the resulting LL is not runnable but is useful for
+        // inspection. The full ABI requires emit_host_with_embedded_device.
+        if (hcf_blob_.empty()) {
+            emit_unsafe(mem);
+            emit(n_groups_def);
+            emit(n_items_def);
+            emit(stream_def);
+            return emit(ret_lam_def);
+        }
+        auto kid = kernel_ids_[lam];
+
+        // dynamic shared memory bytes (literal at compile time per the launch_config)
+        std::uint64_t shared_mem_bytes = 0;
+        if (auto smem_count = Lit::as(m)) {
+            if (smem_count != 1)
+                error("You can only have one dynamic allocation of shared memory per kernel");
+            shared_mem_bytes = Lit::as(world().call(core::trait::size, MT));
+        }
+
+        emit_unsafe(mem);
+        auto n_groups = emit(n_groups_def);
+        auto n_items  = emit(n_items_def);
+        auto stream   = emit(stream_def);
+
+        // Pack each dim3 as the x86-64 SysV coercion clang produces for libacpp-rt:
+        //   dim3{x, y=1, z=1} -> (i64 xy = (1u64 << 32) | x, i32 z = 1)
+        constexpr std::uint64_t y_eq_one_high = 1ull << 32;
+        auto grid_x64 = bb.assign(name + "_grid_x64", "zext i32 {} to i64", n_groups);
+        auto grid_xy  = bb.assign(name + "_grid_xy", "or i64 {}, {}", grid_x64, y_eq_one_high);
+        auto blk_x64  = bb.assign(name + "_blk_x64", "zext i32 {} to i64", n_items);
+        auto blk_xy   = bb.assign(name + "_blk_xy", "or i64 {}, {}", blk_x64, y_eq_one_high);
+
+        print(bb.body().emplace_back(),
+              "call void @__pcudaPushCallConfiguration(i64 {}, i32 1, i64 {}, i32 1, i64 {}, ptr {})",
+              grid_xy, blk_xy, shared_mem_bytes, stream);
+
+        // Flatten kernel args: one alloca + store per logical arg; pack pointers into [N x ptr].
+        auto arg_arity = arg_def->num_projs();
+        auto args_arr  = bb.assign(name + "_args", "alloca [{} x ptr]", arg_arity);
+        for (std::size_t i = 0; i < arg_arity; ++i) {
+            const Def* pi    = (arg_arity > 1) ? arg_def->proj(i) : arg_def;
+            auto pi_val      = emit(pi);
+            auto pi_type     = convert(pi->type());
+            auto slot        = bb.assign(name + "_arg" + std::to_string(i) + "_slot",
+                                         "alloca {}", pi_type);
+            print(bb.body().emplace_back(), "store {} {}, ptr {}", pi_type, pi_val, slot);
+            auto gep = bb.assign(name + "_arg" + std::to_string(i) + "_gep",
+                                 "getelementptr [{} x ptr], ptr {}, i64 0, i64 {}",
+                                 arg_arity, args_arr, i);
+            print(bb.body().emplace_back(), "store ptr {}, ptr {}", slot, gep);
+        }
+
+        // Load the object-id value (NOT its address) — __pcudaKernelCall's third
+        // arg is the runtime-known u64 used to key into hcf_cache.
+        auto obj_id = bb.assign(name + "_objid", "load i64, ptr @__acpp_local_sscp_hcf_object_id");
+        auto launch_res
+            = bb.assign(name + "_kcall",
+                        "call i32 @__pcudaKernelCall("
+                        "ptr @.kname_{}, ptr {}, "
+                        "i64 {}, "
+                        "ptr @.kstorage_{})",
+                        kid, args_arr, obj_id, kid);
+        (void)launch_res;
+
+        return emit(ret_lam_def);
+    }
+
+    return std::nullopt;
+}
+
+// ============================================================================
+// PCUDADeviceEmitter Implementation
+// ============================================================================
+
+std::string PCUDADeviceEmitter::convert(const Def* type) {
+    // Canonical SSCP device IR uses opaque `ptr` for all pointer types — no
+    // explicit address space. The JIT's llvm-to-amdgpu / llvm-to-spirv passes
+    // assign the right addrspace during backend flavoring. Leaving an explicit
+    // addrspace(N) in the stage-1 IR causes the AMDGPU lowering to produce a
+    // kernel that dereferences nil at runtime (GPU page fault).
+    if (auto ptr = Axm::isa<mem::Ptr>(type)) return "ptr";
+    if (auto symptr = Axm::isa<gpu::SymPtr>(type)) return "ptr";
+    return Super::convert(type);
+}
+
+void PCUDADeviceEmitter::start() {
+    for (auto kernel : world().externals().muts()) {
+        auto kernel_lam = kernel->isa_mut<Lam>();
+        assert(kernel_lam && "Expect kernel to be a mutable lambda");
+        kernels_.emplace(kernel_lam);
+    }
+    Super::start();
+
+    // SSCP kernel-discovery metadata: !hipsycl.sscp.annotations names each
+    // kernel function and its dimensionality. Without this, AdaptiveCpp's
+    // llvm-to-amdgpu pass cannot find the kernels in the embedded bitcode and
+    // the JIT silently skips them.
+    //
+    // MimIR's GPU model has a single group_id + item_id per launch (1D), so
+    // we always emit `hipsycl_kernel_dimension = 1`. Multi-dim launches would
+    // need this to come from launch_config instead.
+    std::ostringstream tmp;
+    int next_id = 0;
+    if (!kernels_.empty()) {
+        std::vector<int> md_ids;
+        for (auto k : kernels_) {
+            md_ids.push_back(next_id);
+            print(tmp, "!{} = !{{ptr {}, !\"hipsycl_kernel_dimension\", i32 1}}\n",
+                  next_id, id(k));
+            ++next_id;
+        }
+        print(tmp, "!hipsycl.sscp.annotations = !{{");
+        const char* sep = "";
+        for (auto i : md_ids) {
+            print(tmp, "{}!{}", sep, i);
+            sep = ", ";
+        }
+        print(tmp, "}}\n");
+    }
+
+    // AdaptiveCpp's PTX backend flavoring (LLVMToPtx::toBackendFlavor) appends
+    // the nvvm-reflect ftz / prec-div / prec-sqrt settings via
+    //   M.getModuleFlagsMetadata()->addOperand(...)
+    // and that accessor returns null when the module carries no
+    // !llvm.module.flags — so a device module without module flags makes the
+    // JIT dereference null and segfault. A clang-generated SSCP module always
+    // has module flags; MimIR's minimal device IR does not, so emit a benign
+    // one to give the JIT a node to append to.
+    int uwtable_id = next_id++;
+    print(tmp, "!{} = !{{i32 1, !\"uwtable\", i32 0}}\n", uwtable_id);
+    print(tmp, "!llvm.module.flags = !{{!{}}}\n", uwtable_id);
+
+    ostream() << tmp.str();
+}
+
+std::string PCUDADeviceEmitter::prepare() {
+    auto is_kern = kernels_.contains(root());
+    if (!is_kern) return Super::prepare();
+    auto kernel = root();
+
+    // Generate generic kernel attributes compatible with SSCP/multiple backends
+    // Instead of spir_kernel, use a generic function attribute
+    print(func_impls_, "define {} {}(", convert_ret_pi(kernel->type()->ret_pi()), id(kernel));
+
+    auto [m1, m3, m4, m5, group_id, item_id, symptrs, smem, arg, ret_lam] = kernel->vars<10>();
+
+    auto arg_name     = id(arg);
+    auto arg_type     = arg->type();
+    auto arg_type_str = convert(arg_type);
+    auto arg_arity    = arg->num_projs();
+    locals_[arg]      = arg_name;
+
+    auto& bb = lam2bb_[kernel];
+
+    // SSCP "free kernel" ABI (FreeKernelCall.cpp:142-234): each logical kernel
+    // arg gets its own LLVM parameter — the SSCP plugin's FreeKernelCallPass
+    // generates a launcher that for each arg either passes the value directly
+    // (pointer-to-ByVal-struct) or alloca's the value and passes &alloca, and
+    // the runtime copies bytes per HCF parameter offsets/sizes. This matches
+    // what MimIR already emits on the host side and what HCFBuilder records.
+    if (arg_arity > 1) {
+        // Tuple-typed arg: emit N flat LLVM params, synthesize the aggregate at
+        // function entry via insertvalue chain so the body's existing
+        // extractvalue uses continue to work unchanged.
+        std::vector<std::pair<std::string, std::string>> flat_params;
+        flat_params.reserve(arg_arity);
+        const char* sep = "";
+        for (size_t i = 0; i < arg_arity; ++i) {
+            auto pt = convert(arg->proj(i)->type());
+            auto pn = arg_name + ".f" + std::to_string(i);
+            flat_params.emplace_back(pt, pn);
+            print(func_impls_, "{}{} {}", sep, pt, pn);
+            sep = ", ";
+        }
+        print(func_impls_, ") {{\n");
+
+        std::string prev = "undef";
+        for (size_t i = 0; i < arg_arity; ++i) {
+            bool last      = (i + 1 == arg_arity);
+            auto step_name = last ? arg_name : (arg_name + ".iv" + std::to_string(i));
+            bb.assign(step_name, "insertvalue {} {}, {} {}, {}", arg_type_str, prev,
+                      flat_params[i].first, flat_params[i].second, i);
+            prev = step_name;
+        }
+    } else {
+        // Single-arg kernel: one LLVM param matching the arg type. The body
+        // uses arg_name directly (no extractvalue needed).
+        print(func_impls_, "{} {}) {{\n", arg_type_str, arg_name);
+    }
+
+    auto register_sreg_idx = [&](const Def* def, std::string_view sreg) {
+        auto name        = id(def);
+        auto type        = def->type();
+        auto type_name   = convert(type);
+        auto opt_idx_lit = Idx::isa_lit(type);
+        if (!opt_idx_lit) error("Type of '{}' must have known index type but has {}", def, type);
+        auto idx_lit = opt_idx_lit.value();
+        locals_[def] = name;
+
+        std::string sscp_builtin;
+        if (sreg == "ctaid.x")
+            sscp_builtin = "__acpp_sscp_get_group_id";
+        else if (sreg == "tid.x")
+            sscp_builtin = "__acpp_sscp_get_local_id";
+        else
+            error("pCUDA backend: unknown work-item register '{}'", sreg);
+
+        declare("i32 @{}(i32)", sscp_builtin);
+        if (type_name == "i0") {
+            locals_[def] = "0";
+        } else if (type_name == "i32") {
+            bb.assign(name, "call i32 @{}(i32 0)", sscp_builtin);
+        } else if (idx_lit < (1u << 31)) {
+            auto i32 = bb.assign(name + "i32", "call i32 @{}(i32 0)", sscp_builtin);
+            bb.assign(name, "trunc i32 {} to {}", i32, type_name);
+        } else {
+            error("Work-item index type too large, must fit into I32");
+        }
+    };
+    register_sreg_idx(group_id, "ctaid.x");
+    register_sreg_idx(item_id, "tid.x");
+
+    // Static/pre-declared shared memory: when the kernel's `smem` param is a
+    // %gpu.SharedPtr T (not the empty-tuple default), emit a module-level
+    // `addrspace(3) global T undef` and bind `locals_[smem]` to an
+    // addrspacecast constexpr that yields a generic ptr. The body then uses
+    // mem.lea/mem.store/etc. on the generic ptr unchanged.
+    //
+    // Canonical reference (acpp-built __shared__ var):
+    //   @__acpp_local_mem.k.0 = internal unnamed_addr addrspace(3) global [4 x i32] undef
+    //   ; uses: addrspacecast (ptr addrspace(3) @__acpp_local_mem.k.0 to ptr)
+    if (auto smem_ptr = Axm::isa<mem::Ptr>(smem->type())) {
+        auto [T, a] = smem_ptr->args<2>();
+        auto a_lit  = Lit::as(a);
+        auto gname  = "@" + smem->unique_name();
+        print(vars_decls_, "{} = internal addrspace({}) global {} undef\n", gname, a_lit, convert(T));
+        locals_[smem] = fmt("addrspacecast (ptr addrspace({}) {} to ptr)", a_lit, gname);
+    } else if (auto sigma = smem->type()->isa<Sigma>()) {
+        assert(sigma->num_ops() == 0 && "Expect empty sigma for shared memory variable");
+    }
+
+    // Symbol pointers (gpu.GlobalSymPtr / gpu.ConstSymPtr): declared in the
+    // device IR as externally-initialized globals with the appropriate
+    // addrspace. The host-side runtime is expected to write to them via
+    // gpu.symbol_copy_to_device — which on SSCP/pCUDA is not yet wired through
+    // AdaptiveCpp's runtime, so kernels that USE these will JIT-link with
+    // unresolved external symbols. The declaration alone keeps mim from
+    // hitting the Pi-conversion assertion when the body references symptrs.
+    auto bind_symptr = [&](const Def* var) {
+        auto sym_t = Axm::isa<gpu::SymPtr>(var->type());
+        if (!sym_t) return;
+        auto [id_, T, a] = sym_t->args<3>();
+        auto a_lit  = Lit::as(a);
+        auto id_lit = Lit::as(id_);
+        auto gname  = fmt("@__acpp_sscp_sym_a{}_{}", a_lit, id_lit);
+        if (symbols_.emplace(gname, a_lit).second) {
+            print(vars_decls_, "{} = dso_local addrspace({}) externally_initialized global {} undef\n",
+                  gname, a_lit, convert(T));
+        }
+        locals_[var] = fmt("addrspacecast (ptr addrspace({}) {} to ptr)", a_lit, gname);
+    };
+    if (auto sigma = symptrs->type()->isa<Sigma>()) {
+        if (sigma->num_ops() > 0) {
+            // Multiple symbols — would need to bind each projection. Not
+            // exercised by the current tests; deferred.
+            error("pCUDA backend: kernel symptrs with multiple entries not yet supported");
+        }
+    } else if (Axm::isa<gpu::SymPtr>(symptrs->type())) {
+        bind_symptr(symptrs);
+    }
+
+    return kernel->unique_name();
+}
+
+std::optional<std::string> PCUDADeviceEmitter::isa_targetspecific_intrinsic(BB& bb, const Def* def) {
+    auto name = id(def);
+
+    auto shared_as = Lit::as(world().annex<gpu::addr_space_shared>());
+
+    // Shared-memory `mem.slot` (addrspace=shared): emit a module-level
+    // addrspace(3) global and return an addrspacecast constexpr that yields a
+    // generic ptr. The rest of the body uses normal mem.lea/load/store on the
+    // generic ptr — LLVM verifies fine, and the AMDGPU JIT recognizes the
+    // addrspace(3) global as local memory during backend flavoring.
+    if (auto mslot = Axm::isa<mem::mslot>(def)) {
+        auto [T, a] = mslot->decurry()->args<2>();
+        auto a_lit  = Lit::as(a);
+        if (a_lit == shared_as) {
+            auto gname = "@" + def->unique_name();
+            emit_unsafe(mslot->arg(0));
+            print(vars_decls_, "{} = internal addrspace({}) global {} undef\n", gname, a_lit, convert(T));
+            return fmt("addrspacecast (ptr addrspace({}) {} to ptr)", a_lit, gname);
+        }
+        return std::nullopt;
+    }
+
+    // gpu::symptr2ptr — converts a gpu::SymPtr to a regular pointer. Since
+    // our convert() returns "ptr" for both types, this is a no-op at the IR
+    // level; just pass the inner pointer through.
+    if (auto symptr2ptr = Axm::isa<gpu::symptr2ptr>(def)) {
+        emit_unsafe(symptr2ptr->arg(0));
+        return emit(symptr2ptr->arg(1));
+    }
+
+    // Override base-class `mem::malloc` handling when address-space != 0.
+    // AMDGPU/SSCP has no usable device-side `malloc`/`free` symbol — the
+    // AMDGPU linker errors out with `undefined symbol: malloc` if we leave
+    // the call in. Lower per-thread heap-alloc to `alloca` (stack) instead.
+    // Semantically equivalent for the typical "scratch per-thread" pattern
+    // (e.g. scope_sync_kernel's `local_data`); the lifetime is bounded by
+    // the kernel invocation which is exactly what `alloca` gives us.
+    if (auto mlc = Axm::isa<mem::malloc>(def)) {
+        auto as = mlc->decurry()->arg(1);
+        if (Lit::as(as) != 0) {
+            emit_unsafe(mlc->arg(0));
+            // Element type comes from the result ptr's pointee.
+            auto pointee = Axm::as<mem::Ptr>(def->proj(1)->type())->arg(0);
+            return bb.assign(name, "alloca {}", convert(pointee));
+        }
+        return std::nullopt;
+    } else if (auto fr = Axm::isa<mem::free>(def)) {
+        auto as = fr->decurry()->arg(1);
+        if (Lit::as(as) != 0) {
+            // alloca is automatically released at function return; free is a no-op.
+            emit_unsafe(fr->arg(0));
+            emit_unsafe(fr->arg(1));
+            return name;
+        }
+        return std::nullopt;
+    }
+
+    if (auto sync_work_items = Axm::isa<gpu::sync_work_items>(def)) {
+        // Canonical SSCP barrier (AdaptiveCpp). Args: memory_scope, memory_order.
+        // scope=2 (work_group), order=0 (relaxed) per hipSYCL/sycl/libkernel/memory.hpp.
+        declare("void @__acpp_sscp_work_group_barrier(i32, i32)");
+        emit_unsafe(sync_work_items->arg(0));
+        emit_unsafe(sync_work_items->arg(1));
+        print(bb.body().emplace_back(), "call void @__acpp_sscp_work_group_barrier(i32 2, i32 0)");
+        return name;
+    } else if (auto warp_size = Axm::isa<pcuda::warp_size>(def)) {
+        // Use SSCP JIT reflection for warp size (backend-agnostic)
+        declare("i32 @__acpp_sscp_jit_reflect_warp_size()");
+        assert(name[0] == '%');
+        auto valid_name = name.substr(1);
+        bb.assign(valid_name, "call i32 @__acpp_sscp_jit_reflect_warp_size()");
+        return valid_name;
+    } else if (auto fmaf = Axm::isa<pcuda::fmaf>(def)) {
+        // Fused multiply-add → SSCP builtin, picked by float width. The JIT
+        // resolves __acpp_sscp_fma_f32/f64 against the kernel library for the
+        // active backend, so this stays backend-agnostic.
+        std::string_view builtin;
+        std::string_view ty;
+        switch (math::isa_f(fmaf->arg(0)->type()).value_or(0)) {
+            case 32: builtin = "__acpp_sscp_fma_f32"; ty = "float"; break;
+            case 64: builtin = "__acpp_sscp_fma_f64"; ty = "double"; break;
+            default: error("pCUDA backend: %pcuda.fmaf only supports f32 and f64");
+        }
+
+        auto x = emit(fmaf->arg(0));
+        auto y = emit(fmaf->arg(1));
+        auto z = emit(fmaf->arg(2));
+
+        declare("{} @{}({}, {}, {})", ty, builtin, ty, ty, ty);
+        return bb.assign(name, "call {} @{}({} {}, {} {}, {} {})", ty, builtin, ty, x, ty, y, ty, z);
+    }
+
+    return std::nullopt;
+}
+
+// ============================================================================
+// Backend emission functions
+// ============================================================================
+
+static auto get_setup_stage(World& world) {
+    auto flags         = Annex::base<gpu::setup4backend>();
+    auto stage_funcptr = world.driver().stage(flags);
+    auto stage         = (*stage_funcptr)(world);
+    auto phase         = stage->isa<RWPhase>();
+    if (!phase) error("Found unexpected gpu::setup4backend stage");
+    return std::make_pair(std::move(stage), phase);
+}
+
+void emit_host(World& world, std::ostream& ostream) {
+    auto [stage, setup_phase] = get_setup_stage(world);
+    setup_phase->run();
+
+    PCUDAHostEmitter emitter(setup_phase->old_world(), ostream);
+    emitter.run();
+}
+
+void emit_device(World& world, std::ostream& ostream) {
+    auto [stage, setup_phase] = get_setup_stage(world);
+    setup_phase->run();
+
+    PCUDADeviceEmitter emitter(setup_phase->new_world(), ostream);
+    emitter.run();
+}
+
+namespace {
+
+/// Classify a MimIR type into one of the HCF parameter categories and return
+/// its byte size. Pointers and SymPtrs are always 8 bytes; other types defer
+/// to MimIR's `core::trait::size`.
+std::pair<HCFParamType, std::size_t> classify_param(World& w, const Def* type) {
+    if (Axm::isa<mem::Ptr>(type) || Axm::isa<gpu::SymPtr>(type))
+        return {HCFParamType::Pointer, 8};
+
+    auto size_def    = w.call(core::trait::size, type);
+    std::size_t size = 0;
+    if (auto sz = Lit::isa(size_def)) size = static_cast<std::size_t>(*sz);
+
+    if (Idx::isa_lit(type)) return {HCFParamType::Integer, size};
+    if (math::isa_f(type))  return {HCFParamType::FloatingPoint, size};
+    return {HCFParamType::OtherByValue, size};
+}
+
+} // namespace
+
+void emit_host_with_embedded_device(World& world, std::ostream& ostream) {
+    static constexpr auto dev_ll_name   = "tmp_mimir_pcuda_dev.ll";
+    static constexpr auto dev_bc_name   = "tmp_mimir_pcuda_dev.bc";
+    static constexpr auto dev_link_name = "tmp_mimir_pcuda_dev_linked.bc";
+
+    // SSCP stubs path is baked at configure time from the source tree by CMake.
+    // If find_package(AdaptiveCpp) didn't run (MIM_WITH_ADAPTIVECPP=OFF), the
+    // path is the empty string — fall back to cwd-relative for back-compat.
+    std::string stub_path = config::SSCP_STUBS_PATH;
+    if (stub_path.empty()) stub_path = "scripts/sscp_stubs.ll";
+
+    auto [stage, setup_phase] = get_setup_stage(world);
+    setup_phase->run();
+    auto& new_w = setup_phase->new_world();
+
+    // 1. Emit device LLVM IR to a temp file.
+    {
+        std::ofstream dev_ofs(dev_ll_name);
+        if (!dev_ofs.is_open() || dev_ofs.fail())
+            error("Cannot open temp file '{}' for device LL", dev_ll_name);
+        PCUDADeviceEmitter device_emitter(new_w, dev_ofs);
+        device_emitter.run();
+    }
+
+    // 2a. Convert .ll -> .bc via llvm-as.
+    {
+        auto llvm_as = sys::find_cmd("llvm-as");
+        if (!std::filesystem::exists(llvm_as))
+            error("Could not find command: llvm-as ({})", llvm_as);
+        auto cmd = fmt("{} {} -o {}", llvm_as, dev_ll_name, dev_bc_name);
+        auto rc  = sys::system(cmd);
+        if (rc != 0) error("llvm-as exited with code {}", rc);
+    }
+
+    // 2b. Link in scripts/sscp_stubs.ll so the parameterized SSCP builtins
+    // (__acpp_sscp_get_group_id / get_local_id) are defined inside the embedded
+    // device bitcode. AdaptiveCpp's canonical kernel-library only ships the
+    // per-dimension variants (`_x` / `_y` / `_z`); the stubs dispatch our
+    // parameterized form to those. Without this link the SSCP JIT fails with
+    // "undefined symbol: __acpp_sscp_get_group_id" at runtime on every backend.
+    {
+        auto llvm_link = sys::find_cmd("llvm-link");
+        if (!std::filesystem::exists(llvm_link))
+            error("Could not find command: llvm-link ({})", llvm_link);
+        if (!std::filesystem::exists(stub_path))
+            error("SSCP stubs file not found at '{}' "
+                  "(run from project root or wire ACPP_INSTALL_DIR)", stub_path);
+        auto cmd = fmt("{} {} {} -o {}", llvm_link, dev_bc_name, stub_path, dev_link_name);
+        auto rc  = sys::system(cmd);
+        if (rc != 0) error("llvm-link exited with code {}", rc);
+    }
+
+    // 3. Slurp the linked .bc bytes (post-llvm-link with sscp_stubs).
+    std::string bc_bytes;
+    {
+        std::ifstream bc_ifs(dev_link_name, std::ios::binary);
+        if (!bc_ifs) error("Cannot open generated bitcode '{}'", dev_link_name);
+        std::ostringstream slurp;
+        slurp << bc_ifs.rdbuf();
+        bc_bytes = slurp.str();
+    }
+
+    // 4. Build the HCF blob (kernel metadata + bitcode attachment).
+    HCFBuilder hcf;
+    // TODO: derive object_id from a stable hash of the kernel set rather than a
+    // fixed literal so that two TUs linked into one binary don't collide.
+    constexpr std::uint64_t object_id = 0xACFFB1ECABCDEF01ull;
+    hcf.set_object_id(object_id);
+    hcf.set_device_bitcode(std::move(bc_bytes));
+
+    std::vector<std::string> exported_symbols;
+    for (auto def : new_w.externals().muts()) {
+        auto lam = def->isa_mut<Lam>();
+        if (!lam) continue;
+        auto kname = std::string{lam->sym().str()};
+        exported_symbols.push_back(kname);
+
+        HCFKernel k;
+        k.name = kname;
+
+        // The 9th kernel param (index 8) is the user-data arg; the preceding
+        // slots are mem effects (m1, m3, m4, m5), group_id, item_id, symptrs,
+        // smem. The final slot is the return continuation.
+        //
+        // For **free kernels** (the MimIR model: each logical arg becomes one
+        // LLVM parameter, not packed into a captures struct), each HCF
+        // parameter has byte-offset = 0 within its own host-side `args[i]`
+        // storage. Cf. acpp-built reference for `__global__ void k(int* a, int* b)`:
+        // the HCF emits two parameters both at byte-offset 0, distinguished by
+        // original-index 0 and 1. Cumulative offsets would mean SYCL/captures
+        // semantics (one struct with multiple fields).
+        auto arg   = lam->var(8);
+        auto arity = arg->num_projs();
+        for (std::size_t i = 0; i < arity; ++i) {
+            const Def* pi    = (arity > 1) ? arg->proj(i) : arg;
+            auto [cls, size] = classify_param(new_w, pi->type());
+            k.host_side_parameter_sizes.push_back(size);
+            k.parameters.push_back(HCFParam{0, size, i, cls, {}});
+        }
+        hcf.add_kernel(std::move(k));
+    }
+    hcf.set_exported_symbols(std::move(exported_symbols));
+
+    auto hcf_blob = hcf.serialize();
+
+    // 5. Run the host emitter, with the HCF blob to embed.
+    PCUDAHostEmitter host_emitter(setup_phase->old_world(), ostream);
+    host_emitter.set_hcf_embed(std::move(hcf_blob), object_id);
+    host_emitter.run();
+
+    // 6. Print the link recipe so users don't have to look it up. CMake baked
+    // the AdaptiveCpp install dir into pcuda_config.h at configure time.
+    if (config::ACPP_FOUND) {
+        std::cerr << "[pcuda] emitted host module with embedded device bitcode.\n"
+                     "[pcuda] link recipe:\n"
+                     "[pcuda]   clang++ -O2 <host.ll> -o <app> \\\n"
+                     "[pcuda]     -L"
+                  << config::ACPP_LIB_DIR
+                  << " -lacpp-rt -lacpp-common \\\n"
+                     "[pcuda]     -Wl,-rpath="
+                  << config::ACPP_LIB_DIR << " -pthread -ldl\n";
+    }
+}
+
+} // namespace mim::ll::pcuda
diff --git a/src/mim/plug/pcuda/be/pcuda_config.h.in b/src/mim/plug/pcuda/be/pcuda_config.h.in
new file mode 100644
index 0000000000..4ad94742e7
--- /dev/null
+++ b/src/mim/plug/pcuda/be/pcuda_config.h.in
@@ -0,0 +1,28 @@
+// Auto-generated by CMake from pcuda_config.h.in. Do not edit by hand.
+//
+// This header surfaces AdaptiveCpp / pCUDA paths discovered at configure time
+// so the runtime emitter can find scripts/sscp_stubs.ll regardless of the
+// caller's cwd and so the link recipe baked into diagnostics points at the
+// real AdaptiveCpp install.
+
+#pragma once
+
+namespace mim::ll::pcuda::config {
+
+// Set to 1 when find_package(AdaptiveCpp) succeeded at MimIR configure time.
+constexpr bool ACPP_FOUND = @MIM_ACPP_FOUND@;
+
+// AdaptiveCpp install root (the dir containing lib/libacpp-rt.so etc).
+// Empty string when ACPP_FOUND is 0.
+constexpr const char* ACPP_INSTALL_DIR = "@MIM_ACPP_INSTALL_DIR@";
+
+// AdaptiveCpp lib dir (where libacpp-rt.so + libacpp-common.so live).
+constexpr const char* ACPP_LIB_DIR = "@MIM_ACPP_LIB_DIR@";
+
+// Absolute path to scripts/sscp_stubs.ll (baked from MimIR's source tree).
+// emit_host_with_embedded_device shells `llvm-link` against this file to
+// resolve __acpp_sscp_get_group_id / __acpp_sscp_get_local_id (which the
+// canonical AdaptiveCpp kernel-library only ships in per-dim form).
+constexpr const char* SSCP_STUBS_PATH = "@MIM_SSCP_STUBS_PATH@";
+
+} // namespace mim::ll::pcuda::config
diff --git a/src/mim/plug/pcuda/pcuda.cpp b/src/mim/plug/pcuda/pcuda.cpp
new file mode 100644
index 0000000000..c5685ddbcc
--- /dev/null
+++ b/src/mim/plug/pcuda/pcuda.cpp
@@ -0,0 +1,28 @@
+#include "mim/plug/pcuda/pcuda.h"
+
+#include <mim/pass.h>
+#include <mim/plugin.h>
+
+#include <mim/plug/gpu/gpu.h>
+
+#include "mim/plug/pcuda/be/ll_pcuda.h"
+
+using namespace mim;
+using namespace mim::plug;
+
+void reg_stages(Flags2Stages& stages) {
+    MIM_REPL(stages, pcuda::stream_impl_repl, {
+        auto stream_flags = Annex::base<gpu::Stream>();
+        if (def->flags() == stream_flags) return world().annex<pcuda::Stream>();
+        return {};
+    });
+}
+
+extern "C" MIM_EXPORT Plugin mim_get_plugin() {
+    return {"pcuda", nullptr, reg_stages, [](Backends& backends) {
+                // pCUDA (AdaptiveCpp SSCP - generic multi-backend) backends
+                backends["ll-host-pcuda"]           = &ll::pcuda::emit_host;
+                backends["ll-dev-pcuda"]            = &ll::pcuda::emit_device;
+                backends["ll-host-pcuda-embed-dev"] = &ll::pcuda::emit_host_with_embedded_device;
+            }};
+}
diff --git a/src/mim/plug/pcuda/pcuda.mim b/src/mim/plug/pcuda/pcuda.mim
new file mode 100644
index 0000000000..e35e52803d
--- /dev/null
+++ b/src/mim/plug/pcuda/pcuda.mim
@@ -0,0 +1,56 @@
+/// # The pcuda Plugin {#pcuda}
+///
+/// @see mim::plug::pcuda
+///
+/// A minimal pCUDA plugin targeting AdaptiveCpp's SSCP (generic, multi-backend)
+/// runtime. Unlike @ref nvptx it does not lower streams to a backend-specific
+/// pointer type; the host emitter consumes the @ref gpu stream axioms directly
+/// and maps them onto the `pcudaStream*` runtime ABI.
+///
+/// [TOC]
+///
+/// ## Dependencies
+///
+plugin gpu;
+import compile;
+import mem;
+///
+/// ## Types
+///
+/// ### Streams
+///
+/// The pCUDA runtime ABI represents `pcudaStream_t` as a plain pointer, so the
+/// abstract @ref gpu stream is lowered to a byte pointer (mirrors @ref nvptx).
+let %pcuda.Stream = %mem.Ptr0 I8;
+///
+/// ## Operations
+///
+/// ### Special Registers
+///
+/// Retrieves the warp size on the device. With SSCP this is resolved via JIT
+/// reflection (`__acpp_sscp_jit_reflect_warp_size`) so it stays backend-agnostic.
+axm %pcuda.warp_size: I32;
+// TODO: only allow in device code
+///
+/// ### Arithmetic Functions
+///
+/// #### %%pcuda.fmaf
+///
+/// Computes `x*y+z`, i.e. fused-multiply-add. Lowered to AdaptiveCpp's
+/// backend-agnostic SSCP builtin `__acpp_sscp_fma_f32` / `__acpp_sscp_fma_f64`
+/// depending on the float width of `T`. Unlike @ref nvptx's `%nvptx.fmaf` there
+/// is no rounding-mode selector — SSCP rounds to nearest (the IEEE default).
+axm %pcuda.fmaf: {T: *} → [T, T, T] → T;
+///
+/// ## Stages
+///
+/// ### Repls
+///
+/// Lowers `%gpu.Stream` to `%pcuda.Stream` during the opt pipeline.
+axm %pcuda.stream_impl_repl: %compile.Repl;
+///
+/// ### Pipelines
+///
+let %pcuda.general_repls = %compile.repls (
+    %pcuda.stream_impl_repl,
+);
diff --git a/src/mim/plug/pcuda/sscp_stubs.ll b/src/mim/plug/pcuda/sscp_stubs.ll
new file mode 100644
index 0000000000..e5fec95da1
--- /dev/null
+++ b/src/mim/plug/pcuda/sscp_stubs.ll
@@ -0,0 +1,52 @@
+; SSCP stub dispatchers for MimIR pCUDA backend.
+;
+; MimIR's pCUDA emitter uses a parameterized form for work-item / group ID
+; queries (`__acpp_sscp_get_group_id(i32 dim)`), while AdaptiveCpp's canonical
+; SSCP ABI is per-dimension (`__acpp_sscp_get_group_id_x()` etc., returning
+; i64). This file bridges the two: link it against MimIR-emitted device IR
+; before invoking `acpp`.
+;
+; See: AdaptiveCpp/include/hipSYCL/sycl/libkernel/sscp/builtins/core.hpp
+
+declare i64 @__acpp_sscp_get_group_id_x()
+declare i64 @__acpp_sscp_get_group_id_y()
+declare i64 @__acpp_sscp_get_group_id_z()
+declare i64 @__acpp_sscp_get_local_id_x()
+declare i64 @__acpp_sscp_get_local_id_y()
+declare i64 @__acpp_sscp_get_local_id_z()
+
+define i32 @__acpp_sscp_get_group_id(i32 %dim) {
+entry:
+  switch i32 %dim, label %d0 [i32 1, label %d1
+                              i32 2, label %d2]
+d0:
+  %a = call i64 @__acpp_sscp_get_group_id_x()
+  %ai = trunc i64 %a to i32
+  ret i32 %ai
+d1:
+  %b = call i64 @__acpp_sscp_get_group_id_y()
+  %bi = trunc i64 %b to i32
+  ret i32 %bi
+d2:
+  %c = call i64 @__acpp_sscp_get_group_id_z()
+  %ci = trunc i64 %c to i32
+  ret i32 %ci
+}
+
+define i32 @__acpp_sscp_get_local_id(i32 %dim) {
+entry:
+  switch i32 %dim, label %d0 [i32 1, label %d1
+                              i32 2, label %d2]
+d0:
+  %a = call i64 @__acpp_sscp_get_local_id_x()
+  %ai = trunc i64 %a to i32
+  ret i32 %ai
+d1:
+  %b = call i64 @__acpp_sscp_get_local_id_y()
+  %bi = trunc i64 %b to i32
+  ret i32 %bi
+d2:
+  %c = call i64 @__acpp_sscp_get_local_id_z()
+  %ci = trunc i64 %c to i32
+  ret i32 %ci
+}
diff --git a/submodules/abseil-cpp b/submodules/abseil-cpp
new file mode 160000
index 0000000000..d38452e1ee
--- /dev/null
+++ b/submodules/abseil-cpp
@@ -0,0 +1 @@
+Subproject commit d38452e1ee03523a208362186fd42248ff2609f6
diff --git a/submodules/doxygen-awesome-css b/submodules/doxygen-awesome-css
new file mode 160000
index 0000000000..1f3620084f
--- /dev/null
+++ b/submodules/doxygen-awesome-css
@@ -0,0 +1 @@
+Subproject commit 1f3620084ff75734ed192101acf40e9dff01d848
diff --git a/submodules/fe b/submodules/fe
new file mode 160000
index 0000000000..400c60d4b8
--- /dev/null
+++ b/submodules/fe
@@ -0,0 +1 @@
+Subproject commit 400c60d4b8c1642f644380319262c240e77dff48
diff --git a/submodules/googletest b/submodules/googletest
new file mode 160000
index 0000000000..f8d7d77c06
--- /dev/null
+++ b/submodules/googletest
@@ -0,0 +1 @@
+Subproject commit f8d7d77c06936315286eb55f8de22cd23c188571
diff --git a/submodules/lyra b/submodules/lyra
new file mode 160000
index 0000000000..5a4ef1abaf
--- /dev/null
+++ b/submodules/lyra
@@ -0,0 +1 @@
+Subproject commit 5a4ef1abaf7139ef00204992a2627127ad394499
diff --git a/submodules/nanobind b/submodules/nanobind
new file mode 160000
index 0000000000..2a61ad2494
--- /dev/null
+++ b/submodules/nanobind
@@ -0,0 +1 @@
+Subproject commit 2a61ad2494d09fecb2e13322c1383342c299900d