diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index c0463368..b7ac431a 100644
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@@ -1,14 +1,9 @@
 #python code owners
-nvjitlink/            @rapidsai/nvjitlink-python-codeowners
+pynvjitlink/            @rapidsai/pynvjitlink-write
 
 #cmake code owners
-**/CMakeLists.txt  @rapidsai/nvjitlink-cmake-codeowners
-**/cmake/          @rapidsai/nvjitlink-cmake-codeowners
+**/CMakeLists.txt  @rapidsai/pynvjitlink-write
+**/cmake/          @rapidsai/pynvjitlink-write
 
 #build/ops code owners
-.github/           @rapidsai/ops-codeowners 
-ci/                @rapidsai/ops-codeowners
-conda/             @rapidsai/ops-codeowners
-**/Dockerfile      @rapidsai/ops-codeowners
-**/.dockerignore   @rapidsai/ops-codeowners
-docker/            @rapidsai/ops-codeowners
+.github/           @rapidsai/ops-codeowners
diff --git a/.github/ISSUE_TEMPLATE/bug_report.md b/.github/ISSUE_TEMPLATE/bug_report.md
index 85bcc1e0..d47e0704 100644
--- a/.github/ISSUE_TEMPLATE/bug_report.md
+++ b/.github/ISSUE_TEMPLATE/bug_report.md
@@ -1,6 +1,6 @@
 ---
 name: Bug report
-about: Create a bug report to help us improve nvjitlink
+about: Create a bug report to help us improve pynvjitlink
 title: "[BUG]"
 labels: "? - Needs Triage, bug"
 assignees: ''
@@ -18,9 +18,9 @@ A clear and concise description of what you expected to happen.
 
 **Environment details (please complete the following information):**
  - Environment location: [Bare-metal, Docker, Cloud(specify cloud provider)]
- - Method of nvjitlink install: [conda, Docker, or from source]
+ - Method of pynvjitlink install: [conda, Docker, or from source]
    - If method of install is [Docker], provide `docker pull` & `docker run` commands used
- 
+
 
 **Additional context**
 Add any other context about the problem here.
diff --git a/.github/ISSUE_TEMPLATE/feature-request.md b/.github/ISSUE_TEMPLATE/feature-request.md
index 3457b336..d8efc8f9 100644
--- a/.github/ISSUE_TEMPLATE/feature-request.md
+++ b/.github/ISSUE_TEMPLATE/feature-request.md
@@ -1,6 +1,6 @@
 ---
 name: Feature request
-about: Suggest an idea for nvjitlink
+about: Suggest an idea for pynvjitlink
 title: "[FEA]"
 labels: "? - Needs Triage, feature request"
 assignees: ''
@@ -8,7 +8,7 @@ assignees: ''
 ---
 
 **Is your feature request related to a problem? Please describe.**
-A clear and concise description of what the problem is. Ex. I wish I could use nvjitlink to do [...]
+A clear and concise description of what the problem is. Ex. I wish I could use pynvjitlink to do [...]
 
 **Describe the solution you'd like**
 A clear and concise description of what you want to happen.
diff --git a/.github/ISSUE_TEMPLATE/submit-question.md b/.github/ISSUE_TEMPLATE/submit-question.md
index c1890912..b3b86d1d 100644
--- a/.github/ISSUE_TEMPLATE/submit-question.md
+++ b/.github/ISSUE_TEMPLATE/submit-question.md
@@ -1,6 +1,6 @@
 ---
 name: Submit question
-about: Ask a general question about nvjitlink
+about: Ask a general question about pynvjitlink
 title: "[QST]"
 labels: "? - Needs Triage, question"
 assignees: ''
diff --git a/.github/PULL_REQUEST_TEMPLATE.md b/.github/PULL_REQUEST_TEMPLATE.md
index 4e930c42..03d36a01 100644
--- a/.github/PULL_REQUEST_TEMPLATE.md
+++ b/.github/PULL_REQUEST_TEMPLATE.md
@@ -1,6 +1,6 @@
 <!--
 
-Thank you for contributing to nvjitlink :)
+Thank you for contributing to pynvjitlink :)
 
 Here are some guidelines to help the review process go smoothly.
 
diff --git a/.github/actions/compute-matrix/action.yaml b/.github/actions/compute-matrix/action.yaml
new file mode 100644
index 00000000..4672fba3
--- /dev/null
+++ b/.github/actions/compute-matrix/action.yaml
@@ -0,0 +1,44 @@
+name: Compute matrix
+
+outputs:
+  BUILD_MATRIX:
+    value: ${{ steps.compute-matrix.outputs.BUILD_MATRIX }}
+  TEST_MATRIX:
+    value: ${{ steps.compute-matrix.outputs.TEST_MATRIX }}
+
+runs:
+  using: composite
+  steps:
+    - name: Compute matrix
+      id: compute-matrix
+      shell: bash
+      run: |
+        set -eo pipefail
+
+        export BUILD_MATRIX="
+        - { CUDA_VER: '12.0.1', ARCH: 'amd64', PY_VER: '3.9', LINUX_VER: 'centos7' }
+        - { CUDA_VER: '12.0.1', ARCH: 'amd64', PY_VER: '3.10', LINUX_VER: 'centos7' }
+        - { CUDA_VER: '12.0.1', ARCH: 'amd64', PY_VER: '3.11', LINUX_VER: 'centos7' }
+        - { CUDA_VER: '12.0.1', ARCH: 'arm64', PY_VER: '3.9', LINUX_VER: 'rockylinux8' }
+        - { CUDA_VER: '12.0.1', ARCH: 'arm64', PY_VER: '3.10', LINUX_VER: 'rockylinux8' }
+        - { CUDA_VER: '12.0.1', ARCH: 'arm64', PY_VER: '3.11', LINUX_VER: 'rockylinux8' }
+        "
+
+        export TEST_MATRIX="
+          - { CUDA_VER: '12.0.1', ARCH: 'amd64', PY_VER: '3.9', LINUX_VER: 'ubuntu18.04', gpu: 'v100', driver: 'latest' }
+          - { CUDA_VER: '12.0.1', ARCH: 'amd64', PY_VER: '3.10', LINUX_VER: 'ubuntu18.04', gpu: 'v100', driver: 'latest' }
+          - { CUDA_VER: '12.0.1', ARCH: 'amd64', PY_VER: '3.11', LINUX_VER: 'ubuntu18.04', gpu: 'v100', driver: 'latest' }
+          - { CUDA_VER: '12.0.1', ARCH: 'arm64', PY_VER: '3.9', LINUX_VER: 'ubuntu20.04', gpu: 'a100', driver: 'latest' }
+          - { CUDA_VER: '12.0.1', ARCH: 'arm64', PY_VER: '3.10', LINUX_VER: 'ubuntu20.04', gpu: 'a100', driver: 'latest' }
+          - { CUDA_VER: '12.0.1', ARCH: 'arm64', PY_VER: '3.11', LINUX_VER: 'ubuntu20.04', gpu: 'a100', driver: 'latest' }
+        "
+
+        echo "BUILD_MATRIX=$(
+          yq -n -o json 'env(BUILD_MATRIX)' | \
+          jq -c '.' \
+        )" | tee --append "${GITHUB_OUTPUT}"
+
+        echo "TEST_MATRIX=$(
+          yq -n -o json 'env(TEST_MATRIX)' | \
+          jq -c '.' \
+        )" | tee --append "${GITHUB_OUTPUT}"
diff --git a/.github/copy-pr-bot.yaml b/.github/copy-pr-bot.yaml
new file mode 100644
index 00000000..07ad9ea5
--- /dev/null
+++ b/.github/copy-pr-bot.yaml
@@ -0,0 +1,6 @@
+# Copyright (c) 2023, NVIDIA CORPORATION.
+
+# Configuration file for `copy-pr-bot` GitHub App
+# https://docs.gha-runners.nvidia.com/apps/copy-pr-bot/
+
+enabled: true
diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml
new file mode 100644
index 00000000..a8f251c2
--- /dev/null
+++ b/.github/workflows/build.yaml
@@ -0,0 +1,41 @@
+# Copyright (c) 2023, NVIDIA CORPORATION.
+name: build
+
+on:
+  push:
+    tags:
+      - 'v*'
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.ref }}-release
+  cancel-in-progress: true
+
+jobs:
+  compute-matrix:
+    runs-on: ubuntu-latest
+    outputs:
+      BUILD_MATRIX: ${{ steps.compute-matrix.outputs.BUILD_MATRIX }}
+    steps:
+      - uses: actions/checkout@v4
+      - name: Compute Build Matrix
+        id: compute-matrix
+        uses: ./.github/actions/compute-matrix
+  build-wheels:
+    needs:
+      - compute-matrix
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.02
+    with:
+      build_type: branch
+      script: "ci/build_wheel.sh"
+      matrix_filter: ${{ needs.compute-matrix.outputs.BUILD_MATRIX }}
+  publish-wheels:
+    needs:
+      - build-wheels
+    secrets: inherit
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-publish.yaml@branch-24.02
+    with:
+      build_type: ${{ inputs.build_type || 'branch' }}
+      branch: ${{ inputs.branch }}
+      sha: ${{ inputs.sha }}
+      date: ${{ inputs.date }}
+      package-name: pynvjitlink
diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml
new file mode 100644
index 00000000..0fe57264
--- /dev/null
+++ b/.github/workflows/pr.yaml
@@ -0,0 +1,53 @@
+# Copyright (c) 2023, NVIDIA CORPORATION.
+name: pr
+
+on:
+  push:
+    branches:
+      - "pull-request/[0-9]+"
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.ref }}-${{ github.event_name }}
+  cancel-in-progress: true
+
+jobs:
+  pr-builder:
+    needs:
+      - checks
+      - compute-matrix
+      - build-wheels
+      - test-wheels
+    secrets: inherit
+    uses: rapidsai/shared-workflows/.github/workflows/pr-builder.yaml@branch-24.02
+  checks:
+    secrets: inherit
+    uses: rapidsai/shared-workflows/.github/workflows/checks.yaml@branch-24.02
+    with:
+      enable_check_generated_files: false
+  compute-matrix:
+    runs-on: ubuntu-latest
+    outputs:
+      BUILD_MATRIX: ${{ steps.compute-matrix.outputs.BUILD_MATRIX }}
+      TEST_MATRIX: ${{ steps.compute-matrix.outputs.TEST_MATRIX }}
+    steps:
+      - uses: actions/checkout@v4
+      - name: Compute Build Matrix
+        id: compute-matrix
+        uses: ./.github/actions/compute-matrix
+  build-wheels:
+    needs:
+      - compute-matrix
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.02
+    with:
+      build_type: pull-request
+      script: "ci/build_wheel.sh"
+      matrix_filter: ${{ needs.compute-matrix.outputs.BUILD_MATRIX }}
+  test-wheels:
+    needs:
+      - build-wheels
+      - compute-matrix
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.02
+    with:
+      build_type: pull-request
+      script: "ci/test_wheel.sh"
+      matrix_filter: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
diff --git a/.gitignore b/.gitignore
index fb95f3b5..1ee3bad8 100644
--- a/.gitignore
+++ b/.gitignore
@@ -3,3 +3,7 @@ build
 *.so
 *.egg-info
 .*.swp
+
+# compiled test files
+pynvjitlink/tests/test_device_functions.*
+pynvjitlink/tests/undefined_extern.cubin
diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index c55fc173..0e557d65 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -1,9 +1,15 @@
 repos:
-- repo: https://github.com/PyCQA/flake8
-  rev: 3.7.8
-  hooks:
-  - id: flake8
-- repo: https://github.com/pre-commit/mirrors-clang-format
-  rev: v11.1.0
-  hooks:
-  - id: clang-format
+  - repo: https://github.com/pre-commit/pre-commit-hooks
+    rev: v4.5.0
+    hooks:
+      - id: trailing-whitespace
+      - id: end-of-file-fixer
+  - repo: https://github.com/pre-commit/mirrors-clang-format
+    rev: v16.0.6
+    hooks:
+      - id: clang-format
+  - repo: https://github.com/astral-sh/ruff-pre-commit
+    rev: v0.1.6
+    hooks:
+      - id: ruff
+      - id: ruff-format
diff --git a/.ycm_extra_conf.py b/.ycm_extra_conf.py
index c7483920..f14a0da5 100644
--- a/.ycm_extra_conf.py
+++ b/.ycm_extra_conf.py
@@ -3,28 +3,28 @@
 
 from pathlib import Path
 
-CONDA_PREFIX = os.environ['CONDA_PREFIX']
+CONDA_PREFIX = os.environ["CONDA_PREFIX"]
 NVJITLINK_DIR = Path(__file__).parent
 VER = sys.version_info
-PYTHON_INCLUDE_NAME = 'python3.10'  # Not ideal
+PYTHON_INCLUDE_NAME = "python3.10"  # Not ideal
 
-CONDA_INCLUDE_DIR = Path(CONDA_PREFIX, 'include')
+CONDA_INCLUDE_DIR = Path(CONDA_PREFIX, "include")
 PYTHON_INCLUDE_DIR = Path(CONDA_INCLUDE_DIR, PYTHON_INCLUDE_NAME)
-NVJITLINK_INCLUDE_DIR = Path(NVJITLINK_DIR, 'include')
-CUDA_INCLUDE_DIR = '/usr/local/cuda-12.0/include'
+NVJITLINK_INCLUDE_DIR = Path(NVJITLINK_DIR, "include")
+CUDA_INCLUDE_DIR = "/usr/local/cuda-12.0/include"
 
 flags = [
-    '--cuda-gpu-arch=sm_50',
-    f'-I{CONDA_INCLUDE_DIR}',
-    f'-I{CUDA_INCLUDE_DIR}',
-    f'-I{NVJITLINK_INCLUDE_DIR}',
-    f'-I{PYTHON_INCLUDE_DIR}',
+    "--cuda-gpu-arch=sm_50",
+    f"-I{CONDA_INCLUDE_DIR}",
+    f"-I{CUDA_INCLUDE_DIR}",
+    f"-I{NVJITLINK_INCLUDE_DIR}",
+    f"-I{PYTHON_INCLUDE_DIR}",
 ]
 
 
 def Settings(**kwargs):
-    return {'flags': flags}
+    return {"flags": flags}
 
 
-if __name__ == '__main__':
+if __name__ == "__main__":
     print(Settings())
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 54e30c07..3038ebd3 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -1,4 +1,4 @@
-# nvjitlink 0.0.0 (DD Mon YYYY)
+# pynvjitlink 0.0.0 (DD Mon YYYY)
 
 ## New Features
 
diff --git a/CMakeLists.txt b/CMakeLists.txt
index c5aef5a2..47cd7ad0 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,24 +1,22 @@
 cmake_minimum_required(VERSION 3.26.4 FATAL_ERROR)
 
 project(
-  nvjitlink
-  VERSION 0.1.0
-  LANGUAGES CXX
+  pynvjitlink
+  VERSION 0.1.6
+  LANGUAGES CXX CUDA
 )
 
 find_package(Python COMPONENTS Interpreter Development REQUIRED)
 
-Python_add_library(_nvjitlinklib MODULE nvjitlink/_nvjitlinklib.cpp WITH_SOABI)
+Python_add_library(_nvjitlinklib MODULE pynvjitlink/_nvjitlinklib.cpp WITH_SOABI)
 
 find_package(
-  CUDAToolkit 12.0 REQUIRED
+  CUDAToolkit 12.3 REQUIRED
 )
 target_link_libraries(_nvjitlinklib PRIVATE CUDA::nvJitLink_static CUDA::nvptxcompiler_static)
 
-set(opt_option -O0)
-if(WIN32)
-  set(opt_option -/Od)
-endif()
-target_compile_options(_nvjitlinklib PRIVATE -Werror -Wall ${opt_option})
+target_compile_options(_nvjitlinklib PRIVATE -Werror -Wall)
 
-install(TARGETS _nvjitlinklib LIBRARY DESTINATION nvjitlink)
+target_compile_features(_nvjitlinklib PRIVATE cxx_std_11)
+
+install(TARGETS _nvjitlinklib LIBRARY DESTINATION pynvjitlink)
diff --git a/README.md b/README.md
index cfaeba09..1890c51b 100644
--- a/README.md
+++ b/README.md
@@ -1,25 +1,26 @@
-# <div align="left"><img src="https://rapids.ai/assets/images/rapids_logo.png" width="90px"/>&nbsp;nvjitlink</div>
+# <div align="left"><img src="https://rapids.ai/assets/images/rapids_logo.png" width="90px"/>&nbsp;pynvjitlink</div>
 
-The [RAPIDS](https://rapids.ai) nvjitlink library provides a Python binding for the 
+The [RAPIDS](https://rapids.ai) pynvjitlink library provides a Python binding for the
 [nvJitLink library](https://docs.nvidia.com/cuda/nvJitLink/index.html). It is presently a
 work-in-progress.
 
-**NOTE:** For the latest stable [README.md](https://github.com/rapidsai/nvjitlink/blob/main/README.md) ensure you are on the `main` branch.
+**NOTE:** For the latest stable [README.md](https://github.com/rapidsai/pynvjitlink/blob/main/README.md) ensure you are on the `main` branch.
 
 ## Quick Start
 TODO
 
-## Install nvjitlink
+## Install pynvjitlink
+
 Install with either:
 
-```
-python setup.py develop
+```shell
+python -m pip install -e .
 ```
 
 or
 
-```
-python setup.py install
+```shell
+python -m pip install .
 ```
 
 ### Conda
@@ -30,4 +31,4 @@ TODO
 
 ## Contributing Guide
 
-Review the [CONTRIBUTING.md](https://github.com/rapidsai/nvjitlink/blob/main/CONTRIBUTING.md) file for information on how to contribute code and issues to the project.
+Review the [CONTRIBUTING.md](https://github.com/rapidsai/pynvjitlink/blob/main/CONTRIBUTING.md) file for information on how to contribute code and issues to the project.
diff --git a/ci/build_wheel.sh b/ci/build_wheel.sh
new file mode 100755
index 00000000..79d4a923
--- /dev/null
+++ b/ci/build_wheel.sh
@@ -0,0 +1,14 @@
+#!/bin/bash
+# Copyright (c) 2023-2024, NVIDIA CORPORATION
+
+set -euo pipefail
+
+rapids-logger "Install CUDA Toolkit"
+source "$(dirname "$0")/install_latest_cuda_toolkit.sh"
+
+rapids-logger "Build wheel"
+mkdir -p ./wheel-build
+pip wheel . --wheel-dir=./wheel-build -vvv
+
+rapids-logger "Upload Wheel"
+RAPIDS_PY_WHEEL_NAME="pynvjitlink-cu12" rapids-upload-wheels-to-s3 ./wheel-build
diff --git a/ci/check_style.sh b/ci/check_style.sh
new file mode 100755
index 00000000..84822c4b
--- /dev/null
+++ b/ci/check_style.sh
@@ -0,0 +1,9 @@
+#!/bin/bash
+# Copyright (c) 2023, NVIDIA CORPORATION
+
+set -euo pipefail
+
+pip install pre-commit
+
+# Run pre-commit checks
+pre-commit run --all-files --show-diff-on-failure
diff --git a/ci/install_latest_cuda_toolkit.sh b/ci/install_latest_cuda_toolkit.sh
new file mode 100755
index 00000000..ca4a7813
--- /dev/null
+++ b/ci/install_latest_cuda_toolkit.sh
@@ -0,0 +1,20 @@
+#!/bin/bash
+# Copyright (c) 2024, NVIDIA CORPORATION
+
+# Installs the latest CUDA Toolkit.
+# Supports CentOS 7 and Rocky Linux 8.
+
+yum update -y
+yum install -y epel-release
+
+OS_ID=$(. /etc/os-release; echo $ID)
+if [ "${OS_ID}" == "centos" ]; then
+    yum install -y nvidia-driver-latest-dkms
+elif [ "${OS_ID}" == "rocky" ]; then
+    yum install -y nvidia-driver
+else
+    echo "Error: OS not detected as CentOS or Rocky Linux. Exiting."
+    exit 1
+fi
+
+yum install -y cuda-toolkit-12-3
diff --git a/ci/test_wheel.sh b/ci/test_wheel.sh
new file mode 100755
index 00000000..ed329e16
--- /dev/null
+++ b/ci/test_wheel.sh
@@ -0,0 +1,18 @@
+#!/bin/bash
+# Copyright (c) 2023-2024, NVIDIA CORPORATION
+
+set -euo pipefail
+
+rapids-logger "Download Wheel"
+RAPIDS_PY_WHEEL_NAME="pynvjitlink-cu12" rapids-download-wheels-from-s3 ./wheel-build/
+
+rapids-logger "Install wheel"
+pip install --find-links ./wheel-build pynvjitlink-cu12
+
+rapids-logger "Build Tests"
+cd test_binary_generation && make
+
+rapids-logger "Run Tests"
+cd ..
+pip install pytest
+pytest pynvjitlink/tests
diff --git a/nvjitlink/tests/test_device_functions.a b/nvjitlink/tests/test_device_functions.a
deleted file mode 100644
index ced3797d..00000000
Binary files a/nvjitlink/tests/test_device_functions.a and /dev/null differ
diff --git a/nvjitlink/tests/test_device_functions.cubin b/nvjitlink/tests/test_device_functions.cubin
deleted file mode 100644
index 0b6eae94..00000000
Binary files a/nvjitlink/tests/test_device_functions.cubin and /dev/null differ
diff --git a/nvjitlink/tests/test_device_functions.fatbin b/nvjitlink/tests/test_device_functions.fatbin
deleted file mode 100644
index d61b9e9f..00000000
Binary files a/nvjitlink/tests/test_device_functions.fatbin and /dev/null differ
diff --git a/nvjitlink/tests/test_device_functions.ltoir b/nvjitlink/tests/test_device_functions.ltoir
deleted file mode 100644
index c4838663..00000000
Binary files a/nvjitlink/tests/test_device_functions.ltoir and /dev/null differ
diff --git a/nvjitlink/tests/test_device_functions.o b/nvjitlink/tests/test_device_functions.o
deleted file mode 100644
index ea6efe11..00000000
Binary files a/nvjitlink/tests/test_device_functions.o and /dev/null differ
diff --git a/nvjitlink/tests/test_device_functions.ptx b/nvjitlink/tests/test_device_functions.ptx
deleted file mode 100644
index 08d3e3f6..00000000
--- a/nvjitlink/tests/test_device_functions.ptx
+++ /dev/null
@@ -1,57 +0,0 @@
-//
-// Generated by NVIDIA NVVM Compiler
-//
-// Compiler Build ID: CL-31968024
-// Cuda compilation tools, release 12.0, V12.0.76
-// Based on NVVM 7.0.1
-//
-
-.version 8.0
-.target sm_75
-.address_size 64
-
-	// .globl	_Z13test_add_fp166__halfS_
-
-.visible .func  (.param .align 2 .b8 func_retval0[2]) _Z13test_add_fp166__halfS_(
-	.param .align 2 .b8 _Z13test_add_fp166__halfS__param_0[2],
-	.param .align 2 .b8 _Z13test_add_fp166__halfS__param_1[2]
-)
-{
-	.reg .b16 	%rs<4>;
-
-
-	ld.param.u16 	%rs2, [_Z13test_add_fp166__halfS__param_0];
-	ld.param.u16 	%rs3, [_Z13test_add_fp166__halfS__param_1];
-	// begin inline asm
-	{add.f16 %rs1,%rs2,%rs3;
-}
-	// end inline asm
-	st.param.b16 	[func_retval0+0], %rs1;
-	ret;
-
-}
-	// .globl	_Z13test_cmp_fp166__halfS_
-.visible .func  (.param .b32 func_retval0) _Z13test_cmp_fp166__halfS_(
-	.param .align 2 .b8 _Z13test_cmp_fp166__halfS__param_0[2],
-	.param .align 2 .b8 _Z13test_cmp_fp166__halfS__param_1[2]
-)
-{
-	.reg .pred 	%p<2>;
-	.reg .b16 	%rs<4>;
-	.reg .b32 	%r<2>;
-
-
-	ld.param.u16 	%rs2, [_Z13test_cmp_fp166__halfS__param_0];
-	ld.param.u16 	%rs3, [_Z13test_cmp_fp166__halfS__param_1];
-	// begin inline asm
-	{ .reg .pred __$temp3;
-  setp.eq.f16  __$temp3, %rs2, %rs3;
-  selp.u16 %rs1, 1, 0, __$temp3;}
-	// end inline asm
-	setp.ne.s16 	%p1, %rs1, 0;
-	selp.u32 	%r1, 1, 0, %p1;
-	st.param.b32 	[func_retval0+0], %r1;
-	ret;
-
-}
-
diff --git a/nvjitlink/tests/test_nvjitlink_api.py b/nvjitlink/tests/test_nvjitlink_api.py
deleted file mode 100644
index d012403a..00000000
--- a/nvjitlink/tests/test_nvjitlink_api.py
+++ /dev/null
@@ -1,161 +0,0 @@
-# Copyright (c) 2023, NVIDIA CORPORATION. All rights reserved.
-
-import pytest
-import sys
-import os
-
-from nvjitlink import NvJitLinker, NvJitLinkError
-
-
-@pytest.fixture
-def device_functions_cubin():
-    test_dir = os.path.dirname(os.path.abspath(__file__))
-    cubin_path = os.path.join(test_dir, 'test_device_functions.cubin')
-    with open(cubin_path, 'rb') as f:
-        return f.read()
-
-
-@pytest.fixture
-def device_functions_fatbin():
-    test_dir = os.path.dirname(os.path.abspath(__file__))
-    fatbin_path = os.path.join(test_dir, 'test_device_functions.fatbin')
-    with open(fatbin_path, 'rb') as f:
-        return f.read()
-
-
-@pytest.fixture
-def undefined_extern_cubin():
-    test_dir = os.path.dirname(os.path.abspath(__file__))
-    fatbin_path = os.path.join(test_dir, 'undefined_extern.cubin')
-    with open(fatbin_path, 'rb') as f:
-        return f.read()
-
-
-def test_create_no_arch_error():
-    # nvlink expects at least the architecture to be specified.
-    with pytest.raises(NvJitLinkError,
-                       match='NVJITLINK_ERROR_MISSING_ARCH error'):
-        NvJitLinker()
-
-
-def test_invalid_arch_error():
-    # sm_XX is not a valid architecture
-    with pytest.raises(NvJitLinkError,
-                       match='NVJITLINK_ERROR_UNRECOGNIZED_OPTION error'):
-        NvJitLinker('-arch=sm_XX')
-
-
-def test_invalid_option_type_error():
-    with pytest.raises(TypeError,
-                       match='Expecting only strings'):
-        NvJitLinker('-arch', 53)
-
-
-def test_create_and_destroy():
-    nvjitlinker = NvJitLinker('-arch=sm_53')
-    assert nvjitlinker.handle != 0
-
-
-def test_add_cubin(device_functions_cubin):
-    nvjitlinker = NvJitLinker('-arch=sm_75')
-    name = 'test_device_functions.cubin'
-    nvjitlinker.add_cubin(device_functions_cubin, name)
-
-
-def test_add_incompatible_cubin_arch_error(device_functions_cubin):
-    nvjitlinker = NvJitLinker('-arch=sm_70')
-    name = 'test_device_functions.cubin'
-    with pytest.raises(NvJitLinkError,
-                       match='NVJITLINK_ERROR_INVALID_INPUT error'):
-        nvjitlinker.add_cubin(device_functions_cubin, name)
-
-
-def test_add_fatbin_sm75(device_functions_fatbin):
-    nvjitlinker = NvJitLinker('-arch=sm_75')
-    name = 'test_device_functions.fatbin'
-    nvjitlinker.add_fatbin(device_functions_fatbin, name)
-
-
-def test_add_fatbin_sm70(device_functions_fatbin):
-    nvjitlinker = NvJitLinker('-arch=sm_70')
-    name = 'test_device_functions.fatbin'
-    nvjitlinker.add_fatbin(device_functions_fatbin, name)
-
-
-def test_add_incompatible_fatbin_arch_error(device_functions_fatbin):
-    nvjitlinker = NvJitLinker('-arch=sm_80')
-    name = 'test_device_functions.fatbin'
-    with pytest.raises(NvJitLinkError,
-                       match='NVJITLINK_ERROR_INVALID_INPUT error'):
-        nvjitlinker.add_fatbin(device_functions_fatbin, name)
-
-
-def test_add_cubin_with_fatbin_error(device_functions_fatbin):
-    nvjitlinker = NvJitLinker('-arch=sm_75')
-    name = 'test_device_functions.fatbin'
-    with pytest.raises(NvJitLinkError,
-                       match='NVJITLINK_ERROR_INVALID_INPUT error'):
-        nvjitlinker.add_cubin(device_functions_fatbin, name)
-
-
-def test_add_fatbin_with_cubin(device_functions_cubin):
-    # Adding a cubin with add_fatbin seems to work - this may be expected
-    # behaviour.
-    nvjitlinker = NvJitLinker('-arch=sm_75')
-    name = 'test_device_functions.cubin'
-    nvjitlinker.add_fatbin(device_functions_cubin, name)
-
-
-def test_duplicate_symbols_cubin_and_fatbin(device_functions_cubin,
-                                            device_functions_fatbin):
-    # This link errors because the cubin and the fatbin contain the same
-    # symbols.
-    nvjitlinker = NvJitLinker('-arch=sm_75')
-    name = 'test_device_functions.cubin'
-    nvjitlinker.add_cubin(device_functions_cubin, name)
-    name = 'test_device_functions.fatbin'
-    with pytest.raises(NvJitLinkError,
-                       match="NVJITLINK_ERROR_INVALID_INPUT error"):
-        nvjitlinker.add_fatbin(device_functions_fatbin, name)
-
-
-def test_get_linked_cubin_complete_empty_error():
-    nvjitlinker = NvJitLinker('-arch=sm_75')
-    cubin = nvjitlinker.get_linked_cubin()
-
-    # Linking nothing still gives us an empty ELF back
-    assert cubin[:4] == b'\x7fELF'
-
-
-def test_get_linked_cubin(device_functions_cubin):
-    nvjitlinker = NvJitLinker('-arch=sm_75')
-    name = 'test_device_functions.cubin'
-    nvjitlinker.add_cubin(device_functions_cubin, name)
-    cubin = nvjitlinker.get_linked_cubin()
-
-    # Just check we got something that looks like an ELF
-    assert cubin[:4] == b'\x7fELF'
-
-
-def test_get_error_log(undefined_extern_cubin):
-    nvjitlinker = NvJitLinker('-arch=sm_75')
-    name = 'undefined_extern.cubin'
-    nvjitlinker.add_cubin(undefined_extern_cubin, name)
-    with pytest.raises(NvJitLinkError):
-        nvjitlinker.get_linked_cubin()
-    error_log = nvjitlinker.error_log
-    assert "Undefined reference to '_Z5undefff'" in error_log
-
-
-def test_get_info_log(device_functions_cubin):
-    nvjitlinker = NvJitLinker('-arch=sm_75')
-    name = 'test_device_functions.cubin'
-    nvjitlinker.add_cubin(device_functions_cubin, name)
-    nvjitlinker.get_linked_cubin()
-    info_log = nvjitlinker.info_log
-    # Info log is empty
-    assert "" == info_log
-
-
-if __name__ == '__main__':
-    sys.exit(pytest.main())
diff --git a/nvjitlink/tests/undefined_extern.cubin b/nvjitlink/tests/undefined_extern.cubin
deleted file mode 100644
index 00c220e0..00000000
Binary files a/nvjitlink/tests/undefined_extern.cubin and /dev/null differ
diff --git a/pynvjitlink/VERSION b/pynvjitlink/VERSION
new file mode 100644
index 00000000..c946ee61
--- /dev/null
+++ b/pynvjitlink/VERSION
@@ -0,0 +1 @@
+0.1.6
diff --git a/nvjitlink/__init__.py b/pynvjitlink/__init__.py
similarity index 51%
rename from nvjitlink/__init__.py
rename to pynvjitlink/__init__.py
index 7b99f87d..597f9626 100644
--- a/nvjitlink/__init__.py
+++ b/pynvjitlink/__init__.py
@@ -1,6 +1,6 @@
 # Copyright (c) 2023, NVIDIA CORPORATION.
 
-from nvjitlink.api import NvJitLinker, NvJitLinkError
+from pynvjitlink.api import NvJitLinker, NvJitLinkError
 
 __all__ = (NvJitLinker, NvJitLinkError)
-__version__ = '0.0.0'
+__version__ = "0.0.0"
diff --git a/nvjitlink/_nvjitlinklib.cpp b/pynvjitlink/_nvjitlinklib.cpp
similarity index 99%
rename from nvjitlink/_nvjitlinklib.cpp
rename to pynvjitlink/_nvjitlinklib.cpp
index 83488166..421731d9 100644
--- a/nvjitlink/_nvjitlinklib.cpp
+++ b/pynvjitlink/_nvjitlinklib.cpp
@@ -323,7 +323,7 @@ static PyMethodDef ext_methods[] = {
     {nullptr}};
 
 static struct PyModuleDef moduledef = {
-    PyModuleDef_HEAD_INIT, "nvjitlink",
+    PyModuleDef_HEAD_INIT, "pynvjitlink",
     "Provides access to nvJitLink API methods", -1, ext_methods};
 
 PyMODINIT_FUNC PyInit__nvjitlinklib(void) {
diff --git a/nvjitlink/api.py b/pynvjitlink/api.py
similarity index 89%
rename from nvjitlink/api.py
rename to pynvjitlink/api.py
index 20cac5ed..34da29ee 100644
--- a/nvjitlink/api.py
+++ b/pynvjitlink/api.py
@@ -1,7 +1,7 @@
 # Copyright (c) 2023, NVIDIA CORPORATION.
 
 from enum import Enum
-from nvjitlink import _nvjitlinklib
+from pynvjitlink import _nvjitlinklib
 
 import weakref
 
@@ -43,7 +43,7 @@ def error_log(self):
 
     def add_data(self, input_type, data, name):
         if self._complete:
-            raise NvJitLinkError('Cannot add data to already-completeted link')
+            raise NvJitLinkError("Cannot add data to already-completeted link")
 
         try:
             _nvjitlinklib.add_data(self.handle, input_type.value, data, name)
@@ -53,23 +53,23 @@ def add_data(self, input_type, data, name):
             raise NvJitLinkError(f"{e}\n{self.error_log}")
 
     def add_cubin(self, cubin, name=None):
-        name = name or 'unnamed-cubin'
+        name = name or "unnamed-cubin"
         self.add_data(InputType.CUBIN, cubin, name)
 
     def add_ptx(self, ptx, name=None):
-        name = name or 'unnamed-ptx'
+        name = name or "unnamed-ptx"
         self.add_data(InputType.PTX, ptx, name)
 
     def add_ltoir(self, ltoir, name=None):
-        name = name or 'unnamed-ltoir'
+        name = name or "unnamed-ltoir"
         self.add_data(InputType.LTOIR, ltoir, name)
 
     def add_object(self, object_, name=None):
-        name = name or 'unnamed-object'
+        name = name or "unnamed-object"
         self.add_data(InputType.OBJECT, object_, name)
 
     def add_fatbin(self, fatbin, name=None):
-        name = name or 'unnamed-fatbin'
+        name = name or "unnamed-fatbin"
         self.add_data(InputType.FATBIN, fatbin, name)
 
     def add_library(self, library, name=None):
diff --git a/nvjitlink/patch.py b/pynvjitlink/patch.py
similarity index 51%
rename from nvjitlink/patch.py
rename to pynvjitlink/patch.py
index 784ef27b..e5f6cb3c 100644
--- a/nvjitlink/patch.py
+++ b/pynvjitlink/patch.py
@@ -1,5 +1,5 @@
 # Copyright (c) 2023, NVIDIA CORPORATION.
-from nvjitlink.api import NvJitLinker, NvJitLinkError
+from pynvjitlink.api import NvJitLinker, NvJitLinkError
 
 import os
 import pathlib
@@ -9,15 +9,19 @@
 
 required_numba_ver = (0, 57)
 
-mvc_docs_url = ("https://numba.readthedocs.io/en/stable/cuda/"
-                "minor_version_compatibility.html")
+mvc_docs_url = (
+    "https://numba.readthedocs.io/en/stable/cuda/" "minor_version_compatibility.html"
+)
 
 try:
     import numba
+
     ver = numba.version_info.short
     if ver < required_numba_ver:
-        _numba_error = (f"version {numba.__version__} is insufficient for "
-                        "patching - %s.%s is needed." % required_numba_ver)
+        _numba_error = (
+            f"version {numba.__version__} is insufficient for "
+            "patching - %s.%s is needed." % required_numba_ver
+        )
     else:
         _numba_version_ok = True
 except ImportError as ie:
@@ -25,8 +29,12 @@
 
 if _numba_version_ok:
     from numba.core import config
-    from numba.cuda.cudadrv.driver import (FILE_EXTENSION_MAP, NvrtcProgram,
-                                           Linker, LinkerError)
+    from numba.cuda.cudadrv.driver import FILE_EXTENSION_MAP, Linker, LinkerError
+
+    if ver < (0, 58):
+        from numba.cuda.cudadrv.driver import NvrtcProgram
+    else:
+        from numba.cuda.cudadrv.nvrtc import NvrtcProgram
 else:
     # Prevent the definition of PatchedLinker failing if we have no Numba
     # Linker - it won't be used anyway.
@@ -34,7 +42,14 @@
 
 
 class PatchedLinker(Linker):
-    def __init__(self, max_registers=None, lineinfo=False, cc=None):
+    def __init__(
+        self,
+        max_registers=None,
+        lineinfo=False,
+        cc=None,
+        lto=False,
+        additional_flags=None,
+    ):
         if cc is None:
             raise RuntimeError("PatchedLinker requires CC to be specified")
         if not any(isinstance(cc, t) for t in [list, tuple]):
@@ -42,13 +57,18 @@ def __init__(self, max_registers=None, lineinfo=False, cc=None):
 
         sm_ver = f"{cc[0] * 10 + cc[1]}"
         arch = f"-arch=sm_{sm_ver}"
-        opts = [arch]
+        options = [arch]
         if max_registers:
-            opts.append(f"-maxrregcount={max_registers}")
+            options.append(f"-maxrregcount={max_registers}")
         if lineinfo:
-            opts.append("-lineinfo")
+            options.append("-lineinfo")
+        if lto:
+            options.append("-lto")
+        if additional_flags is not None:
+            options.extend(additional_flags)
 
-        self._linker = NvJitLinker(*opts)
+        self._linker = NvJitLinker(*options)
+        self.options = options
 
     @property
     def info_log(self):
@@ -58,25 +78,36 @@ def info_log(self):
     def error_log(self):
         return self._linker.error_log
 
-    def add_ptx(self, ptx, name='<cudapy-ptx>'):
+    def add_ptx(self, ptx, name="<cudapy-ptx>"):
         self._linker.add_ptx(ptx, name)
 
+    def add_fatbin(self, fatbin, name="<external-fatbin>"):
+        self._linker.add_fatbin(fatbin, name)
+
+    def add_ltoir(self, ltoir, name="<external-ltoir>"):
+        self._linker.add_ltoir(ltoir, name)
+
+    def add_object(self, obj, name="<external-object>"):
+        self._linker.add_object(obj, name)
+
     def add_file(self, path, kind):
         try:
-            with open(path, 'rb') as f:
+            with open(path, "rb") as f:
                 data = f.read()
         except FileNotFoundError:
-            raise LinkerError(f'{path} not found')
+            raise LinkerError(f"{path} not found")
 
         name = pathlib.Path(path).name
-        if kind == FILE_EXTENSION_MAP['cubin']:
+        if kind == FILE_EXTENSION_MAP["cubin"]:
             fn = self._linker.add_cubin
-        elif kind == FILE_EXTENSION_MAP['fatbin']:
+        elif kind == FILE_EXTENSION_MAP["fatbin"]:
             fn = self._linker.add_fatbin
-        elif kind == FILE_EXTENSION_MAP['a']:
+        elif kind == FILE_EXTENSION_MAP["a"]:
             raise LinkerError("Don't know how to link archives")
-        elif kind == FILE_EXTENSION_MAP['ptx']:
+        elif kind == FILE_EXTENSION_MAP["ptx"]:
             return self.add_ptx(data, name)
+        elif kind == FILE_EXTENSION_MAP["o"]:
+            fn = self._linker.add_object
         else:
             raise LinkerError(f"Don't know how to link {kind}")
 
@@ -89,13 +120,13 @@ def add_cu(self, cu, name):
         program = NvrtcProgram(cu, name)
 
         if config.DUMP_ASSEMBLY:
-            print(("ASSEMBLY %s" % name).center(80, '-'))
+            print(("ASSEMBLY %s" % name).center(80, "-"))
             print(program.ptx.decode())
-            print('=' * 80)
+            print("=" * 80)
 
         # Link the program's PTX using the normal linker mechanism
         ptx_name = os.path.splitext(name)[0] + ".ptx"
-        self.add_ptx(program.ptx.rstrip(b'\x00'), ptx_name)
+        self.add_ptx(program.ptx.rstrip(b"\x00"), ptx_name)
 
     def complete(self):
         try:
@@ -106,8 +137,16 @@ def complete(self):
             raise LinkerError from e
 
 
-def new_patched_linker(max_registers=0, lineinfo=False, cc=None):
-    return PatchedLinker(max_registers, lineinfo, cc)
+def new_patched_linker(
+    max_registers=0, lineinfo=False, cc=None, lto=False, additional_flags=None
+):
+    return PatchedLinker(
+        max_registers=max_registers,
+        lineinfo=lineinfo,
+        cc=cc,
+        lto=lto,
+        additional_flags=additional_flags,
+    )
 
 
 def patch_numba_linker():
diff --git a/pynvjitlink/tests/test_patch.py b/pynvjitlink/tests/test_patch.py
new file mode 100644
index 00000000..14a57a5b
--- /dev/null
+++ b/pynvjitlink/tests/test_patch.py
@@ -0,0 +1,96 @@
+# Copyright (c) 2023, NVIDIA CORPORATION.
+
+import pytest
+import sys
+
+from pynvjitlink import patch, NvJitLinkError
+from pynvjitlink.patch import (
+    PatchedLinker,
+    patch_numba_linker,
+    new_patched_linker,
+    required_numba_ver,
+    _numba_version_ok,
+)
+from unittest.mock import patch as mock_patch
+
+
+def test_numba_patching_numba_not_ok():
+    with mock_patch.multiple(patch, _numba_version_ok=False, _numba_error="<error>"):
+        with pytest.raises(RuntimeError, match="Cannot patch Numba: <error>"):
+            patch_numba_linker()
+
+
+@pytest.mark.skipif(
+    not _numba_version_ok,
+    reason=f"Requires Numba == {required_numba_ver[0]}.{required_numba_ver[1]}",
+)
+def test_numba_patching():
+    # We import the linker here rather than at the top level because the import
+    # may fail if if Numba is not present or an unsupported version.
+    from numba.cuda.cudadrv.driver import Linker
+
+    patch_numba_linker()
+    assert Linker.new is new_patched_linker
+
+
+def test_create():
+    patched_linker = PatchedLinker(cc=(7, 5))
+    assert "-arch=sm_75" in patched_linker.options
+
+
+def test_create_no_cc_error():
+    # nvJitLink expects at least the architecture to be specified.
+    with pytest.raises(RuntimeError, match="PatchedLinker requires CC to be specified"):
+        PatchedLinker()
+
+
+def test_invalid_arch_error():
+    # CC 0.0 is not a valid compute capability
+    with pytest.raises(
+        NvJitLinkError, match="NVJITLINK_ERROR_UNRECOGNIZED_OPTION error"
+    ):
+        PatchedLinker(cc=(0, 0))
+
+
+def test_invalid_cc_type_error():
+    with pytest.raises(TypeError, match="`cc` must be a list or tuple of length 2"):
+        PatchedLinker(cc=0)
+
+
+@pytest.mark.parametrize("max_registers", (None, 32))
+@pytest.mark.parametrize("lineinfo", (False, True))
+@pytest.mark.parametrize("lto", (False, True))
+@pytest.mark.parametrize("additional_flags", (None, ("-g",), ("-g", "-time")))
+def test_ptx_compile_options(max_registers, lineinfo, lto, additional_flags):
+    patched_linker = PatchedLinker(
+        cc=(7, 5),
+        max_registers=max_registers,
+        lineinfo=lineinfo,
+        lto=lto,
+        additional_flags=additional_flags,
+    )
+
+    assert "-arch=sm_75" in patched_linker.options
+
+    if max_registers:
+        assert f"-maxrregcount={max_registers}" in patched_linker.options
+    else:
+        assert "-maxrregcount" not in patched_linker.options
+
+    if lineinfo:
+        assert "-lineinfo" in patched_linker.options
+    else:
+        assert "-lineinfo" not in patched_linker.options
+
+    if lto:
+        assert "-lto" in patched_linker.options
+    else:
+        assert "-lto" not in patched_linker.options
+
+    if additional_flags:
+        for flag in additional_flags:
+            assert flag in patched_linker.options
+
+
+if __name__ == "__main__":
+    sys.exit(pytest.main())
diff --git a/nvjitlink/tests/test_nvjitlink.py b/pynvjitlink/tests/test_pynvjitlink.py
similarity index 60%
rename from nvjitlink/tests/test_nvjitlink.py
rename to pynvjitlink/tests/test_pynvjitlink.py
index 39910b4a..e4c7e2fc 100644
--- a/nvjitlink/tests/test_nvjitlink.py
+++ b/pynvjitlink/tests/test_pynvjitlink.py
@@ -3,114 +3,111 @@
 import os
 import pytest
 
-from nvjitlink import _nvjitlinklib
-from nvjitlink.api import InputType
+from pynvjitlink import _nvjitlinklib
+from pynvjitlink.api import InputType
 
 
 def read_test_file(filename):
     test_dir = os.path.dirname(os.path.abspath(__file__))
     path = os.path.join(test_dir, filename)
-    with open(path, 'rb') as f:
+    with open(path, "rb") as f:
         return filename, f.read()
 
 
 @pytest.fixture
 def device_functions_cubin():
-    return read_test_file('test_device_functions.cubin')
+    return read_test_file("test_device_functions.cubin")
 
 
 @pytest.fixture
 def device_functions_fatbin():
-    return read_test_file('test_device_functions.fatbin')
+    return read_test_file("test_device_functions.fatbin")
 
 
 @pytest.fixture
 def device_functions_ltoir():
-    return read_test_file('test_device_functions.ltoir')
+    return read_test_file("test_device_functions.ltoir")
 
 
 @pytest.fixture
 def device_functions_object():
-    return read_test_file('test_device_functions.o')
+    return read_test_file("test_device_functions.o")
 
 
 @pytest.fixture
 def device_functions_archive():
-    return read_test_file('test_device_functions.a')
+    return read_test_file("test_device_functions.a")
 
 
 @pytest.fixture
 def device_functions_ptx():
-    return read_test_file('test_device_functions.ptx')
+    return read_test_file("test_device_functions.ptx")
 
 
 @pytest.fixture
 def undefined_extern_cubin():
-    return read_test_file('undefined_extern.cubin')
+    return read_test_file("undefined_extern.cubin")
 
 
 def test_create_no_arch_error():
     # nvjitlink expects at least the architecture to be specified.
-    with pytest.raises(RuntimeError,
-                       match='NVJITLINK_ERROR_MISSING_ARCH error'):
+    with pytest.raises(RuntimeError, match="NVJITLINK_ERROR_MISSING_ARCH error"):
         _nvjitlinklib.create()
 
 
-@pytest.mark.skip('Causes fatal error and exit(1)')
 def test_invalid_arch_error():
     # sm_XX is not a valid architecture
-    with pytest.raises(RuntimeError,
-                       match='NVJITLINK_ERROR_UNRECOGNIZED_OPTION error'):
-        _nvjitlinklib.create('-arch=sm_XX')
+    with pytest.raises(RuntimeError, match="NVJITLINK_ERROR_UNRECOGNIZED_OPTION error"):
+        _nvjitlinklib.create("-arch=sm_XX")
 
 
 def test_unrecognized_option_error():
-    with pytest.raises(RuntimeError,
-                       match='NVJITLINK_ERROR_UNRECOGNIZED_OPTION error'):
-        _nvjitlinklib.create('-fictitious_option')
+    with pytest.raises(RuntimeError, match="NVJITLINK_ERROR_UNRECOGNIZED_OPTION error"):
+        _nvjitlinklib.create("-fictitious_option")
 
 
 def test_invalid_option_type_error():
-    with pytest.raises(TypeError,
-                       match='Expecting only strings'):
-        _nvjitlinklib.create('-arch', 53)
+    with pytest.raises(TypeError, match="Expecting only strings"):
+        _nvjitlinklib.create("-arch", 53)
 
 
 def test_create_and_destroy():
-    handle = _nvjitlinklib.create('-arch=sm_53')
+    handle = _nvjitlinklib.create("-arch=sm_53")
     assert handle != 0
     _nvjitlinklib.destroy(handle)
 
 
 def test_complete_empty():
-    handle = _nvjitlinklib.create('-arch=sm_75')
+    handle = _nvjitlinklib.create("-arch=sm_75")
     _nvjitlinklib.complete(handle)
     _nvjitlinklib.destroy(handle)
 
 
-@pytest.mark.parametrize('input_file,input_type', [
-    ('device_functions_cubin', InputType.CUBIN),
-    ('device_functions_fatbin', InputType.FATBIN),
-    # XXX: LTOIR input type needs debugging - results in
-    # NVJITLINK_ERROR_INTERNAL.
-    pytest.param('device_functions_ltoir', InputType.LTOIR,
-                 marks=pytest.mark.xfail),
-    ('device_functions_ptx', InputType.PTX),
-    ('device_functions_object', InputType.OBJECT),
-    # XXX: Archive type needs debugging - results in a segfault.
-    pytest.param('device_functions_archive', InputType.LIBRARY,
-                 marks=pytest.mark.skip),
-])
+@pytest.mark.parametrize(
+    "input_file,input_type",
+    [
+        ("device_functions_cubin", InputType.CUBIN),
+        ("device_functions_fatbin", InputType.FATBIN),
+        # XXX: LTOIR input type needs debugging - results in
+        # NVJITLINK_ERROR_INTERNAL.
+        pytest.param(
+            "device_functions_ltoir", InputType.LTOIR, marks=pytest.mark.xfail
+        ),
+        ("device_functions_ptx", InputType.PTX),
+        ("device_functions_object", InputType.OBJECT),
+        ("device_functions_archive", InputType.LIBRARY),
+    ],
+)
 def test_add_file(input_file, input_type, request):
     filename, data = request.getfixturevalue(input_file)
 
-    handle = _nvjitlinklib.create('-arch=sm_75')
+    handle = _nvjitlinklib.create("-arch=sm_75")
     _nvjitlinklib.add_data(handle, input_type.value, data, filename)
     _nvjitlinklib.destroy(handle)
 
 
 def test_get_error_log(undefined_extern_cubin):
-    handle = _nvjitlinklib.create('-arch=sm_75')
+    handle = _nvjitlinklib.create("-arch=sm_75")
     filename, data = undefined_extern_cubin
     input_type = InputType.CUBIN.value
     _nvjitlinklib.add_data(handle, input_type, data, filename)
@@ -118,15 +115,14 @@ def test_get_error_log(undefined_extern_cubin):
         _nvjitlinklib.complete(handle)
     error_log = _nvjitlinklib.get_error_log(handle)
     _nvjitlinklib.destroy(handle)
-    # XXX: The error message in the log is strange. The actual expected error
-    # message appears on the terminal:
-    #     error   : Undefined reference to '_Z5undefff' in
-    #               'undefined_extern.cubin'
-    assert "ERROR 9: finish" in error_log
+    assert (
+        "Undefined reference to '_Z5undefff' "
+        "in 'undefined_extern.cubin'" in error_log
+    )
 
 
 def test_get_info_log(device_functions_cubin):
-    handle = _nvjitlinklib.create('-arch=sm_75')
+    handle = _nvjitlinklib.create("-arch=sm_75")
     filename, data = device_functions_cubin
     input_type = InputType.CUBIN.value
     _nvjitlinklib.add_data(handle, input_type, data, filename)
@@ -138,7 +134,7 @@ def test_get_info_log(device_functions_cubin):
 
 
 def test_get_linked_cubin(device_functions_cubin):
-    handle = _nvjitlinklib.create('-arch=sm_75')
+    handle = _nvjitlinklib.create("-arch=sm_75")
     filename, data = device_functions_cubin
     input_type = InputType.CUBIN.value
     _nvjitlinklib.add_data(handle, input_type, data, filename)
@@ -147,16 +143,15 @@ def test_get_linked_cubin(device_functions_cubin):
     _nvjitlinklib.destroy(handle)
 
     # Just check we got something that looks like an ELF
-    assert cubin[:4] == b'\x7fELF'
+    assert cubin[:4] == b"\x7fELF"
 
 
 def test_get_linked_cubin_link_not_complete_error(device_functions_cubin):
-    handle = _nvjitlinklib.create('-arch=sm_75')
+    handle = _nvjitlinklib.create("-arch=sm_75")
     filename, data = device_functions_cubin
     input_type = InputType.CUBIN.value
     _nvjitlinklib.add_data(handle, input_type, data, filename)
-    with pytest.raises(RuntimeError,
-                       match="NVJITLINK_ERROR_INTERNAL error"):
+    with pytest.raises(RuntimeError, match="NVJITLINK_ERROR_INTERNAL error"):
         _nvjitlinklib.get_linked_cubin(handle)
     _nvjitlinklib.destroy(handle)
 
@@ -168,14 +163,14 @@ def test_get_linked_cubin_from_lto(device_functions_ltoir):
     # requested. So we need to use the OBJECT input type, and the linker
     # retrieves the LTO IR from it because we passed the -lto flag.
     input_type = InputType.OBJECT.value
-    handle = _nvjitlinklib.create('-arch=sm_75', '-lto')
+    handle = _nvjitlinklib.create("-arch=sm_75", "-lto")
     _nvjitlinklib.add_data(handle, input_type, data, filename)
     _nvjitlinklib.complete(handle)
     cubin = _nvjitlinklib.get_linked_cubin(handle)
     _nvjitlinklib.destroy(handle)
 
     # Just check we got something that looks like an ELF
-    assert cubin[:4] == b'\x7fELF'
+    assert cubin[:4] == b"\x7fELF"
 
 
 def test_get_linked_ptx_from_lto(device_functions_ltoir):
@@ -185,19 +180,18 @@ def test_get_linked_ptx_from_lto(device_functions_ltoir):
     # requested. So we need to use the OBJECT input type, and the linker
     # retrieves the LTO IR from it because we passed the -lto flag.
     input_type = InputType.OBJECT.value
-    handle = _nvjitlinklib.create('-arch=sm_75', '-lto', '-ptx')
+    handle = _nvjitlinklib.create("-arch=sm_75", "-lto", "-ptx")
     _nvjitlinklib.add_data(handle, input_type, data, filename)
     _nvjitlinklib.complete(handle)
-    ptx = _nvjitlinklib.get_linked_ptx(handle)
+    _nvjitlinklib.get_linked_ptx(handle)
     _nvjitlinklib.destroy(handle)
 
 
 def test_get_linked_ptx_link_not_complete_error(device_functions_ltoir):
-    handle = _nvjitlinklib.create('-arch=sm_75', '-lto', '-ptx')
+    handle = _nvjitlinklib.create("-arch=sm_75", "-lto", "-ptx")
     filename, data = device_functions_ltoir
     input_type = InputType.OBJECT.value
     _nvjitlinklib.add_data(handle, input_type, data, filename)
-    with pytest.raises(RuntimeError,
-                       match="NVJITLINK_ERROR_INTERNAL error"):
+    with pytest.raises(RuntimeError, match="NVJITLINK_ERROR_INTERNAL error"):
         _nvjitlinklib.get_linked_ptx(handle)
     _nvjitlinklib.destroy(handle)
diff --git a/pynvjitlink/tests/test_pynvjitlink_api.py b/pynvjitlink/tests/test_pynvjitlink_api.py
new file mode 100644
index 00000000..807a0711
--- /dev/null
+++ b/pynvjitlink/tests/test_pynvjitlink_api.py
@@ -0,0 +1,156 @@
+# Copyright (c) 2023, NVIDIA CORPORATION. All rights reserved.
+
+import pytest
+import sys
+import os
+
+from pynvjitlink import NvJitLinker, NvJitLinkError
+
+
+@pytest.fixture
+def device_functions_cubin():
+    test_dir = os.path.dirname(os.path.abspath(__file__))
+    cubin_path = os.path.join(test_dir, "test_device_functions.cubin")
+    with open(cubin_path, "rb") as f:
+        return f.read()
+
+
+@pytest.fixture
+def device_functions_fatbin():
+    test_dir = os.path.dirname(os.path.abspath(__file__))
+    fatbin_path = os.path.join(test_dir, "test_device_functions.fatbin")
+    with open(fatbin_path, "rb") as f:
+        return f.read()
+
+
+@pytest.fixture
+def undefined_extern_cubin():
+    test_dir = os.path.dirname(os.path.abspath(__file__))
+    fatbin_path = os.path.join(test_dir, "undefined_extern.cubin")
+    with open(fatbin_path, "rb") as f:
+        return f.read()
+
+
+def test_create_no_arch_error():
+    # nvlink expects at least the architecture to be specified.
+    with pytest.raises(NvJitLinkError, match="NVJITLINK_ERROR_MISSING_ARCH error"):
+        NvJitLinker()
+
+
+def test_invalid_arch_error():
+    # sm_XX is not a valid architecture
+    with pytest.raises(
+        NvJitLinkError, match="NVJITLINK_ERROR_UNRECOGNIZED_OPTION error"
+    ):
+        NvJitLinker("-arch=sm_XX")
+
+
+def test_invalid_option_type_error():
+    with pytest.raises(TypeError, match="Expecting only strings"):
+        NvJitLinker("-arch", 53)
+
+
+def test_create_and_destroy():
+    nvjitlinker = NvJitLinker("-arch=sm_53")
+    assert nvjitlinker.handle != 0
+
+
+def test_add_cubin(device_functions_cubin):
+    nvjitlinker = NvJitLinker("-arch=sm_75")
+    name = "test_device_functions.cubin"
+    nvjitlinker.add_cubin(device_functions_cubin, name)
+
+
+def test_add_incompatible_cubin_arch_error(device_functions_cubin):
+    nvjitlinker = NvJitLinker("-arch=sm_70")
+    name = "test_device_functions.cubin"
+    with pytest.raises(NvJitLinkError, match="NVJITLINK_ERROR_INVALID_INPUT error"):
+        nvjitlinker.add_cubin(device_functions_cubin, name)
+
+
+def test_add_fatbin_sm75(device_functions_fatbin):
+    nvjitlinker = NvJitLinker("-arch=sm_75")
+    name = "test_device_functions.fatbin"
+    nvjitlinker.add_fatbin(device_functions_fatbin, name)
+
+
+def test_add_fatbin_sm70(device_functions_fatbin):
+    nvjitlinker = NvJitLinker("-arch=sm_70")
+    name = "test_device_functions.fatbin"
+    nvjitlinker.add_fatbin(device_functions_fatbin, name)
+
+
+def test_add_incompatible_fatbin_arch_error(device_functions_fatbin):
+    nvjitlinker = NvJitLinker("-arch=sm_80")
+    name = "test_device_functions.fatbin"
+    with pytest.raises(NvJitLinkError, match="NVJITLINK_ERROR_INVALID_INPUT error"):
+        nvjitlinker.add_fatbin(device_functions_fatbin, name)
+
+
+def test_add_cubin_with_fatbin_error(device_functions_fatbin):
+    nvjitlinker = NvJitLinker("-arch=sm_75")
+    name = "test_device_functions.fatbin"
+    with pytest.raises(NvJitLinkError, match="NVJITLINK_ERROR_INVALID_INPUT error"):
+        nvjitlinker.add_cubin(device_functions_fatbin, name)
+
+
+def test_add_fatbin_with_cubin_error(device_functions_cubin):
+    nvjitlinker = NvJitLinker("-arch=sm_75")
+    name = "test_device_functions.cubin"
+    with pytest.raises(NvJitLinkError, match="NVJITLINK_ERROR_INVALID_INPUT error"):
+        nvjitlinker.add_fatbin(device_functions_cubin, name)
+
+
+def test_duplicate_symbols_cubin_and_fatbin(
+    device_functions_cubin, device_functions_fatbin
+):
+    # This link errors because the cubin and the fatbin contain the same
+    # symbols.
+    nvjitlinker = NvJitLinker("-arch=sm_75")
+    name = "test_device_functions.cubin"
+    nvjitlinker.add_cubin(device_functions_cubin, name)
+    name = "test_device_functions.fatbin"
+    with pytest.raises(NvJitLinkError, match="NVJITLINK_ERROR_INVALID_INPUT error"):
+        nvjitlinker.add_fatbin(device_functions_fatbin, name)
+
+
+def test_get_linked_cubin_complete_empty_error():
+    nvjitlinker = NvJitLinker("-arch=sm_75")
+    cubin = nvjitlinker.get_linked_cubin()
+
+    # Linking nothing still gives us an empty ELF back
+    assert cubin[:4] == b"\x7fELF"
+
+
+def test_get_linked_cubin(device_functions_cubin):
+    nvjitlinker = NvJitLinker("-arch=sm_75")
+    name = "test_device_functions.cubin"
+    nvjitlinker.add_cubin(device_functions_cubin, name)
+    cubin = nvjitlinker.get_linked_cubin()
+
+    # Just check we got something that looks like an ELF
+    assert cubin[:4] == b"\x7fELF"
+
+
+def test_get_error_log(undefined_extern_cubin):
+    nvjitlinker = NvJitLinker("-arch=sm_75")
+    name = "undefined_extern.cubin"
+    nvjitlinker.add_cubin(undefined_extern_cubin, name)
+    with pytest.raises(NvJitLinkError):
+        nvjitlinker.get_linked_cubin()
+    error_log = nvjitlinker.error_log
+    assert "Undefined reference to '_Z5undefff'" in error_log
+
+
+def test_get_info_log(device_functions_cubin):
+    nvjitlinker = NvJitLinker("-arch=sm_75")
+    name = "test_device_functions.cubin"
+    nvjitlinker.add_cubin(device_functions_cubin, name)
+    nvjitlinker.get_linked_cubin()
+    info_log = nvjitlinker.info_log
+    # Info log is empty
+    assert "" == info_log
+
+
+if __name__ == "__main__":
+    sys.exit(pytest.main())
diff --git a/pyproject.toml b/pyproject.toml
index fcf1792d..5c380377 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -1,18 +1,17 @@
 [tool.scikit-build]
 cmake.minimum-version = "3.26.4"
+cmake.verbose = true
 ninja.make-fallback = true
 build-dir = "build/{wheel_tag}"
-# Workaround for
-# https://github.com/scikit-build/scikit-build-core/issues/437
-sdist.exclude = [".gitignore"]
+wheel.packages = ["pynvjitlink"]
 
 [build-system]
 requires = ["scikit-build-core"]
 build-backend = "scikit_build_core.build"
 
 [project]
-name = "nvjitlink"
-version = "0.1.0"
+name = "pynvjitlink-cu12"
+version = "0.1.6"
 description = "nvJitLink Python binding"
 readme = { file = "README.md", content-type = "text/markdown" }
 authors = [
@@ -20,3 +19,9 @@ authors = [
 ]
 license = { text = "Apache 2.0" }
 requires-python = ">=3.9"
+
+[tool.setuptools]
+license-files = ["LICENSE"]
+
+[tool.setuptools.dynamic]
+version = {file = "pynvjitlink/VERSION"}
diff --git a/test_binary_generation/Makefile b/test_binary_generation/Makefile
index 4d772556..3671d566 100644
--- a/test_binary_generation/Makefile
+++ b/test_binary_generation/Makefile
@@ -1,4 +1,4 @@
-# Generates the input files used by the nvjitlink binding test suite
+# Generates the input files used by the pynvjitlink binding test suite
 
 NVCC_FLAGS := -O3 -rdc true
 CUBIN_FLAGS := -gencode arch=compute_75,code=sm_75 --cubin
@@ -9,7 +9,7 @@ OBJECT_FLAGS := -gencode arch=compute_75,code=compute_75 -dc
 LIBRARY_FLAGS := -gencode arch=compute_75,code=compute_75 -lib
 
 
-OUTPUT_DIR := ../nvjitlink/tests
+OUTPUT_DIR := ../pynvjitlink/tests
 
 all:
 	nvcc $(NVCC_FLAGS) $(CUBIN_FLAGS) -o $(OUTPUT_DIR)/undefined_extern.cubin undefined_extern.cu
diff --git a/test_binary_generation/test.cu b/test_binary_generation/test.cu
index c7a37953..00b52061 100644
--- a/test_binary_generation/test.cu
+++ b/test_binary_generation/test.cu
@@ -2,14 +2,10 @@
 
 extern __device__ bool __heq(__half arg1, __half arg2);
 
-__device__
-__half test_add_fp16(__half arg1, __half arg2)
-{
+__device__ __half test_add_fp16(__half arg1, __half arg2) {
   return __hadd(arg1, arg2);
 }
 
-__device__
-bool test_cmp_fp16(__half arg1, __half arg2)
-{
+__device__ bool test_cmp_fp16(__half arg1, __half arg2) {
   return __heq(arg1, arg2);
 }
diff --git a/test_binary_generation/undefined_extern.cu b/test_binary_generation/undefined_extern.cu
index 86b07d85..9e0a0cad 100644
--- a/test_binary_generation/undefined_extern.cu
+++ b/test_binary_generation/undefined_extern.cu
@@ -1,6 +1,3 @@
 extern __device__ float undef(float a, float b);
 
-__global__ void f(float* r, float *a, float *b)
-{
-  r[0] = undef(a[0], b[0]);
-}
+__global__ void f(float *r, float *a, float *b) { r[0] = undef(a[0], b[0]); }