CUDA 13.4 Update 1 brings useful fixes, but the real upgrade work lies in reconciling drivers, framework wheels, build scripts and edge deployments.
NVIDIA released CUDA Toolkit 13.4 Update 1, package version 13.4.1, on September 9, 2026, and refreshed the release notes on September 15. The early problems teams are reporting are less about the toolkit itself and more about the seams around it: drivers, framework wheels, older build scripts, and edge platforms.
This guide starts with those reports, then moves through a pre-upgrade checklist, the known issues most relevant to AI workloads, and the release background for teams that want it.
The short answer
The public threads reviewed for this guide don’t point to a broadly broken toolkit. They do show five recurring sources of friction: a driver that is no longer included with the toolkit, PyTorch wheels that lag the installed CUDA version by design, stale library wheels that can cause an empty import error, dependencies that still hard-code GPU targets CUDA 13 removed, and a narrower prerelease path for Jetson users.
None of these issues are unusual on their own, but they can still cost a team an afternoon when the upgrade plan doesn’t account for them.
What early adopters are running into
We looked at the public threads from the 13.4 Developer Preview in July through the first week after GA. The problems appear in five main places.
1. The driver is no longer bundled, and the minimum version moved
Starting with CUDA 13.4 on Linux and CUDA 13.1 on Windows, the toolkit installer no longer includes the NVIDIA driver. You install it separately, and the release notes now state a minimum driver branch rather than a minimum driver version. For CUDA 13.4, that branch is R615. NVIDIA’s CUDA 13.4 Update 1 release notes document both the separate driver installation and the R615 branch requirement.
The change is straightforward in theory, but it moves a question that the installer used to answer onto the developer or platform team. A NVIDIA Developer Forum thread on CUDA 13.x driver compatibility remained active in September as developers asked how to interpret the new branch-only requirement. Earlier release notes listed a distinct minimum driver version for each point release, while the newer notes identify the required driver branch.
There’s also a related trap for anyone testing Rubin. The CUTLASS 4.8 development notes state that running Rubin, or sm_107, kernels requires an R615 driver, and that the R610 driver that shipped alongside the CUDA 13.4 Developer Preview isn’t sufficient. Teams that set up Rubin porting branches in July on the preview driver will need to update it before those kernels run.
For teams that build container images, this may be a Dockerfile change. For teams with developers installing on laptops and workstations, it adds another dependency that needs to be documented and checked.
2. PyTorch wheels lag the toolkit, and always have
A PyTorch forum user upgraded to CUDA 13.3 and found that the newest official wheel targeted CUDA 13.2. The maintainer’s response applies to every toolkit version: pip wheels ship their own CUDA runtime and do not use the toolkit installed on the machine. If a team needs a newer toolkit for a build, it needs to build from source.
Installing CUDA 13.4 won’t, on its own, make a pip-installed PyTorch package faster or slower. It affects what the team can compile locally, but it does not change the CUDA runtime bundled with an official wheel.
3. Leftover library wheels break a 13.4 build with an empty error
A PyTorch CI fix from September 10 described a libcublasLt.so.13 from a CUDA 13.0 wheel lingering in site-packages after a cleanup step removed torch but not its nvidia-* dependencies. PyTorch preloads component wheels it finds there, so a build against CUDA 13.4 headers referenced a grouped-GEMM symbol, cublasLtGroupedMatrixLayoutCreate, that the CUDA 13.0 cuBLAS library did not export. import torch then failed with an empty ImportError.
If import torch fails with an empty ImportError after upgrading, check for stray nvidia-* wheels before changing driver settings or compiler flags.
4. CUDA 13 dropped pre-Turing architectures, and older build scripts may not know it
CUDA 13 removed offline compilation and library support for Maxwell, Pascal, and Volta architectures. NVIDIA’s CUDA 13 documentation notes that applications can still be built for those targets with CUDA 12.x, but newer toolkits can no longer target them.
An NVIDIA Cosmos pull request from August hit nvcc fatal: Unsupported gpu architecture 'compute_70' because a pinned version of Apex hard-coded compute_70 into its own gencode list while ignoring TORCH_CUDA_ARCH_LIST. The fix was to update Apex to a commit that derives architectures from the detected CUDA version.
Other dependencies that have not been updated for CUDA 13 may fail in the same way. This is not a CUDA 13.4 defect; CUDA 13 removed support for these targets, and some codebases are only encountering the change now.
5. Edge and Jetson users are on a narrower path
A detailed JetPack 7.2 write-up for Orin boards running CUDA 13.2 identifies three constraints likely to remain relevant as teams plan for CUDA 13.4: the cu132 aarch64 wheels were prerelease and invisible to pip without --pre, Python 3.12 was the only interpreter with matching wheels, and the tested Python 3.10 wheels did not work. A freshly flashed board also had the CUDA runtime installed but not on PATH, which could look like a broken installation until the environment was checked.
The Python limitation also aligns with the CUDA 13.4 Update 1 release notes, which deprecate Python 3.10 across CUDA Python packages. If images still pin Python 3.10, the edge ecosystem is showing where the desktop stack is likely headed.
A note on CUDA Tile
CUDA 13.4 extends CUDA Tile with Programmatic Dependent Launch support and new strided and gather-scatter views in the C++ API. The Tile programming model has drawn comparisons with OpenAI’s Triton since it appeared, and a founding member of the original CUDA team has publicly suggested that countering Triton was one motivation.
An April 2026 evaluation of CUDA Tile on Hopper and Blackwell found that Triton supports all recent NVIDIA architectures while CUDA Tile does not yet, which is a portability trade-off to weigh before committing kernels to it. NVIDIA’s own TileGym benchmarks show the two approaches trading wins on fused attention depending on sequence length.
If the team already has Triton kernels in production, CUDA 13.4 does not fundamentally change the trade-off. For new Blackwell or Rubin work, weigh CUDA Tile’s platform fit against Triton’s broader portability. A comparison using your own kernels and workloads will be more useful than either project’s benchmark results.
A checklist before you upgrade
Each item below maps to one of the reports above or to a known issue in the release notes.
Confirm the driver branch. nvidia-smi should report a 615-series driver on Linux before you install the 13.4 toolkit. On an older driver, you’ll be running under minor version compatibility, and the new 13.4 features will not be available.
If you set up Rubin porting on the Developer Preview, move off the R610 driver. sm_107 kernels need R615.
Don’t expect pip install torch to change. Official wheels carry their own runtime. If you need 13.4 in PyTorch today, use the cu134 nightly index or build from source, and plan for the stable index later.
Check site-packages for nvidia-* wheels from earlier CUDA versions before building anything against 13.4 headers. An empty ImportError on import torch is one symptom.
Search build scripts and third-party dependencies for compute_70, sm_70, and other pre-Turing targets. CUDA 13 does not compile them, and some dependencies ignore TORCH_CUDA_ARCH_LIST.
On Grace Hopper, Grace Blackwell, or Vera Rubin systems, decide on CDMM versus NUMA memory mode before you touch the driver. 13.4 changes the default, and switching back is a node-wide setting that needs a reboot.
If you use grouped GEMM with per-batch scales, audit for groups with k = 0 and pass valid scale pointers. This is an open issue.
If Compute Sanitizer runs in CI, add an allowlist entry for the documented cp.async false positive so real failures do not get lost in the noise.
Plan the Python 3.10 retirement in any image that uses CUDA Python packages.
On Jetson, wait for the JetPack release that carries the toolkit version you want, and stay on Python 3.12 unless a wheel for another interpreter is published.
Known issues and fixes to read before you upgrade
The CUDA Toolkit 13.4 Update 1 release notes are extensive. The following entries are the most likely to affect AI workloads, based on the cuBLAS, cuSOLVER, cuFFT, and image-library sections.
Still open in Update 1
Compute Sanitizer (up to 13.4) reports false-positive out-of-bounds reads inside cuBLAS kernels that use cp.async with a source size smaller than the copy size. The hardware zero-pads the remainder, so these reports can be ignored. Your CI will not know that unless you tell it.
cuBLASLt grouped GEMM with per-batch tensor-wide scales causes an invalid memory access for groups where k = 0. The workaround is to always pass valid scale pointers for every group. This has been present since 13.1.
cuSPARSE mixed-precision CSR and COO SpMM isn’t supported in some cases.
Fixed in Update 1
cuBLASLt grouped GEMM with CUBLAS_POINTER_MODE_HOST could return incorrect results on Hopper GPUs when the wave count exceeded 2. This was introduced in 13.2 Update 1.
cuBLASLt grouped GEMM with NVFP4 inputs could return incorrect results on B300 and Rubin GPUs under a specific algorithm configuration. Introduced in 13.4 GA.
cublasLtMatmul() split-K could return incorrect output when the split count did not evenly divide K, on compute capability 9.0, 10.x, and 11.x. This one dates back to CUDA 12.6 Update 2.
cublasLtMatmul() with NVFP4 inputs could produce incorrect results on B300 and Rubin under algorithm 66 with a particular stage setting. Introduced in 13.3.
cuSOLVER routines could hang inside a CUDA Green Context because the handle cached the physical SM count at creation time.
nvJPEG could crash or hang on malformed JPEG bitstreams. Relevant to any pipeline that decodes user-uploaded images.
Applications built against an older CUDA Runtime (back to 11.2) running on the R615 driver could print benign but alarming error messages about IMEX channels.
Kernels using the fabric.try_put.tensor, fabric.try_red.tensor, or fabric.try_atom PTX instructions did not set the fabric-write grid attribute, which could leave fabric writes unfenced when the kernel ran inside a CUDA graph.
Fixed in 13.4 GA
cuFFT real-side LTO callback kernels for R2C and C2R transforms had a correctness bug for large even sizes, first identified in 13.3 Update 1.
NPP's nppiNV12ToRGB color-twist conversion could render black NV12 frames as blue. This regression was introduced in CUDA 12.9 and matters to anyone doing video preprocessing.
CUDA Math fixed fp128 off-by-one results in edge cases, with a 2x to 8x speedup for fp128 operations on GB200 as a side effect.
Several entries follow the same pattern: an issue is introduced in one CUDA 13.x point release and fixed two or three releases later. Teams should review release notes even for an upgrade that appears incremental, such as moving from CUDA 13.3 to CUDA 13.4.
Background: what’s actually in this release
This section is for teams that want to understand the toolkit changes behind the upgrade.
Update 1 versus 13.4 GA
If the environment already runs CUDA 13.4 GA, the Update 1 delta is PTX ISA 9.4 support, the runtime and compiler fixes listed above, and library fixes in cuBLAS 13.8.0.4, cuSOLVER 12.3.4.7, and nvJPEG 13.2.3.58. Most other “New Features” sections in the Update 1 notes list “None.” That is typical for an update release and means teams that already validated CUDA 13.4 GA have a smaller delta to review.
The broader CUDA 13.4 feature set arrived with the September 9 release: Windows on Arm support, a Rubin GPU architecture preview with compute capability 10.7, Multi-Process Service V3, and faster grouped GEMMs for mixture-of-experts workloads. NVIDIA’s CUDA 13.4 release announcement provides the full feature overview.
Any CUDA 13.x application can run on R580-and-later drivers under minor-version compatibility. New CUDA 13.4 features and newly enabled platforms require R615 or later.
What's new in 13.4 that AI teams will notice
Faster grouped GEMMs for mixture-of-experts models. cuBLASLt now schedules grouped GEMM computations dynamically across streaming multiprocessors on Blackwell data center GPUs, to smooth out the load imbalance that is typical of MoE routing. NVIDIA reports improvements of up to 20% for grouped GEMM calls with a large number of groups (32 is the example given), with additional gains when grouped GEMMs run concurrently with other kernels. If you serve or train MoE models on B200 or B300 hardware, this is the most direct performance reason to upgrade.
Emulated FP64 gets faster, and now runs on Rubin. cuBLAS has been building out fixed-point emulation of double-precision matrix multiplication since 13.0 Update 2. In 13.4, the library switches between the Ozaki-I and Ozaki-II schemes based on which is faster, and NVIDIA quotes up to 175 TFLOPS of emulated DGEMM on B200 and up to 212 TFLOPS on Rubin. Update 1 adds a further 25 percent to emulated ZGEMM peak performance on Rubin and moves the emulation workspace into a cuBLAS-managed memory pool. This matters more for scientific computing than for transformer inference, but it is relevant to any AI team with a physics or finance workload sharing the same cluster.
Multi-Process Service V3. MPS V3 replaces the old control interface with a scriptable CLI, named server instances, TOML configuration files, SM partition controls, and GPU memory limits that integrate with cgroups. For platform teams running many small inference jobs on a shared GPU, this is a real quality-of-life change.
CUDA Python and CCCL. cuda.core 1.1.0 adds texture and surface programming, NUMA-aware managed memory, and .pyi type stubs for every public API. The stubs are a small item with an outsized effect: IDE autocompletion and coding agents can finally see signatures and return types. cuda.compute 1.1 adds ahead-of-time compilation of CCCL algorithms for multiple GPU architectures, on build machines that have no GPU at all, and a serialize() and deserialize() pair for shipping the result. CCCL 3.4 includes a warp-specialized cub::DeviceScan for Blackwell that NVIDIA benchmarks at up to 92% memory-bandwidth utilization, up from around 50 percent in the earlier implementation.
Rubin preview and Windows on Arm. CUDA 13.4 adds a functional preview of the Rubin architecture (compute capability 10.7, target sm_107). NVIDIA is explicit that the preview is for porting and testing, not production or benchmarking. Windows on Arm is now a supported CUDA development platform, both natively and by cross-compiling from x86-64.
A memory-management default change on coherent platforms. On Grace Hopper, Grace Blackwell, and Vera Rubin systems, the driver now defaults to Coherent Driver-based Memory Management instead of NUMA mode. NUMA mode is still supported, but it is a node-wide kernel module setting that needs a driver reload or reboot, and NVIDIA recommends choosing it before you upgrade rather than after.
Also in the box. Compute Fabric Transport (a driver-level API for communication-library authors; NVIDIA points most application teams to NCCL or NVSHMEM instead), programmatic locality domains, a cudaMemGetLocationInfo API to query where unified memory currently resides, GCC 16 and Clang 22 host compiler support, Nsight Python 1.0, and a --pytorch=functions-trace-shapes option in Nsight Systems that records tensor shapes alongside traced functions.
Should your team upgrade to CUDA 13.4 Update 1 now?
It depends on what you run and what you build. Here is how the trade-offs fall for common situations.
<div style="width:100%; margin:28px 0; overflow:hidden; border:1px solid rgba(255,255,255,0.22); border-radius:14px; background:rgba(255,255,255,0.025);"> <table style="width:100%; table-layout:fixed; border-collapse:collapse; font-family:inherit; color:#ffffff; font-size:15px; line-height:1.45;"> <colgroup> <col style="width:22%;"> <col style="width:22%;"> <col style="width:23%;"> <col style="width:33%;"> </colgroup> <thead> <tr> <th style="padding:16px 22px; text-align:left; vertical-align:middle; font-size:16px; font-weight:700; line-height:1.25; background:rgba(255,255,255,0.08); border-right:1px solid rgba(255,255,255,0.18); border-bottom:1px solid rgba(255,255,255,0.22);">Your situation</th> <th style="padding:16px 22px; text-align:left; vertical-align:middle; font-size:16px; font-weight:700; line-height:1.25; background:rgba(255,255,255,0.08); border-right:1px solid rgba(255,255,255,0.18); border-bottom:1px solid rgba(255,255,255,0.22);">What 13.4 gives you</th> <th style="padding:16px 22px; text-align:left; vertical-align:middle; font-size:16px; font-weight:700; line-height:1.25; background:rgba(255,255,255,0.08); border-right:1px solid rgba(255,255,255,0.18); border-bottom:1px solid rgba(255,255,255,0.22);">What it costs you</th> <th style="padding:16px 22px; text-align:left; vertical-align:middle; font-size:16px; font-weight:700; line-height:1.25; background:rgba(255,255,255,0.08); border-bottom:1px solid rgba(255,255,255,0.22);">Reasonable timing</th> </tr> </thead> <tbody> <tr> <td style="padding:20px 22px; text-align:left; vertical-align:top; font-weight:600; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">PyTorch or JAX inference on Hopper or Blackwell, using official wheels</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Nothing until cu134 stable wheels ship</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">A driver update to R615 with no framework benefit yet</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-bottom:1px solid rgba(255,255,255,0.14);">Update the driver when convenient; move the toolkit when your framework does</td> </tr> <tr> <td style="padding:20px 22px; text-align:left; vertical-align:top; font-weight:600; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Custom CUDA or CUTLASS kernels, MoE serving on B200 or B300</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Up to 20 percent faster grouped GEMMs, NVFP4 correctness fixes</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Re-validation of kernels against new cuBLAS heuristics</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-bottom:1px solid rgba(255,255,255,0.14);">Now, behind a canary deployment</td> </tr> <tr> <td style="padding:20px 22px; text-align:left; vertical-align:top; font-weight:600; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Preparing for Rubin hardware</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">sm_107 target, R615 driver, Rubin-enabled math libraries</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Preview status; not for production or benchmarking</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-bottom:1px solid rgba(255,255,255,0.14);">Now, in a separate porting branch</td> </tr> <tr> <td style="padding:20px 22px; text-align:left; vertical-align:top; font-weight:600; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Shared GPU clusters with many small jobs</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">MPS V3 with cgroup memory limits and SM partitioning</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Rewriting MPS automation to the new CLI and TOML config</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-bottom:1px solid rgba(255,255,255,0.14);">Plan a quarter; test on one node first</td> </tr> <tr> <td style="padding:20px 22px; text-align:left; vertical-align:top; font-weight:600; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Grace Hopper or Grace Blackwell systems</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Locality domains, unified memory residency queries</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">CDMM becomes the default memory mode; NUMA mode needs a reboot to restore</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-bottom:1px solid rgba(255,255,255,0.14);">Decide the memory mode before touching the driver</td> </tr> <tr> <td style="padding:20px 22px; text-align:left; vertical-align:top; font-weight:600; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Jetson and edge deployments</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Tracks JetPack, not the desktop toolkit</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14); border-bottom:1px solid rgba(255,255,255,0.14);">Prerelease wheels, Python 3.12 only in practice</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-bottom:1px solid rgba(255,255,255,0.14);">Wait for the JetPack release that carries it</td> </tr> <tr> <td style="padding:20px 22px; text-align:left; vertical-align:top; font-weight:600; border-right:1px solid rgba(255,255,255,0.14);">Windows on Arm laptops for developers</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14);">Native CUDA development for the first time</td> <td style="padding:20px 22px; text-align:left; vertical-align:top; border-right:1px solid rgba(255,255,255,0.14);">New platform; expect early gaps in third-party libraries</td> <td style="padding:20px 22px; text-align:left; vertical-align:top;">Now, for developer machines; hold on production tooling</td> </tr> </tbody> </table> </div>
If you land in the first row, which most AI product teams do, 13.4 is a driver update to make now and a toolkit update to schedule.
Where FullStack fits
Toolkit upgrades like this one are rarely a single engineer's job. The driver rollout usually belongs to one team and the container images to another, and the coordination between them is where the afternoons disappear. FullStack's nearshore AI engineering teams work inside client environments on exactly this kind of platform work, from validating grouped GEMM changes on Blackwell to rebuilding CI images for a new toolkit. If your team is planning a 13.4 or Rubin migration and wants engineers who have done it before, we are happy to talk through the plan.
Learn more
Frequently Asked Questions
What is the difference between CUDA 13.4 and CUDA 13.4 Update 1?
13.4 GA (September 9, 2026) added Windows on Arm, the Rubin preview, MPS V3, and the library performance work. Update 1 (13.4.1) adds PTX ISA 9.4 and a set of correctness fixes, mostly in cuBLAS and cuSOLVER. The two share the R615 driver branch.
Which NVIDIA driver do I need for CUDA 13.4?
The R615 branch or later for full 13.4 functionality. Existing 13.x applications continue to run on drivers from R580 onward. The driver is no longer included in the Linux toolkit installer, so download it separately.
Why does `import torch` fail with an empty ImportError after upgrading to 13.4?
The most common cause reported so far is a leftover nvidia-* wheel from an earlier CUDA version in site-packages. PyTorch preloads those libraries, and an older cuBLAS does not export symbols that a 13.4 build expects. Remove the stale wheels and reinstall.
Does CUDA 13.4 support the Rubin GPU?
As a preview. Compute capability 10.7 and the sm_107 target are available for porting and functional testing. NVIDIA states the preview is not for production deployment or benchmarking. Rubin kernels require the R615 driver; the R610 driver from the Developer Preview is not sufficient.
Does CUDA 13.4 work with PyTorch?
Yes, with the usual caveat. Official PyTorch wheels carry their own CUDA runtime and do not use your installed toolkit. As of mid-September 2026, cu134 was available in the PyTorch nightly index but not the stable index. Building PyTorch from source against 13.4 works; installing the 13.4 toolkit does not change a pip-installed torch.
AI is changing software development.
The Engineer's AI-Enabled Development Handbook is your guide to incorporating AI into development processes for smoother, faster, and smarter development.
Enjoyed the article? Get new content delivered to your inbox.
Subscribe below and stay updated with the latest developer guides and industry insights.
Thank you! Your submission has been received!
Oops! Something went wrong while submitting the form.
Iframe is blocked. Accept cookies to load it.
Thank you! Your submission has been received!
Oops! Something went wrong while submitting the form.
We use cookies to provide our services, to allow us to better understand our audience, and to provide and serve personalized ads or content. By using our website, you consent to the terms of our Privacy Policy and our Cookie Policy, and the use of cookies, pixels, and other technology as described more fully therein