skills.rest
shakfu
shakfuCOMMUNITY
Explore all skills by this author →
GitHub

mojo-gpu-fundamentals

Guides writing Mojo GPU kernels for NVIDIA, AMD, and Apple accelerators.

Updated Sep 15, 2026
Search Tags:#kernel#mojo#cuda#simd#accelerator#gpu programming#tiletensor
One-click install
npx skills add https://github.com/shakfu/mdsp --skill mojo-gpu-fundamentals-shakfu
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: mojo-gpu-fundamentals
Source: https://github.com/shakfu/mdsp/tree/main/.claude/skills/mojo-gpu-fundamentals
Command: npx skills add https://github.com/shakfu/mdsp --skill mojo-gpu-fundamentals-shakfu

SYSTEM DOCUMENTATION & REQUIREMENTS

What problem does it solve? Pretrained models default to CUDA syntax and outdated Mojo APIs when writing GPU code, producing kernels that fail to compile. This Skill corrects those misconceptions with the actual Mojo GPU programming model, covering TileTensor, DeviceContext, kernel launch, shared memory, and warp operations. ## Core Features & Use Cases - CUDA-to-Mojo mapping: Translates familiar CUDA concepts (__global__, cudaMalloc, __syncthreads, threadIdx) into their Mojo equivalents like plain def kernels, ctx.enqueue_create_buffer, barrier(), and global_idx. - TileTensor and layout guidance: Explains row_major layouts, tiling, vectorization, rebind for cross-layout element mismatches, and the mandatory comptime assert flat_rank pattern for tensor indexing. - Complete working examples: Provides full vector-addition and tiled shared-memory matmul programs, plus reduction, benchmarking, and architecture-detection patterns. - Use Case: You ask an AI to write a Mojo kernel that offloads a matrix multiplication to an NVIDIA GPU; this Skill ensures the generated code uses enqueue_function with bound comptime parameters instead of invalid CUDA-style launch syntax. ## Quick Start Write a Mojo GPU kernel that adds two float32 vectors on my GPU using the mojo-gpu-fundamentals conventions.

Frequently Asked Questions about mojo-gpu-fundamentals

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I write a GPU kernel in Mojo?▼

Mojo GPU kernels are plain functions with no decorator or special return type. Define a function taking TileTensor arguments, launch it with ctx.enqueue_function[kernel](args, grid_dim=..., block_dim=...), and index threads with global_idx or thread_idx from max.gpu.

What is the Mojo equivalent of CUDA syntax like __global__ and cudaMalloc?▼

Mojo has no CUDA syntax. __global__ becomes a plain def function, cudaMalloc becomes ctx.enqueue_create_buffer[dtype](count), cudaMemcpy becomes ctx.enqueue_copy, __syncthreads becomes barrier(), and threadIdx.x becomes thread_idx.x.

Does Mojo GPU programming support AMD and Apple GPUs?▼

Yes, Mojo targets NVIDIA, AMD, and Apple silicon GPUs. Use has_accelerator() from host code to check for any GPU, and is_nvidia_gpu(), is_amd_gpu(), or is_apple_gpu() inside GPU-compiled code to dispatch architecture-specific paths.

Why does my Mojo kernel fail with 'no matching method' on enqueue_function?▼

This happens when the kernel has comptime parameters like a TensorLayout trait parameter. Bind them first with comptime kernel = my_kernel[type_of(layout)], then pass the bound symbol to ctx.enqueue_function.

How do I use shared memory in a Mojo GPU kernel?▼

Allocate shared memory inside the kernel with stack_allocation from the layout package, passing address_space=AddressSpace.SHARED and a row_major layout. It returns a TileTensor; call barrier() after writing before other threads read the data.

Why do I get 'lacking evidence to prove correctness' when indexing a TileTensor?▼

Subscripting a TileTensor requires a comptime assert on its rank, such as comptime assert tensor.flat_rank == 2, in any function that indexes it. Derived tensors from tile() or vectorize() need their own re-assert before indexing.

Related Skills

Explore More Software Engineering Skills→
better-mojobetter-mojo
6

mojo-gpu-fundamentals

Translate CUDA kernel definitions and launches into Mojo GPU programming constructs.

Official
Intermediate
Harmeet10000Harmeet10000
7

mojo-gpu-fundamentals

Explain CUDA-free Mojo GPU programming with LayoutTensor and shared memory patterns.

Community
Advanced
alibabaalibaba
71

gpu-kernel-bottleneck-analysis

Analyzes GPU kernel performance using Roofline models and profiling data.

Official
Advanced
tensormuxtensormux
54

Port CUDA Kernel to HIP

Port CUDA kernels to HIP for ROCm-enabled AMD GPUs.

Official
Advanced